514
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終於入門了！
最後更新：2017-04-03 16:48:43