class - Cuda warp非法地址
问题描述
我在使用 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)>>> ()
你们有谁知道我犯了错误吗?
解决方案
您发布的代码中有很多损坏,但错误的核心来源是您试图访问内核内的主机指针(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;
}
请注意,您必须基本上在主机内存中构建类的设备副本,然后将其复制到设备以使其正常工作(这是指针或结构数组和包含指针的类的非常常见的设计模式,尽管它是由于复杂性和性能原因,几乎从不推荐)。
推荐阅读
- azure - Azure 函数 - 将所有请求发送到同一个函数
- vb.net - 流没有结束
- ios - 需要在 iMessage Extension 中的 insertText 方法上调用 didStartSending
- python - 为python安装了一个模块,但我不能调用它,它似乎被锁定了
- java - 使用 Java 谓词和 Lambda
- unity3d - 如何销毁 ARCORE 场景中的对象并在同一场景中部署另一个对象?
- maven - Sonarqube 和构建工具之间的关系
- git - Bitbucket - SourceTree 正在发挥作用
- outlook-web-addins - 如何知道 Outlook 插件正在使用 jquery 或 javascript 在浏览器或 Outlook App 中运行?
- python - Python - 将 PyInt 转换为 C int