并非所有块都调用 CUDA 内核

Question

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

现在如果我删除 printf 函数，只执行第一个块，我得到 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 不会同步块。 atomicAdd 可以解决这个并发问题。

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