文章目錄

1. 流並行

執行緒流中可以有多個執行緒塊，執行緒塊中可以有多個執行緒。執行緒塊和執行緒流只能處理單個函式，執行緒流可以處理多個函式和同一個函式的不同引數。

• cudaStreamCreate(cudaStream_t *pStream) 建立一個執行緒流。
• cudaStreamDestroy(cudaStream_t stream) 銷燬執行緒流。

下面是關於流並行的簡單示例，效果同上一節的執行緒並行和塊並行：

__global__ void addKernel(int *c, int *a, int *b)
{
	int i = threadIdx.x;
	c[i] = a[
 
i] + b[i];
}

// 流並行
void myTestCalcStream(void)
{
	int pDataA[5] = { 1, 2, 3, 4, 5 };
	int pDataB[5] = { 11, 22, 33, 44, 55 };
	int pDataC[5] = { 0 };

	// 申請A、B、C的記憶體
	int *pDevDataA = nullptr, *pDevDataB = nullptr, *pDevDataC = nullptr;
	cudaMalloc(&pDevDataA, sizeof(int) * 5);
	cudaMalloc(&pDevDataB,
 
 sizeof(int) * 5);
	cudaMalloc(&pDevDataC, sizeof(int) * 5);

	// 記憶體拷貝
	cudaMemcpy(pDevDataA, pDataA, sizeof(int) * 5, cudaMemcpyHostToDevice);
	cudaMemcpy(pDevDataB, pDataB, sizeof(int) * 5, cudaMemcpyHostToDevice);

	cudaStream_t streams[5];
	for (int i = 0; i < 5; ++i)
		cudaStreamCreate(streams +
 
 i);

	for (int i=0; i<5; ++i)
		addKernel <<<1, 1, 0, streams[i]>>>(pDevDataC + i, pDevDataA + i, pDevDataB + i);

	cudaDeviceSynchronize();
	cudaThreadSynchronize();
	cudaMemcpy(pDataC, pDevDataC, sizeof(int) * 5, cudaMemcpyDeviceToHost);

	printf("Stream Cala Result is: %d, %d, %d, %d, %d\n", pDataC[0], pDataC[1], pDataC[2], pDataC[3], pDataC[4]);

	for (int i = 0; i < 5; ++i)
		cudaStreamDestroy(streams[i]);

	cudaFree(pDevDataA);
	cudaFree(pDevDataB);
	cudaFree(pDevDataC);
}

函式呼叫比之前的多兩個引數addKernel << <1, 1, 0, streams[i]>> > , 前兩個還表示執行緒塊的數目和每個執行緒塊中執行緒的個數，第三個引數表示每個塊的共享記憶體的大小，第四個引數為流。

2. 執行緒同步

• 使用函式 __syncthreads()，實現執行緒的同步
• __shared__ 表示共享記憶體

下面是關於執行緒同步的示例，示例中分別計算陣列的和、平方和、乘積

__global__ void addKernel2(int *src, int *dest)
{
	int threadIndex = threadIdx.x;
	extern __shared__ int sharedMemory[5];
	sharedMemory[threadIndex] = src[threadIndex];
	__syncthreads();

	if (threadIndex == 0)
	{
		src[threadIndex] = 0;
		for (int i = 0; i < 5; ++i)
			src[threadIndex] += sharedMemory[i];
	}
	else if (threadIndex == 1)
	{
		src[threadIndex] = 0;
		for (int i = 0; i < 5; ++i)
			src[threadIndex] += sharedMemory[i] * sharedMemory[i];
	}
	else if (threadIndex == 2)
	{
		src[threadIndex] = 1;
		for (int i = 0; i < 5; ++i)
			src[threadIndex] *= sharedMemory[i];
	}
}

void threadSyncTest(void)
{
	int srcArray[5] = { 1, 2, 3, 4, 5 };
	int *pDataSrc = nullptr;
	cudaMalloc(&pDataSrc, sizeof(int) * 5);
	cudaMemcpy(pDataSrc, srcArray, sizeof(int) * 5, cudaMemcpyHostToDevice);
	
	addKernel2<<<1, 5, sizeof(int) * 5, 0>>>(pDataSrc, pDataSrc);
	cudaThreadSynchronize();

	int destArray[3] = { 0 };
	cudaMemcpy(destArray, pDataSrc, sizeof(int) * 3, cudaMemcpyDeviceToHost);

	printf("The Thread is : %d, %d, %d\n", destArray[0], destArray[1], destArray[2]);
}

程式執行結果為：
The Thread is : 15, 55, 120