cuda - 从设备复制回主机时如何修复 cudaError 77

问题描述

我正在编写一个简单的示例程序来测试更大程序的 memCpy 和内核运行并发性。在编写此示例时，我偶然发现了错误 77，即 cudaErrorIllegalAddress。

我在某个地方读到了来自内核访问无效地址的内容，而不是 memcpy 本身。所以我试图索引我的输入数组（0）的最低元素。错误仍然存在。

由于它只是一个小示例程序，我将提供整个代码；

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

#define BLOCKS 32
#define THREADS 16

__global__ void kernel(double *d_in, double *d_out) {
    int index = threadIdx.x + blockDim.x * blockIdx.x;
    d_out[index] = d_in[index] + 5;
}

int main() {
    const int GPU_N = 2;
    const int data_size = 2048;
    const int cycles = 2;

    double *h_in, *h_out, *d_in, *d_out;

    h_in = (double*)malloc(sizeof(double) * data_size);
    h_out = (double*)malloc(sizeof(double) * data_size);

    for (int i = 0; i < data_size; i++) {
        h_in[i] = 21;
    }

    cudaError_t error;

    printf("1\n");
    for (int i = 0; i < cycles; i++) {
        //cuMalloc
        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            cudaMalloc((void**)&d_in, sizeof(double) * data_size / 4);
            cudaMalloc((void**)&d_out, sizeof(double) * data_size / 4);

            printf("2\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            cudaMemcpyAsync(d_in, h_in, sizeof(double) * data_size / 4, cudaMemcpyHostToDevice);
            printf("3\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            kernel<<< BLOCKS, THREADS, 0, 0 >>>(d_in, d_out);
            error = cudaGetLastError();
            printf("4\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            error = cudaMemcpyAsync(h_out, d_out, sizeof(double) * data_size / 4, cudaMemcpyDeviceToHost);
            printf("D2H %i\n", error);
            printf("5\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            cudaFree(d_in);
            cudaFree(d_out);
            printf("6\n");
        }
    }

    for (int i = 0; i < data_size; i++) {
        printf("%i\n", h_out[i]);
    }

    getchar();

}

所以输出应该是这样的：

然后是结果的垃圾邮件。它这样做直到它必须打印5，然后它输出error 77. 此外，结果的输出并不26像预期的那样，但是-842150451

标签： cudaruntime-error

解决方案

这段代码有几个问题。

正如评论中已经指出的那样，printf这里的格式说明符 ( %i) 是错误的：
```
printf("%i\n", h_out[i]);
```
打印的数量是double数量，适当的格式说明符是%f.
此代码将不起作用（GPU_N大于 1）：
```
for (int j = 0; j < GPU_N; j++) {
    cudaSetDevice(j);
    cudaMalloc((void**)&d_in, sizeof(double) * data_size / 4);
    cudaMalloc((void**)&d_out, sizeof(double) * data_size / 4);

    printf("2\n");
}
```
d_in并且d_out是个体变量。你不能以这种方式重用它们。当此循环进行第二次（或以后）迭代时，它将覆盖先前分配的指针值。稍后这将导致代码问题，因为对于至少一个内核启动，您将传递指向不驻留在该特定 GPU 上的数据的指针（问题的这个特定方面是错误 77 的近端原因报告。）

一种解决方案是提供指针数组来完成这项工作。
您在循环中发出的某些 CUDA 活动可能是异步的。因此，为确保您的最终打印输出h_out显示预期结果，您应该等待 GPU 上的所有工作完成。实现此目的的一种方法是调用另一组cudaDeviceSynchronize(). （我不想争论是否cudaFree是否异步。我认为这个项目是一个明智的建议，值得注意。如果你觉得你可以跳过这个项目，做你想做的事。出于学习目的，我认为指出这一点很重要。）由于以下评论中指出的原因，此项目对于获得此特定代码的预期结果不是必需的/强制性的。这个答案并不是关于异步工作发布的完整论文；为此，我建议进一步研究cuda标签上的任何相关问题，和/或研究相关的 CUDA 示例代码。

这是一个修改后的代码，解决了上述问题（我缩短了最终的打印输出循环）：

$ cat t1477.cu
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

#define BLOCKS 32
#define THREADS 16

__global__ void kernel(double *d_in, double *d_out) {
    int index = threadIdx.x + blockDim.x * blockIdx.x;
    d_out[index] = d_in[index] + 5;
}

int main() {
    const int GPU_N = 2;
    const int data_size = 2048;
    const int cycles = 2;

    double *h_in, *h_out, *d_in[GPU_N], *d_out[GPU_N];

    h_in = (double*)malloc(sizeof(double) * data_size);
    h_out = (double*)malloc(sizeof(double) * data_size);

    for (int i = 0; i < data_size; i++) {
        h_in[i] = 21;
    }

    cudaError_t error;

    printf("1\n");
    for (int i = 0; i < cycles; i++) {
        //cuMalloc
        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            cudaMalloc((void**)(&(d_in[j])), sizeof(double) * data_size / 4);
            cudaMalloc((void**)(&(d_out[j])), sizeof(double) * data_size / 4);

            printf("2\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            cudaMemcpyAsync(d_in[j], h_in, sizeof(double) * data_size / 4, cudaMemcpyHostToDevice);
            printf("3\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            kernel<<< BLOCKS, THREADS, 0, 0 >>>(d_in[j], d_out[j]);
            error = cudaGetLastError();
            printf("4\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            error = cudaMemcpyAsync(h_out, d_out[j], sizeof(double) * data_size / 4, cudaMemcpyDeviceToHost);
            printf("D2H %i\n", error);
            printf("5\n");
        }

        for (int j = 0; j < GPU_N; j++) {
            cudaSetDevice(j);
            cudaFree(d_in[j]);
            cudaFree(d_out[j]);
            printf("6\n");
        }
    }
    for (int i = 0; i < GPU_N; i++){
        cudaSetDevice(i);
        cudaDeviceSynchronize();}
    for (int i = 0; i < 10; i++) {
        printf("%f\n", h_out[i]);
    }


}
$ nvcc -o t1477 t1477.cu
$ cuda-memcheck ./t1477
========= CUDA-MEMCHECK
1
2
2
3
3
4
4
D2H 0
5
D2H 0
5
6
6
2
2
3
3
4
4
D2H 0
5
D2H 0
5
6
6
26.000000
26.000000
26.000000
26.000000
26.000000
26.000000
26.000000
26.000000
26.000000
26.000000
========= ERROR SUMMARY: 0 errors
$

cuda - 从设备复制回主机时如何修复 cudaError 77

问题描述

解决方案

推荐阅读