Cuda warp非法地址答案

【问题标题】：Cuda warp illegal addressCuda warp非法地址
【发布时间】：2019-01-14 09:08:50
【问题描述】：

我在使用 CUDA 和将类传递给内核时遇到了一些问题。我有一些函数可以为 GPU 上的类分配内存，传递它，并且工作正常。但是，还有另一个是行不通的。我注意到只有在我使用数组时才会发生这种情况。这是一个例子。

文件1.hh

#ifndef PROVA1_HH
#define PROVA1_HH

#include <cstdio>

class cls {
public:
    int *x, y;
    cls();  
    void kernel();
};

#endif

文件1.cu

#include "Prova1.hh"

__global__ void kernel1(cls* c){
    printf("%d\n", c->y);
    c->y=2;
    printf("%d\n", c->y);
    c->x[0]=0; c->x[1]=1;
    printf("%d %d\n", c->x[0], c->x[1]);
}

void cls::kernel(){
    cls* dev_c; cudaMalloc(&dev_c, sizeof(cls));
    cudaMemcpy(dev_c, this, sizeof(cls), cudaMemcpyHostToDevice);
    printf("(%d, %d)\n", x[0], x[1]);
    kernel1<<<1, 1>>> (dev_c);
    cudaDeviceSynchronize();
    cudaMemcpy(this, dev_c, sizeof(cls), cudaMemcpyDeviceToHost);
    printf("(%d, %d)\n", x[0], x[1]);
}

cls::cls(){
    y=3;
    x=(int*) malloc(sizeof(int)*2);
    x[0]=1; x[1]=2;
}

文件.cu

#include<cstdio>
#include "Prova1.hh"

int main(){
    cls c=cls();
    c.kernel();

    return 0;
}

我正在编译：

nvcc -std=c++11 -arch=sm_35 -rdc=true -c -o File1.o File1.cu
nvcc -std=c++11 -arch=sm_35 -rdc=true -g -G -o File.out File1.o File.cu

当我简单运行它时，输出将是：

(1, 2)
3
2
(1, 2)

当我调试它时，我得到：

Starting program: 
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/aarch64-linux-gnu/libthread_db.so.1".
[New Thread 0x7fb10eb1e0 (LWP 806)]
(1, 2)

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x84fa10

Thread 1 "File.out" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
0x000000000084fad0 in kernel1(ciao*)<<<(1,1,1),(1,1,1)>>> ()

你们有谁知道我犯了错误吗？

【问题讨论】：

您如何想象dev_c->x 会为它分配设备内存？
我在 cls::kernel() 中分配了内存。这还不够吗？
不，你没有为dev_c->x分配内存。在您的主机构造函数中，有一个对 cls.x 的 malloc 调用。 dev_c 没有等价物。
啊，现在我明白了。所以我应该为 cls ad dev_c 分配内存，而不是为 dev_c.x... 对吗？
不，不是。 dev_c 分配到设备上后，不能从主机修改。

标签： class memory cuda

【解决方案1】：

您发布的代码中有很多损坏，但错误的核心来源是您正试图访问内核内的主机指针（设备上没有分配内存给 x 和值也不复制）。除非您使用托管内存，否则这显然是行不通的。

你可以把你的例子改成这样：

#include <cstdio>

class cls {
public:
    int *x, y;

    __host__ __device__
    cls(int *x_, int y_) : x(x_), y(y_) {};  

    void kernel();
};

__global__ void kernel1(cls* c){
    printf("%d\n", c->y);
    c->y=2;
    printf("%d\n", c->y);
    c->x[0]=0; c->x[1]=1;
    printf("%d %d\n", c->x[0], c->x[1]);
}

void cls::kernel(){

    int* dev_x; cudaMalloc(&dev_x, sizeof(int)*2);
    cudaMemcpy(dev_x, x, sizeof(int)*2, cudaMemcpyHostToDevice);

    cls h_dev_c(dev_x, y);
    cls* dev_c; cudaMalloc(&dev_c, sizeof(cls));
    cudaMemcpy(dev_c, &h_dev_c, sizeof(cls), cudaMemcpyHostToDevice);

    printf("(%d)\n", y);
    printf("(%d, %d)\n", x[0], x[1]);
    kernel1<<<1, 1>>> (dev_c);
    cudaDeviceSynchronize();

    cudaMemcpy(&y, &(dev_c->y), sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(x, dev_x, sizeof(int)*2, cudaMemcpyDeviceToHost);
    printf("(%d)\n", y);
    printf("(%d, %d)\n", x[0], x[1]);
}

int main(){

    int y=3;
    int* x=(int*) malloc(sizeof(int)*2);
    x[0]=1; x[1]=2;

    cls c(x,y);
    c.kernel();

    return 0;
}

请注意，您必须基本上在主机内存中构建类的设备副本，然后将其复制到设备以使其正常工作（这是指针或结构数组和包含指针的类的非常常见的设计模式，尽管出于复杂性和性能方面的原因，几乎从不推荐它）。

【讨论】：