3、线程同步
3.1、共享内存
共享内存位于芯片内部，因此它比全局内存快得多。（CUDA里面存储器的快慢有两方面，一个是延迟低，一个是带宽大。这里特指延迟低），相比没有经过缓存的全局内存访问，共享内存大约在延迟上低100倍。同一个块中的线程可以访问相同的一段共享内存（注意：不同块中的线程所见到的共享内存中的内容是不相同的），这在许多线程需要与其他线程共享它们的结果的应用程序中非常有用。但是如果不同步，也可能会造成混乱或错误的结果。如果某线程的计算结果在写入到共享内存完成之前被其他线程读取，那么将会导致错误。因此，应该正确地控制或管理内存访问。这是由__syncthreads()指令完成的，该指令确保在继续执行程序之前完成对内存的所有写入操作。这也被称为barrier。barrier的含义是块中的所有线程都将到达该代码行，然后在此等待其他线程完成。当所有线程都到达了这里之后，它们可以一起继续往下执行。
#include <stdio.h> __global__ void gpu_shared_memory(float *d_a) { // Defining local variables which are private to each thread int i, index = threadIdx.x; float average, sum = 0.0f; //Define shared memory __shared__ float sh_arr[10]; sh_arr[index] = d_a[index]; __syncthreads(); // This ensures all the writes to shared memory have completed for (i = 0; i<= index; i++) { sum += sh_arr[i]; } average = sum / (index + 1.0f); d_a[index] = average; sh_arr[index] = average; } int main(int argc, char **argv) { //Define Host Array float h_a[10]; //Define Device Pointer float *d_a; for (int i = 0; i < 10; i++) { h_a[i] = i; } // allocate global memory on the device cudaMalloc((void **)&d_a, sizeof(float) * 10); // now copy data from host memory to device memory cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 10, cudaMemcpyHostToDevice); gpu_shared_memory << <1, 10 >> >(d_a); // copy the modified array back to the host memory cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost); printf("Use of Shared Memory on GPU: \n"); //Printing result on console for (int i = 0; i < 10; i++) { printf("The running average after %d element is %f \n", i, h_a[i]); } return 0; }
在main函数中，当分配好主机和设备上的数组后，用0.0到9.0填充主机上的数组，然后将这个数组复制到显存。内核将对显存中的数据进行读取，计算并保存结果。最后结果从显存中传输到内存，然后在控制台上输出。控制台上的输出结果如图所示：
这个程序还含有额外的一个CUDA函数调用：cudaDeviceSynchronize()。为何要加这句？这是因为启动内核是一个异步操作，只要发布了内核启动命令，不等内核执行完成，控制权就会立刻返回给调用内核的CPU线程。在上述的代码中，CPU线程返回，继续执行的下一句是printf()。而再之后，在内核完成之前，进程就会结束，终止控制台窗口。所以，如果不加上这句同步函数，你就看不到任何的内核执行结果输出。在程序退出后内核生成的输出结果，将没有地方可去，你没法看到它们，因此，如果我们不包含这个指令，你将不会看到任何内核执行的printf语句的输出结果。要能看到内核生成的输出结果，我们必须包含这句同步函数。这样，内核的结果将通过可用的标准输出显示，而应用程序则会在内核执行完成之后才退出。
3.2、原子操作
考虑当大量的线程需要试图修改一段较小的内存区域的情形，这是（在日常的算法实现中）常发生的现象。当我们试图进行“读取-修改-写入”操作序列的时候，这种情形经常会带来很多麻烦。
一个例子是代码d_out[i]++，这代码首先将d_out[i]的原值从存储器中读取出来，然后执行了+1操作，再将结果回写到存储器。然而，如果多个线程试图在同一个内存区域中进行这个操作，则可能会得到错误的结果。
假设某内存区域中有初始值6，两个线程p和q分别试图将这段区域中的内容+1，则最终的结果应当是8。但是在实际执行的时候，可能p和q两个线程同时读取了这个初始值，两者都得到了6，执行+1操作都得到了7，然后它们将7写回这个内存区域。这样，和正确的结果8不同，我们得到的最终结果是7，这是错误的。这种错误是如何的危险，我们通过ATM取现操作来演示。假设你的账户余额为5000卢比，你的账户下面开了两张银行卡，你和你的朋友同时去2个不同的ATM上取现4000卢比，你俩在同一瞬间刷卡取现。所以，当两个ATM检查余额的时候，都将显示5000卢比的余额。当你俩同时取现4000卢比的时候，两个ATM机都只根据初始值5000卢比判断，要取的现金4000卢比小于当前余额。所以两个机器将会给你们每人4000卢比。即使你之前只有5000卢比的余额，你们也能得到8000卢比，这很危险。为了示范一下这种情形，做了一个很多线程试图同时访问一个小数组的例子：
#include <stdio.h> #define NUM_THREADS 10000 #define SIZE 10 #define BLOCK_WIDTH 100 __global__ void gpu_increment_without_atomic(int *d_a) { // Calculate thread id for current thread int tid = blockIdx.x * blockDim.x + threadIdx.x; // each thread increments elements wrapping at SIZE variable tid = tid % SIZE; d_a[tid] += 1; } int main(int argc, char **argv) { printf("%d total threads in %d blocks writing into %d array elements\n", NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE); // declare and allocate host memory int h_a[SIZE]; const int ARRAY_BYTES = SIZE * sizeof(int); // declare and allocate GPU memory int * d_a; cudaMalloc((void **)&d_a, ARRAY_BYTES); //Initialize GPU memory to zero cudaMemset((void *)d_a, 0, ARRAY_BYTES); gpu_increment_without_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a); // copy back the array to host memory cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost); printf("Number of times a particular Array index has been incremented without atomic add is: \n"); for (int i = 0; i < SIZE; i++) { printf("index: %d --> %d times\n ", i, h_a[i]); } cudaFree(d_a); return 0; }
可能如同已经猜到的那样，每次运行你的程序，每个内存区域中的元素值都可能会不同。这是设备上不定顺序的多线程执行导致的。
为了解决这个问题，CUDA提供了atomicAdd这种原子操作函数。该函数会从逻辑上保证，每个调用它的线程对相同的内存区域上的“读取旧值-累加-回写新值”操作是不可被其他线程扰乱的原子性的整体完成的。使用atomicAdd进行原子累加的内核函数代码如下：
#include <stdio.h> #define NUM_THREADS 10000 #define SIZE 10 #define BLOCK_WIDTH 100 __global__ void gpu_increment_atomic(int *d_a) { // Calculate thread id for current thread int tid = blockIdx.x * blockDim.x + threadIdx.x; // each thread increments elements wrapping at SIZE variable tid = tid % SIZE; atomicAdd(&d_a[tid], 1); } int main(int argc, char **argv) { printf("%d total threads in %d blocks writing into %d array elements\n", NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE); // declare and allocate host memory int h_a[SIZE]; const int ARRAY_BYTES = SIZE * sizeof(int); // declare and allocate GPU memory int * d_a; cudaMalloc((void **)&d_a, ARRAY_BYTES); //Initialize GPU memory to zero cudaMemset((void *)d_a, 0, ARRAY_BYTES); gpu_increment_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a); // copy back the array to host memory cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost); printf("Number of times a particular Array index has been incremented is: \n"); for (int i = 0; i < SIZE; i++) { printf("index: %d --> %d times\n ", i, h_a[i]); } cudaFree(d_a); return 0; }
在main函数中，具有10个元素的数组被初始化成0值，然后传递给了内核，但现在，内核中的代码将执行原子累加操作。所以，这个程序输出的结果将是对的，数组中的每个元素将被累加1000。运行结果显示如图：
如果你测量一下这个程序的运行时间，相比之前的那个简单地在全局内存上直接进行加法操作的程序它用的时间更长。这是因为使用原子操作后程序具有更大的执行代价。可以通过使用共享内存来加速这些原子累加操作。如果线程规模不变，但原子操作的元素数量扩大，则这些同样次数的原子操作会更快地完成。这是因为更广泛的分布范围上的原子操作有利于利用多个能执行原子操作的单元，以及每个原子操作单元上面的竞争性的原子事务也相应减少了。
