CUDA程序将两个整数相加不起作用

Question

这是 Jason Sanders 和 Edward Kandrot 所著《CUDA By Examples》一书中的示例程序“simple_kernel_params.cu”。

#include <cstdio>

static void HandleError(cudaError_t err, const char *file, int line)
{
    if (err != cudaSuccess)
    {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

__global__ void add(int a, int b, int *c)
{
    *c = a + b;
}

int main(void)
{
    int c;
    int *dev_c;
    HANDLE_ERROR(cudaMalloc((void **)&dev_c, sizeof(int)));
    add<<<1, 1>>>(2, 7, dev_c);
    HANDLE_ERROR(cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost));
    printf("2 + 7 = %d\n", c); // prints "2 + 7 = 0"
    HANDLE_ERROR(cudaFree(dev_c));
    return 0;
}

输出应该是

2 + 7 = 9

，但实际上是

2 + 7 = 0

。我知道一个问题可能是 NVIDIA 驱动程序尚未初始化，因此这里是

sudo nvidia-smi -a

的（删节的）输出。 （我无法粘贴整个输出，因为 StackOverflow 抱怨“您的帖子主要是代码；请添加更多详细信息”）

==============NVSMI LOG==============

Driver Version                            : 535.104.12
CUDA Version                              : 12.2

Attached GPUs                             : 1
GPU 00000000:00:1E.0
    Product Name                          : Tesla T4
    Product Brand                         : NVIDIA
    Product Architecture                  : Turing
    Display Mode                          : Disabled
    Display Active                        : Disabled
    Persistence Mode                      : Enabled
    Addressing Mode                       : None
    MIG Mode
        Current                           : N/A
        Pending                           : N/A
    Accounting Mode                       : Disabled
    Accounting Mode Buffer Size           : 4000
    Driver Model
        Current                           : N/A
        Pending                           : N/A
    Serial Number                         : 1563820001764
    GPU UUID                              : GPU-c4636f1c-e8b5-a646-1c6d-14eb842ae426
    Minor Number                          : 0
    VBIOS Version                         : 90.04.96.00.02
    MultiGPU Board                        : No
    Board ID                              : 0x1e
    Board Part Number                     : 900-2G183-0000-001
    GPU Part Number                       : 1EB8-895-A1
    FRU Part Number                       : N/A
    Module ID                             : 1
    Inforom Version
        Image Version                     : G183.0200.00.02
        OEM Object                        : 1.1
        ECC Object                        : 5.0
        Power Management Object           : N/A
    GPU Operation Mode
        Current                           : N/A
        Pending                           : N/A
    GSP Firmware Version                  : 535.104.12
    GPU Virtualization Mode
        Virtualization Mode               : Pass-Through
        Host VGPU Mode                    : N/A
    GPU Reset Status
        Reset Required                    : No
        Drain and Reset Recommended       : N/A
    IBMNPU
        Relaxed Ordering Mode             : N/A

为什么我的程序不能运行？

Answer 1

add<<<1, 1>>>(2, 7, dev_c);

异步执行。

由于启动 GPU 代码在设置设备时会产生一些开销（大约 25 微秒），因此它在单独的 CPU 线程中启动并在 GPU 完成之前返回。

参见示例：

Cuda：内核启动队列

有关内核启动的哪些部分是同步的、哪些部分是异步的解释。
您需要在异步设备调用后添加

cudaDeviceSynchronize()

，以便CPU等待您的GPU代码完成。如果您这样做，它就会按预期工作。

它还有助于在设备代码中添加

printf

语句，这样您就可以看到那里发生了什么。

我已经在 Godbolt 中测试了

更新的

代码：https://cuda.godbolt.org/z/q1xTGvYKa 在那里工作得很好。
固定代码：

#include <cuda.h> //not needed, but helps to clarify that this is cuda code
#include <cstdio>

static void HandleError(cudaError_t err, const char *file, int line) {
    if (err != cudaSuccess) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

__global__ void add(const int a, const int b, int *c) {
    *c = a + b;
//#ifdef _DEBUG
    printf("device: *c: %i = a: %i + b: %i\n", *c, a, b); //optional: diagnostics on device
//#endif
}

int main(void) {
    int c;
    int *dev_c;
    static_assert(sizeof(c) == sizeof(*dev_c));
    HANDLE_ERROR(cudaMalloc((void **)&dev_c, sizeof(int)));
    add<<<1, 1>>>(2, 7, dev_c);
    HANDLE_ERROR(cudaDeviceSynchronize()); //wait for device to finish
    HANDLE_ERROR(cudaMemcpy(&c, dev_c, sizeof(c), cudaMemcpyDeviceToHost));
    printf("2 + 7 = %d\n", c); // sometimes, but not always prints "2 + 7 = 0", if you leave out the cudaDeviceSynchronize() 
    HANDLE_ERROR(cudaFree(dev_c));
    return 0;
}