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

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

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

不会飞的纸飞机

扫一扫二维码，了解我的更多动态。