CUDA从入门到精通（九）：线程通信实例

最新推荐文章于 2024-05-26 12:43:02 发布

转载最新推荐文章于 2024-05-26 12:43:02 发布 · 1.3k 阅读

本文介绍了一个使用CUDA共享内存和线程同步技术的示例。该示例演示了如何计算一组固定数字的和、平方和及乘积。文章详细解释了共享内存的使用，包括其优势及如何在GPU上优化数据读取。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

接着上一节，我们利用刚学到的共享内存和线程同步技术，来做一个简单的例子。先看下效果吧：

很简单，就是分别求出1~5这5个数字的和，平方和，连乘积。相信学过C语言的童鞋都能用for循环做出同上面一样的效果，但为了学习CUDA共享内存和同步技术，我们还是要把简单的东西复杂化(^_^)。

简要分析一下，上面例子的输入都是一样的，1,2,3,4,5这5个数，但计算过程有些变化，而且每个输出和所有输入都相关，不是前几节例子中那样，一个输出只和一个输入有关。所以我们在利用CUDA编程时，需要针对特殊问题做些让步，把一些步骤串行化实现。

输入数据原本位于主机内存，通过cudaMemcpy API已经拷贝到GPU显存（术语为全局存储器，Global Memory），每个线程运行时需要从Global Memory读取输入数据，然后完成计算，最后将结果写回Global Memory。当我们计算需要多次相同输入数据时，大家可能想到，每次都分别去Global Memory读数据好像有点浪费，如果数据很大，那么反复多次读数据会相当耗时间。索性我们把它从Global Memory一次性读到SM内部，然后在内部进行处理，这样可以节省反复读取的时间。

有了这个思路，结合上节看到的SM结构图，看到有一片存储器叫做Shared Memory，它位于SM内部，处理时访问速度相当快（差不多每个时钟周期读一次），而全局存储器读一次需要耗费几十甚至上百个时钟周期。于是，我们就制定A计划如下：

线程块数：1，块号为0；（只有一个线程块内的线程才能进行通信，所以我们只分配一个线程块，具体工作交给每个线程完成）

线程数：5，线程号分别为0~4；（线程并行，前面讲过）

共享存储器大小：5个int型变量大小（5*sizeof(int））。

步骤一：读取输入数据。将Global Memory中的5个整数读入共享存储器，位置一一对应，和线程号也一一对应，所以可以同时完成。

步骤二：线程同步，确保所有线程都完成了工作。

步骤三：指定线程，对共享存储器中的输入数据完成相应处理。

代码如下：

[cpp]view plaincopy 
   
print?
 #include "cuda_runtime.h"  
 #include "device_launch_parameters.h"  
   
 #include <stdio.h>  
   
 cudaError_t addWithCuda(int *c, const int *a, size_t size);  
   
 __global__ void addKernel(int *c, const int *a)  
 {  
     int i = threadIdx.x;  
 <span style="font-size:24px;"><strong>  extern __shared__ int smem[];</strong>  
 </span> smem[i] = a[i];  
     __syncthreads();  
     if(i == 0)  //0号线程做平方和  
     {  
         c[0] = 0;  
         for(int d = 0;d<5;d++)  
         {  
             c[0] += smem[d]*smem[d];  
         }  
     }  
     if(i == 1)//1号线程做累加  
     {  
         c[1] = 0;  
         for(int d = 0;d<5;d++)  
         {  
             c[1] += smem[d];  
         }  
     }  
     if(i == 2)  //2号线程做累乘  
     {  
         c[2] = 1;  
         for(int d = 0;d<5;d++)  
         {  
             c[2] *= smem[d];  
         }  
     }  
 }  
   
 int main()  
 {  
     const int arraySize = 5;  
     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
     int c[arraySize] = { 0 };  
     // Add vectors in parallel.  
     cudaError_t cudaStatus = addWithCuda(c, a, arraySize);  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "addWithCuda failed!");  
         return 1;  
     }  
     printf("\t1+2+3+4+5 = %d\n\t1^2+2^2+3^2+4^2+5^2 = %d\n\t1*2*3*4*5 = %d\n\n\n\n\n\n", c[1], c[0], c[2]);  
     // cudaThreadExit must be called before exiting in order for profiling and  
     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
     cudaStatus = cudaThreadExit();  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaThreadExit failed!");  
         return 1;  
     }  
     return 0;  
 }  
   
 // Helper function for using CUDA to add vectors in parallel.  
 cudaError_t addWithCuda(int *c, const int *a,  size_t size)  
 {  
     int *dev_a = 0;  
     int *dev_c = 0;  
     cudaError_t cudaStatus;  
   
     // Choose which GPU to run on, change this on a multi-GPU system.  
     cudaStatus = cudaSetDevice(0);  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
         goto Error;  
     }  
   
     // Allocate GPU buffers for three vectors (two input, one output)    .  
     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaMalloc failed!");  
         goto Error;  
     }  
   
     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaMalloc failed!");  
         goto Error;  
     }  
     // Copy input vectors from host memory to GPU buffers.  
     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaMemcpy failed!");  
         goto Error;  
     }  
     // Launch a kernel on the GPU with one thread for each element.  
 <span style="font-size:24px;"><strong>    addKernel<<<1, size, size*sizeof(int), 0>>>(dev_c, dev_a);</strong>  
 </span>  
     // cudaThreadSynchronize waits for the kernel to finish, and returns  
     // any errors encountered during the launch.  
     cudaStatus = cudaThreadSynchronize();  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
         goto Error;  
     }  
   
     // Copy output vector from GPU buffer to host memory.  
     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
     if (cudaStatus != cudaSuccess)   
     {  
         fprintf(stderr, "cudaMemcpy failed!");  
         goto Error;  
     }  
   
 Error:  
     cudaFree(dev_c);  
     cudaFree(dev_a);      
     return cudaStatus;  
 }  