-1

我有一个奇怪的问题,所以我想我会问一下,看看是否有比我更有经验的人能找到解决方案。

我正在用 CUDA C/C++ 编写一个程序,我有一些常量整数来指定各种东西,比如计算边界的坐标等。目前我只有这些东西在全局设备内存中。它们被每个内核调用中的每个线程访问,所以我认为如果它们在全局内存中,那么它们永远不会被缓存或广播(对吗?)。所以这些小整数占用了很多(相对)开销,并且有很多“读取冗余”。

所以我在标题中声明:

__constant__ int* number;

我包括那个标题,并且,当我做记忆的东西时,我会:

cutilSafeCall( cudaMemcpyToSymbol(number, &(some_host_int), sizeof(int) );

然后我传入number我所有的内核:

__global__ void magical_kernel(int* number, ...){

   //and I access 'number' like this
   int data_thingy = big_array[ *number ];

}

我的代码崩溃了。有了全局内存中的数字,就可以了。我已经确定它在访问内核中的数字时会崩溃。这意味着我访问或分配错误。如果它持有错误的值,它也会导致崩溃,因为它是用来索引到数组中的。

最后,我会问几个问题。首先,我做错了什么?作为奖励:有没有比常量内存更好的方法来完成这个任务 - 我不知道number编译时的值,所以一个简单的 #define 不起作用。恒定内存是否会完全加快代码速度,或者它是否一直被缓存和广播?我能否以某种方式将每个线程块的数据放入共享内存中,并通过多个内核调用将其保留在共享内存中?

4

1 回答 1

1

这里有几个问题:

  1. 您已声明number为指针,但从未为其分配一个值,该值是 GPU 内存中的有效地址
  2. 您有一个变量范围冲突:int * number定义的参数变量与定义为编译单元范围magic_kernel 的变量不同。__constant__ int * variable
  3. 调用的第一个参数cudaMemcpyToSymbol几乎可以肯定是不正确的。

如果您不明白为什么前两点中的任何一点都是正确的,那么您需要对 C++ 中的指针和范围进行一些修改。

根据您对现已删除的答案的回复,我怀疑您实际上想要做的是:

__constant__ int number;

__global__ void magical_kernel(...){

   int data_thingy = big_array[ number ];

}

cudaMemcpyToSymbol("number", &(some_host_int), sizeof(int));

ienumber旨在成为常量内存中的整数,而不是指针,也不是内核参数。


编辑:这是一个例子,它显示了这一点:

#include <cstdio>

__constant__ int number;

__global__ void magical_kernel(int * out)
{
   out[threadIdx.x] = number;
}

int main()
{
    const int value = 314159;
    const size_t sz = size_t(32) * sizeof(int);
    cudaMemcpyToSymbol("number", &value, sizeof(int));

    int * _out, * out;

    out = (int *)malloc(sz);
    cudaMalloc((void **)&_out, sz);

    magical_kernel<<<1,32>>>(_out);

    cudaMemcpy(out, _out, sz, cudaMemcpyDeviceToHost);
    for(int i=0; i<32; i++)
        fprintf(stdout, "%d %d\n", i, out[i]);

    return 0;
}

您应该能够自己运行它并确认它像宣传的那样工作。

于 2012-06-17T06:52:35.557 回答