Nvidia CUDA - 通过指针传递结构体

Question

我有一个问题，需要将指向结构体的指针传递给设备函数。 我想在本地内存中创建一个结构体（我知道这会很慢，但这只是一个例子），并通过指针将其传递给另一个函数。问题是，当我使用memcheck进行调试时，出现错误：

程序接收到CUDA_EXCEPTION_1信号，地址非法。 切换焦点到CUDA kernel 0，网格1，块(0,0,0)，线程(0,0,0)，设备0，sm 7，warp 0，lane 0 0x0000000000977608在foo(st=0x3fffc38)测试.cu:15 15 st->m_tx = 99;

如果我不使用memcheck进行调试，它可以正常工作并返回预期结果。 我的操作系统是RedHat 6.3 64位，内核为2.6.32-220。 我使用GTX680，CUDA 5.0，并使用sm=30编译程序。

下面是我用于测试此功能的代码：

typedef struct __align__(8) {
    int m_x0;
    int m_tx;
} myStruct;

__device__ void foo(myStruct *st) {
    st->m_tx = 99;
    st->m_x0 = 123;
}
__global__ void myKernel(){
    myStruct m_struct ;
    m_struct.m_tx = 45;
    m_struct.m_x0 = 90;
    foo(&m_struct);
}
int main(void) {
    myKernel  <<<1,1 >>>();
    cudaThreadSynchronize();
    return 0;
}

有什么建议吗？非常感谢任何帮助。

- unnamed

---

您的设备代码不会在死代码删除后保留，因为没有任何内容对全局内存写入做出贡献。如果您在此示例中查看除空内核以外的任何内容，我会感到惊讶。 - talonmies

我假设你所说的“CUDA 5.0”是指CUDA 5.0的候选版本（而不是早期的预览）？ - njuffa

Answer 1

你的示例代码被编译器完全优化掉了，因为代码对全局内存写入没有任何贡献。可以通过将内核编译成cubin文件并使用cuobjdump反汇编结果来轻松证明这一点：

$ nvcc -arch=sm_20 -Xptxas="-v" -cubin struct.cu 
ptxas info    : Compiling entry function '_Z8myKernelv' for 'sm_20'
ptxas info    : Function properties for _Z8myKernelv
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 2 registers, 32 bytes cmem[0]

$ cuobjdump -sass struct_dumb.cubin 

    code for sm_20
        Function : _Z8myKernelv
    /*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
    /*0008*/     /*0x00001de780000000*/     EXIT;
        .............................

例如，内核完全为空。调试器无法调试您想要调查的代码，因为它在编译器/汇编器发出的内容中不存在。如果我们对您的代码进行一些修改：

typedef struct __align__(8) {
    int m_x0;
    int m_tx;
} myStruct;
__device__ __noinline__ void foo(myStruct *st) {
    st->m_tx = 99;
    st->m_x0 = 123;
}
__global__ void myKernel(int dowrite, int *output){
    myStruct m_struct ;
    m_struct.m_tx = 45;
    m_struct.m_x0 = 90;
    if (dowrite) {
        foo(&m_struct);
        output[threadIdx.x] = m_struct.m_tx + m_struct.m_x0;
    }
}
int main(void) {
    int * output;
    cudaMalloc((void **)(&output), sizeof(int));
    myKernel  <<<1,1 >>>(1, output);
    cudaThreadSynchronize();
    return 0;
}

当您重复相同的编译和反汇编步骤时，情况会有所不同：

$ nvcc -arch=sm_20 -Xptxas="-v" -cubin struct_dumb.cu 
ptxas info    : Compiling entry function '_Z8myKerneliPi' for 'sm_20'
ptxas info    : Function properties for _Z8myKerneliPi
    8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z3fooP8myStruct
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 5 registers, 40 bytes cmem[0]
$ /usr/local/cuda/bin/cuobjdump -sass struct_dumb.cubin 

    code for sm_20
        Function : _Z8myKerneliPi
    /*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
    /*0008*/     /*0x20105d034800c000*/     IADD R1, R1, -0x8;
    /*0010*/     /*0x68009de218000001*/     MOV32I R2, 0x5a;
    /*0018*/     /*0xb400dde218000000*/     MOV32I R3, 0x2d;
    /*0020*/     /*0x83f1dc23190e4000*/     ISETP.EQ.AND P0, pt, RZ, c [0x0] [0x20], pt;
    /*0028*/     /*0x00101c034800c000*/     IADD R0, R1, 0x0;
    /*0030*/     /*0x00109ca5c8000000*/     STL.64 [R1], R2;
    /*0038*/     /*0x000001e780000000*/     @P0 EXIT;
    /*0040*/     /*0x10011c0348004000*/     IADD R4, R0, c [0x0] [0x4];
    /*0048*/     /*0xc001000750000000*/     CAL 0x80;
    /*0050*/     /*0x00009ca5c0000000*/     LDL.64 R2, [R0];
    /*0058*/     /*0x84011c042c000000*/     S2R R4, SR_Tid_X;
    /*0060*/     /*0x90411c4340004000*/     ISCADD R4, R4, c [0x0] [0x24], 0x2;
    /*0068*/     /*0x0c201c0348000000*/     IADD R0, R2, R3;
    /*0070*/     /*0x00401c8590000000*/     ST [R4], R0;
    /*0078*/     /*0x00001de780000000*/     EXIT;
    /*0080*/     /*0x8c00dde218000001*/     MOV32I R3, 0x63;
    /*0088*/     /*0xec009de218000001*/     MOV32I R2, 0x7b;
    /*0090*/     /*0x1040dc8590000000*/     ST [R4+0x4], R3;
    /*0098*/     /*0x00409c8590000000*/     ST [R4], R2;
    /*00a0*/     /*0x00001de790000000*/     RET;
        ...............................

我们在汇编输出中获得实际的代码。你可能会在调试器中更容易地找到它。

- talonmies

---

我刚刚测试了你的代码，它给我返回了与memcheck相同的错误。 - unnamed

Answer 2

我来自CUDA开发工具团队。当编译时使用设备端调试（即-G）时，原始代码将不会被优化掉。该问题看起来像是一个memcheck的错误。感谢您发现了这个问题。我们会进一步研究它。

- Vyas

---

问题可能是平台特定的，因为我无法在本地重现它（我使用 -g -G 编译了代码）。 - njuffa