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

Question

我在将指向结构的指针传递给设备函数时遇到问题。我想在本地内存中创建一个结构（我知道它很慢，这只是一个例子），并通过指针将其传递给另一个函数。问题是，当我在打开memcheck的情况下调试它时，我会得到错误：

Program received signal CUDA_EXCEPTION_1, Lane Illegal Address. Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 7, warp 0, lane 0 0x0000000000977608 in foo (st=0x3fffc38) at test.cu:15 15 st->m_tx = 99;

如果我在没有打开memcheck的情况下调试它，它会正常工作，并给出预期的结果。我的操作系统是RedHat 6.3 64位内核2.6.32-220。我使用GTX680，CUDA5.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;
}

有什么建议吗？谢谢你的帮助。

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;
        ...............................

我们在汇编程序输出中得到实际的代码。你可能会在调试器中有更多的运气。

Answer 2

我来自CUDA开发工具团队。当编译用于设备端调试（即-G）时，原始代码将不会被优化。该问题看起来像是内存检查错误。谢谢你找到这个。我们将对此进行调查。