c++ - CUDA 内核未被所有 block 调用

Question

我在尝试运行简单的 vector 加法时遇到了一个奇怪的行为。 如果我使用 printf 函数运行下面的代码，一切运行正常，我得到了预期的结果，5050。

现在，如果我删除 printf 函数，则只执行第一个 block ，我得到 2080，这是总和达到 64 的预期结果。

有人知道这里发生了什么吗？

预先感谢您的帮助。

vecSum.cu:

#include <cuda.h>
#include <cuda_runtime.h>
#include <iostream>
#include <math.h>

#define BLOCK_SIZE 64

__global__
void vecSumKernel(int N, float *d_v, float *d_out)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int t = threadIdx.x;

    printf("Processing block #: %i\n", blockIdx.x);

    __shared__ float partialSum[BLOCK_SIZE];
    if(idx < N)
        partialSum[t] = d_v[idx];
    else
        partialSum[t] = 0;


    for(unsigned int stride=1; stride < BLOCK_SIZE; stride *= 2)
    {
        __syncthreads();
        if(t % (2*stride) == 0)
            partialSum[t] += partialSum[t + stride];
    }

    __syncthreads();
    *d_out += partialSum[0];
}

void vecSum_wrapper(int N, float *v, float &out, cudaDeviceProp devProp)
{
    float *d_v;
    float *d_out;
    size_t size = N*sizeof(float);

    cudaMalloc(&d_v, size);
    cudaMalloc(&d_out, sizeof(float));

    cudaMemcpy(d_v, v, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_out, &out, sizeof(float), cudaMemcpyHostToDevice);

    int nbrBlocks = ceil((float)N / (float)BLOCK_SIZE);
    vecSumKernel<<<nbrBlocks, BLOCK_SIZE>>>(N, d_v, d_out);

    cudaDeviceSynchronize();

    cudaMemcpy(&out, d_out, sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(d_v);
}

ma​​in.cpp:

int main()
{
    ...

    int N = 100;

    float *vec = new float[N];

    for(int i=0; i < N; ++i)
        vec[i] = i + 1;

    std::chrono::time_point<timer> start = timer::now();

    float result = 0;
    vecSum_wrapper(N, vec, result, devProp);

    std::cout << "Operation executed in " << std::chrono::duration_cast<chrono>(timer::now() - start).count() << " ms \n";

    std::cout << "Result: " << result << '\n';

    delete[] vec;

    return 0;
}

Answer 1

您的内核 *d_out += partialSum[0] 的最后一行似乎可能会暴露一些并发问题，因为您肯定知道 __syncthreads 不会同步块(synchronized block)。 atomicAdd可能会解决这个并发问题。

至于为什么它与printf一起工作得更好，我假设printf需要一些同步，因此 block 不会同时进入最后一条指令，但我没有什么可以证明的这个。