CUDA線程協作之共享存儲器“shared”&&“__syncthreads()”

2019-11-14 09:07:33

字體：大中小

來源：轉載

供稿：網友

在GPU并行編程中，一般情況下，各個處理器都需要了解其他處理器的執行狀態，在各個并行副本之間進行通信和協作，這涉及到不同線程間的通信機制和并行執行線程的同步機制。

共享內存“share”

CUDA中的線程協作主要是通過共享內存實現的。使用關鍵字“__share__”聲明共享變量，將使這個變量駐留在共享內存中，該變量具有以下特征：位于線程塊的共享存儲器空間中與線程塊具有相同的生命周期僅可通過塊內的所有線程訪問對于GPU上啟動的每個線程塊，CUDA C編譯器都將創建該變量的一個副本。線程塊中的每個線程都共享這塊內存，但線程卻無法看到也不能修改其他線程塊的變量副本。這就使得一個線程塊中的多個線程能夠在計算上進行通信和協作。而且，共享內存緩沖區駐留在物理GPU上，在訪問共享內存時的延遲要遠遠低于訪問普通緩沖區的延遲，使得共享內存的訪問非常高效。

線程同步機制“__syncthreads()”

關鍵字“__share__”只是聲明了共享變量，位于同一個線程塊中的不同線程都可以訪問該變量，如果沒有同步機制，將會發生競態條件 (Race Condition)，導致錯誤的運行結果。CUDA確保同步的方法是調用“__syncthreads()”。__syncthreads()將確保線程塊中的每個線程都執行完 __syncthreads()前面的語句后，才會執行下一條語句。以下是CUDA和OpenCV的應用中，繪制一幅圖像，Grid的尺寸大小是60*60，Block的尺寸大小是10*10，在各個線程塊內聲明了一個共享變量sharedMem：

#include "cuda_runtime.h"    #include <highgui.hpp>    using namespace cv;#define DIM 600   //圖像長寬#define PI 3.1415926535897932f  __global__ void kernel(unsigned char *ptr){	// map from blockIdx to pixel position    	int x = threadIdx.x + blockIdx.x * blockDim.x;	int y = threadIdx.y + blockIdx.y * blockDim.y;	int offset = x + y * blockDim.x * gridDim.x;	__shared__  float sharedMem[16][16];	const float period = 128.0f;	sharedMem[threadIdx.x][threadIdx.y] =		255 * (sinf(x*2.0f*PI / period) + 1.0f) *		(sinf(y*2.0f*PI / period) + 1.0f) / 4.0f;	__syncthreads();	ptr[offset * 3 + 0] = 0;	ptr[offset * 3 + 1] = sharedMem[15 - threadIdx.x][15 - threadIdx.y];	ptr[offset * 3 + 2] = 0;}// globals needed by the update routine    struct DataBlock{	unsigned char   *dev_bitmap;};int main(void){	DataBlock   data;	cudaError_t error;	Mat image = Mat(DIM, DIM, CV_8UC3, Scalar::all(0));	data.dev_bitmap = image.data;	unsigned char    *dev_bitmap;	error = cudaMalloc((void**)&dev_bitmap, 3 * image.cols*image.rows);	data.dev_bitmap = dev_bitmap;	dim3    grid(DIM / 10, DIM / 10);	dim3   block(10, 10);	//DIM*DIM個線程塊  	kernel << <grid, block >> > (dev_bitmap);	error = cudaMemcpy(image.data, dev_bitmap,		3 * image.cols*image.rows,		cudaMemcpyDeviceToHost);	error = cudaFree(dev_bitmap);	imshow("__share__ and __syncthreads()", image);	waitKey();}如果線程間不加入__syncthreads()同步機制，同一線程塊內不同線程訪問sharedMem，獲取的結果可能是不一樣的，生成的圖像如下，有散亂的雜點：加入__syncthreads()同步機制，保證了同一線程塊中不同的線程都執行完成__syncthreads()這個集合點之前的部分之后，才繼續往下執行，所以不同的線程訪問sharedMem獲取的結果是一致的，圖像無雜散點，是一個規律的排布：