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

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

 

很简单，就是分别求出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个整数读入共享存储器，位置一一对应，和线程号也一一对应，所以可以同时完成。

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

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

 

代码如下：

 

#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;
	extern __shared__ int smem[];
	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.
    addKernel<<<1, size, size*sizeof(int), 0>>>(dev_c, dev_a);

    // 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;
}

从代码中看到执行配置<<<>>>中第三个参数为共享内存大小（字节数），这样我们就知道了全部4个执行配置参数的意义。恭喜，你的CUDA终于入门了！