|
|
|---|---|
|
|
When A, B, and the carry-in are all HIGH, a sum of 1 and a carry-out are produced. First, consider A and B. When both are HIGH, the output of gate 1 is LOW, and the output of gate 2 is HIGH, giving us a carry-out at gate 5. The carry-in produces a 1 output at gate 3, giving us a sum of 1. The output of the full adder is 112. The sum of 12 plus 12 plus 12 is 11
PARALLEL ADDERS The adders discussed in the previous section have been limited to adding single-digit binary numbers and carries. The largest sum that can be obtained using a full adder is 112. Parallel adders let us add multiple-digit numbers. If we place full adders in parallel, we can add two- or four-digit numbers or any other size desired. Figure 3-9 uses STANDARD SYMBOLS to show a parallel adder capable of adding two, two-digit binary numbers. In previous discussions we have depicted circuits with individual logic gates shown. Standard symbols (blocks) allow us to analyze circuits with inputs and outputs only. One standard symbol may actually contain many and various types of gates and circuits. The addend would be input on the A inputs (A2 = MSD, A1 = LSD), and the augend input on the B inputs (B2 = MSD, B1 = LSD). For this explanation we will assume there is no input to C0 (carry from a previous circuit).  Figure 3-9. —Parallel binary adder. Now let’s add some two-digit numbers. To add 102 (addend) and 012 (augend), assume there are numbers at the appropriate inputs. The addend inputs will be 1 on A2 and 0 on A1. The augend inputs will be 0 on B2 and 1 on B1. Working from right to left, as we do in normal addition, let’s calculate the outputs of each full adder. With A1 at 0 and B1 at 1, the output of adder 1 will be a sum (S1) of 1 with no carry (C1). Since A2 is 1 and B2 is 0, we have a sum (S2) of 1 with no carry (C2) from adder 1. To determine the sum, read the outputs (C2, S 2, and S1) from left to right. In this case, C2 = 0, S2 = 1, and S1 = 1. The sum, then, of 102 and 012 is 0112 or 112. To add 112 and 012, assume one number is applied to A1 and A2, and the other to B1 and B2, as shown in figure 3-10. Adder 1 produces a sum (S1) of 0 and a carry (C1) of 1. Adder 2 gives us a sum (S2) Source from http://www.tpub.com/content/neets/14185/css/14185_126.htm |
|
|
加法和減法的運算，可以在一個加法器中結合起來，做成一個二進位的並行加法器
。要達成此一目的，就在每一個全加器中增加一個Exclusive-OR閘。下圖所示為一4-bit 的加減法器。Ｓ輸入為運算控制；當Ｓ＝０時做加法運算，Ｓ＝１時做減法運算。  Source from http://www.cis.nctu.edu.tw/~info24/advanced/add_sub.htm |
|
|
|
|
|
呃，忘了关注密码学小组的帖子，原来rockinuk版主两天前就已经帮我找了这么多资料。正好三天假，我先下载下来好好研究下，看能不能有所启发！
感谢 
|
|
|
Parallelization - Parallel Adder
The following is a description of a basic single-processor code and its transition into a parallel code. We hope this example of a simple parallelization process serves to demonstrate what steps and calls are needed for codes whose work can be performed independently of one another. A simple "divide and conquer" approach is appropriate here. To the experienced parallel programmer, this example is a "trivially parallelizable" code because the entire work of the code can be partitioned between the processors very simply. That is the point of this discussion: we are demonstrating how one might go about converting this kind of single-processor code into a parallel code. This technique could apply to your situation in at least three possible ways. First, this technique is applicable to to problems such as minmax searches for the best move in a game such as chess, "monte carlo" techniques, or computationally-intensive searches through a large parameter space. If you are fortunate enough to have a code of this category, then you need to look little further than this technique. Second, selected components of a code could also fall under this category, so this technique could be useful to apply to those portions of a single-processor code. Finally, if this is your first look at parallel code, this example serves as an excellent start to learn how to go about writing and thinking about parallel codes. Before: Single-Processor adder  We begin with a code utilizing single-processor execution. This example, adder, sums the square of each integer from 1 to 10000, inclusive. It is organized into three routines: 1) kernelroutine - performs the elemental work of squaring each integer 2) computeloop - allocates memory, loops over the integers, calls kernelroutine, and saves and sums their results 3) main - performs any required initialization, calls computeloop, saves its results into a file, printf's the sum, and deallocates memory The code is structured in such a way to clearly separate the pieces that perform the work that is secondary yet essential to solving the problem. These include allocating the appropriate memory and saving data to an output file. Besides identifying where these functions are being performed, the structure makes it obvious that a much more complicated problem can substitute for the kernelroutine shown. We also intend this explicit structure to represent corresponding portions of a much larger code. The C source is shown in Listing 1. Listing 1 - adder.c
See also adder.f90
--------------------------------------------------------------------------------
#include <stdio.h> /* standard I/O routines */
#include <stdlib.h> /* standard library routines, such as memory */
#include <math.h> /* math libraries */
/* A routine performing elemental work.
This can be replaced with a much larger routine. */
double kernelroutine(double input);
double kernelroutine(double input)
{
return (input+1)*(input+1);
}
/* HighestIndex specifies how highest index to sample */
#define HighestIndex 10000L
void computeloop(double *theSum, double **theArray);
void computeloop(double *theSum, double **theArray)
{
/* local copies of the data to be output */
double *myArray, mySum=0;
/* limit of the loop */
long loopEnd=HighestIndex;
/* allocate an array to hold all the results */
myArray=malloc(sizeof(double)*loopEnd);
if (myArray) {
long index;
/* loop over indicies */
for(index=0; index<loopEnd; index++) {
/* call the kernel routine for each index, and save into the array */
myArray[index]=kernelroutine(index);
/* sum as we go */
mySum+=myArray[index];
}
}
/* return the sum and the array */
*theSum=mySum;
*theArray=myArray;
}
int main(int argc, char *argv[])
{
/* main copies of the sum and the array */
double theSum, *theArray=NULL;
printf("Beginning computation...\n");
computeloop(&theSum, &theArray);
if (theArray) {/* error checking */
FILE *fp;
/* save the array into a data file */
fp=fopen("output", "w");
if (fp) {
printf("writing array...\n");
fwrite(theArray, sizeof(double), HighestIndex, fp);
fclose(fp);
}
printf("the sum is %f\n", theSum);
/* clean up after myself */
free(theArray);
}
else
printf("memory allocation failure\n");
return 0;
}
When this code is run with its HighestIndex constant set to 10000, the code states that the sum is 333383335000 and produces a 80000 character long binary file containing the squares of integers 1 through 10000. After: paralleladder  A key to parallelizing an application is choosing the appropriate partitioning of the problem. The obvious choice here is to divide the number of integers by the number of processors. However, a few details must be worked out: What if the number of processors doesn't divide evenly into the number of integers? How does each processor know which integers to work on so they don't do the same integer twice or miss one? Once they have the answers, how do the sum numbers on one processor with those on another, and how do these processors save the data into one file? The parallelization process here is primarily a matter of managing and delegating data and computation. Changes to main are limited to "boilerplate" code that prepare and tear down the parallel environment and pass the sufficient information to computeloop so that it can organize the work and a if test so only processor 0 creates an output file. The most important modifications are in computeloop to perform the coordination and calculate the partitioning between processors. Given the identification number of the processor it is running on and the number of processors in the system, computeloop calculates how to partition the problem among the processors. To minimize bottlenecks, the executable running on each processor calculates its assignment itself, without a central authority delegating assignments. Once it finishes its work, it collects the data back to processor 0 so main can write the output in one convenient place for the user. The detailed answer is in the code. Listing 2 shows the C source of paralleladder.c, with the changes relative to adder.c underlined. Listing 2 - paralleladder.c, with changes compared to adder.c underlined
See also paralleladder.f90
--------------------------------------------------------------------------------
#include <stdio.h> /* standard I/O routines */
#include <stdlib.h> /* standard library routines, such as memory */
#include <math.h> /* math libraries */
#include "mpi.h" /* MPI library */
/* A routine performing elemental work.
This can be replaced with a much larger routine. */
double kernelroutine(double input);
double kernelroutine(double input)
{
return (input+1)*(input+1);
}
/* HighestIndex specifies how highest index to sample */
#define HighestIndex 10000L
void computeloop(double *theSum, double **theArray, int idproc, int nproc);
void computeloop(double *theSum, double **theArray, int idproc, int nproc)
{
/* local copies of the data to be output */
double *myArray, mySum=0;
/* limit of the loop */
long loopEnd=(HighestIndex+nproc-1)/nproc;
/* just this proc's piece */
/* this processor's index offset */
long offset=idproc*loopEnd;
/* allocate an array to hold all the results */
myArray=malloc(sizeof(double)*loopEnd);
if (myArray) {
long index;
/* loop over indicies */
for(index=0; index<loopEnd; index++) {
/* call the kernel routine for each index, and save into the array */
myArray[index]=kernelroutine(index+offset);
/* sum as we go */
if (index+offset<HighestIndex)
/* limit to the desired indicies */
mySum+=myArray[index];
}
/* proc 0 needs to hold the entire array */
{double *bigArray=idproc?NULL:malloc(sizeof(double)*loopEnd*nproc);
/* gathers the data from the other arrays ... */
MPI_Gather(myArray, loopEnd, MPI_DOUBLE,
bigArray, loopEnd, MPI_DOUBLE,
/* ... to proc 0 */ 0, MPI_COMM_WORLD);
if (!idproc) {
free(myArray);
myArray=bigArray;
}
}
/* performs a parallel sum across processors
and saves the result ... */
MPI_Reduce(&mySum, &mySum, 1, MPI_DOUBLE, MPI_SUM,
/* ... at proc 0 */ 0, MPI_COMM_WORLD);
}
/* return the sum and the array */
*theSum=mySum;
*theArray=myArray;
}
int ppinit(int argc, char *argv[], int *idproc, int *nproc);
void ppexit(void);
int main(int argc, char *argv[])
{
/* main copies of the sum and the array */
double theSum, *theArray=NULL;
int idproc, nproc, ierr;
/* initialize parallel processing */
ierr = ppinit(argc, argv, &idproc, &nproc);
if (ierr) return ierr; /* stop right there if there's a problem */
printf("I'm processor #%d in a %d-processor cluster.\n", idproc, nproc);
printf("Beginning computation...\n");
computeloop(&theSum, &theArray, idproc, nproc);
if (theArray) {/* error checking */
if (!idproc) {/* only processor 0 */
FILE *fp;
/* save the array into a data file */
fp=fopen("output", "w");
if (fp) {
printf("writing array...\n");
fwrite(theArray, sizeof(double), HighestIndex, fp);
fclose(fp);
}
printf("the sum is %f\n", theSum);
}
/* clean up after myself */
free(theArray);
}
else
printf("memory allocation failure\n");
/* only proc 0 pauses for user exit */
if (!idproc) {
printf("press return to continue\n");
getc(stdin);
}
ppexit();
return 0;
}
#ifdef __MWERKS__
/* only for Metrowerks CodeWarrior */
#include <SIOUX.h>
#endif
int ppinit(int argc, char *argv[], int *idproc, int *nproc)
{
/* this subroutine initializes parallel processing
idproc = processor id
nproc = number of real or virtual processors obtained */
int ierr;
*nproc=0;
/* initialize the MPI execution environment */
ierr = MPI_Init(&argc, &argv);
if (!ierr) {
/* determine the rank of the calling process in the communicator */
ierr = MPI_Comm_rank(MPI_COMM_WORLD, idproc);
/* determine the size of the group associated with the communicator */
ierr = MPI_Comm_size(MPI_COMM_WORLD, nproc);
#ifdef __MWERKS__
/* only for Metrowerks CodeWarrior */
SIOUXSettings.asktosaveonclose=0;
SIOUXSettings.autocloseonquit=1;
#endif
}
return ierr;
}
void ppexit(void)
{
/* this subroutine terminates parallel processing */
int ierr;
/* terminate MPI execution environment */
ierr = MPI_Finalize();
}
The changes from adder to paralleladder: mpi.h - the header file for the MPI library, required to access information about the parallel system and perform communication idproc, nproc - nproc describes how many processors are currently running this job and idproc identifies the designation, labeled from 0 to nproc - 1, of this processor. This information is sufficient to identify exactly which part of the problem this instance of the executable should work on. In this case, these variables are supplied to computeloop by main. loopEnd=(HighestIndex+nproc-1)/nproc; - loopEnd describes how many integers this processor should perform kernelroutine as it was in adder. Here, we choose to make have each processor perform the same amount of work, yet the total amount of work is at least that necessary to complete the problem. Depending on HighestIndex and nproc, the last processor might do a little too much, but that processor would end up waiting for the others to catch up if it did skip the excess work and it's only a small waste if HighestIndex is sufficiently large. offset=idproc*loopEnd; - By shifting the start of the sampling for this particular processor, offset is how each processor knows not to overlap work with other processors, without skipping the needed work. With this variable and loopEnd, this is sufficient information to specify the partition of work assigned to this processor. the malloc - performs the memory allocation for each processor. Except for processor 0, the memory allocation is the amount sufficient to hold the results of this processor's partition of work. In the case of processor 0, it must create an array big enough hold the results of its own work and that of every other processor to save the data later. if (index+offset<HighestIndex) - limit the summation to only the indicies we want in the sum MPI_COMM_WORLD - MPI defines communicator worlds or communicators that define a set of processors that can communicate with each other. At initialization, one communicator, MPI_COMM_WORLD, covers all the processors in the system. Other MPI calls can define arbitrary subsets of MPI_COMM_WORLD, making it possible to confine a code to a particular processor subset just by passing it the appropriate communicator. In simpler cases such as this, using MPI_COMM_WORLD is sufficient. MPI_Gather - fills bigArray of processor 0 with the data in myArray from the other processors in preparation for saving into one data file. myArray, loopEnd, and MPI_DOUBLE specify the first element, size, and data type of the array. 0 specifies where the data is to be collected. This call can be considered a "convenience" routine that simplifies the process of "gathering" regularly organized data from other processors. MPI_Reduce - computes the sum, specified by MPI_SUM, of each mySum variable of each processor and collects the result in processor 0. This call can sum an entire array of values, so mySum is passed as an array of length one. When supplied MPI_MAX or MPI_MIN, it can also perform other operations such as comparing values across processors and retaining the maximum or minimum values. ppinit - performs initialization of the parallel computing environment. Part of the boilerplate for MPI codes, this routine calls: MPI_Init - performs the actual initialization of MPI. It returns the error code; zero means no error. MPI_Comm_size - accesses the processor count of the parallel system MPI_Comm_rank - accesses the identification number of this particular processor ppexit - by calling MPI_Finalize, terminates the parallel computing environment. Also part of the boilerplate of MPI codes, a call to MPI_Finalize is needed to properly cleanup and close the connections between codes running on other processors and release control. __MWERKS__ - these #ifdef blocks merely compensate for a pecularity about the Metrowerks CodeWarrior compiler. These blocks allows the executable on other processors terminate automatically. If you are using another compiler, you don't have to worry about this code. Note that kernelroutine did not need to be altered at all. That reflects the independence of the work performed by this routine. The code that organizes and coordinates the work between processors is in computeloop. main's changes merely perform the boilerplate work of setting up and tearing down the parallel environment, plus an if statement so only processor 0 writes the output file. Something to note is that this code does not explicitly call MPI_Send or MPI_Recv or any other elemental communication calls. Instead it calls MPI_Gather and MPI_Reduce. There isn't anything wrong with using elemental versus collective calls, if either will do the job. In fact, the MPI_Gather could just as easily be replaced with the following: Listing 3 - an alternative to the above MPI_Gather call
--------------------------------------------------------------------------------
if (idproc)
MPI_Send(myArray, loopEnd, MPI_DOUBLE, 0, idproc, MPI_COMM_WORLD);
else {
long otherProc;
MPI_Status status;
for(index=0; index<loopEnd; index++) bigArray[index]=myArray[index];
for(otherProc=1; otherProc<nproc; otherProc++)
MPI_Recv(&bigArray[otherProc*loopEnd], loopEnd, MPI_DOUBLE,
otherProc, otherProc, MPI_COMM_WORLD, &status);
}
References Using MPI by William Gropp, Ewing Lusk and Anthony Skjellum is a good reference for the MPI library. In addition, many of the techniques and style of the above parallel code were adopted from Dr. Viktor K. Decyk. He expresses these ideas in a few references, including: V. K. Decyk, "How to Write (Nearly) Portable Fortran Programs for Parallel Computers", Computers In Physics, 7, p. 418 (1993). V. K. Decyk, "Skeleton PIC Codes for Parallel Computers", Computer Physics Communications, 87, p. 87 (1995). You're welcome to read more about parallel computing via the Tutorials page. sample Code: adder.c adder.f90 paralleladder.c paralleladder.f90 Source from http://daugerresearch.com/vault/paralleladder.shtml |
|
|
|
|
|
顺着这个 L2_Notes.pdf 說的找了一下，收获还不错。虽然并行加法上思路还不是很清晰（看了几个方法，里面讲的Carry-Save Addition、carry look-ahead addition似乎都是针对有限位数计算时进位的优化，感觉这么复杂的逻辑拿到GPU上不一定特别有用），不过找到个不错的并行乘法
 原文链接：Comba multiplication By Paul G. Comba 原理如下，设下划线表进位，我们一般做乘法： 2 3 x 8 9 ======= 2_0_7 1_8_4 ======= 2_0 4 7 Comba的计算方式：（保留进位，算完再加） 2 3
