CUDA学习笔记（3）- 流并行和线程同步

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

不会飞的纸飞机

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