CUDA学习笔记（2）- 线程并行和块并行

---

1. 获取显卡设备信息

有些显卡支持CUDA有些不支持，那么如何确定主机的显卡设备是否支持CUDA呢。可以使用下面的函数获取显卡的相关信息。

• cudaError_t cudaGetDeviceCount(int *count) 获取支持CUDA的显卡设备数量。
• cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) 获取显卡设备的相关属性信息，如显卡名称、显示的最大核心数目、显存大小等信息。
• cudaSetDevice(int device) 设置当前使用的显卡设备。

下面是关于这几个函数的简单使用：

cudaError_t cudaStatus;
int number = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&number);
for (int i = 0; i < number; ++i)
{
	cudaGetDeviceProperties(&prop, i);
	cudaStatus = cudaSetDevice(0);
	if (cudaStatus == cudaSuccess)
	{
		printf("Used GPU %s\n", prop.name);
	}
}

我的显卡为GeForce GTX 960M，所以会打印下面的信息
Used GPU GeForce GTX 960M

2. 线程并行

一个线程块中，可以含有多个线程。本例中将两个数组中的对应数字相加，获取相加后的数据，代码如下：

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

// 线程并行
void myTestCalc(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);

	addKernel<<<1, 5>>>(pDevDataC, pDevDataA, pDevDataB);
	cudaThreadSynchronize();
	cudaMemcpy(pDataC, pDevDataC, sizeof(int) * 5, cudaMemcpyDeviceToHost);

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

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

这里因为传入的线程数为一维的，所以threadIdx.x就可以直接获取当前线程的所在索引号。程序调用函数 myTestCalc() 后，程序执行结果为：
Thread Cala Result is: 12, 24, 36, 48, 60

3. 块并行

使用块并行与上面的代码类似，代码如下:

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

// 块并行
void myTestCalc(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);

	//addKernel<<<1, 5>>>(pDevDataC, pDevDataA, pDevDataB);
	addKernel <<<5, 1 >>>(pDevDataC, pDevDataA, pDevDataB);
	cudaThreadSynchronize();
	cudaMemcpy(pDataC, pDevDataC, sizeof(int) * 5, cudaMemcpyDeviceToHost);

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

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

上面代码相比线程并行的不同之处在于

1. 将threadIdx.x替换为blockIdx.x，来表示当前块的ID。
2. 调用部分修改为addKernel <<<5, 1 >>>(pDevDataC, pDevDataA, pDevDataB);，表示使用5个线程块来处理，每个线程块中有一个线程。