x 8 9
===========
18 27
16 24
===========
2 _0 _4 _7这个方法提出已经是20多年前了，后人还做了一个更好的改进。首先把两个数写在两张纸的角落，倒转其中一张纸和写个另一张纸对接，比如23倒转后是32，然后如下操作： |
|3 2
`----
----. => 3*9 = 27 => 2047
8 9| / |
| / |
<________/ |
| / carry 2 |
|3 2 / |
`---- / |
----. => 2 + 3*8 + 2*9 = 44
8 9| / |
| / |
<________/ |
| / carry 4 |
|3 2 / |
`---- / |
----. => 4 + 2*8 = 20
8 9|
|这样做的话可以减少乘法次数，对于两个大数相乘优势很大 顺便贴两个我找到的资料，都在这里汇总一下好了： Carry-Save Addition www.ece.tamu.edu/~sshakkot/courses/ecen248/csa-notes.pdf Conditional-Sum Addition Chap8 - Conditional Sum Adders |
|
|
并行加法上，借鉴Comba这个保留进位的思路，我想出一种不太完美的解决办法。
其实并行加法需要解决的问题就是carry往哪放的问题，高位计算是需要低位的进位才导致了无法并行。那么如果能把进位结果保存在当前位上，这个问题就解决了。比如计算 999 + 1： 9 9 9 + 0 0 1 ======== 9 9 0 + 1_0 <= 加上进位 ======== 9 0 0 + 1_0 <= 又进位了，这是最棘手的问题 解决办法，进位后把结果直接保留在该位上： 9 9 9 + 0 0 1 ========== 9 9 10 其结果上 9×100 ＋ 9×10 ＋ 10 = 1000 是一样的。 这个理论很简单，放到GPU中就是（假设计算 c[n] = a[n] + b[n] ）： 1. 并行计算 c = a + b，将进位结果存入u 当然这样c[i]保存数的范围会变成一个可变值，这会带来很多不方便的地方，我在考虑是否运用超前进位加法器的思路，当需要统一范围的c[i]时，有选择的在尽可能少的循环次数内把c[i] = c[i] + u[i-1]这里的进位给消除掉。这样兼顾了运算速度并且其他算法实现也会简单点（至少考虑乘法计算的溢出时少了一个不确定因素）。 再次感谢R大的资料，顺着参考资料、关键词搜索容易很多 
|
|
|
|
|
|
|
|
|
好的，我去找来看看先。
不是倍数也没有关系，比如512个线程计算一个300位的加法，那么让多出的212线程闲着好了，虽然没有最大限度利用GPU性能  用原始数据自行对齐？这个怎么做？刚学CUDA... |
|
|
|
|
|
所以在编程时要注意一些地方如； 还是没明白为什么最好不要用线程号  ... 上面那句话说的，应该是将矩阵线性化表示后，使用线程号对其遇到的问题（这也取决于具体算法）难道这样计算出来的idx会出现重复？ int idx = blockIdx.x*blockDim.x + threadIdx.x; if (idx<N) c[idx] = a[idx]+ b[idx]; 并且如果只有一个block时，threadIdx.x应该是唯一的。 share memory和global memory速度差异这个我知道一点，CSDN的CUDA版上看到过几个类似讨论。 目前CUDA大数库方面资料很少，国内我只搜到《GPU高性能运算之CUDA》作者OpenHero的的一个开发计划，其中有大数计算的事项，只不过好像还在开发中，没有进一步的消息 
|
|
|
|
|
|
|
