为什么从 __global__ 函数中引用外部内存会搞砸一切？答案

【问题标题】：Why referencing to outer memory from within __global__ function messes everything up?为什么从 __global__ 函数中引用外部内存会搞砸一切？
【发布时间】：2012-06-04 08:18:17
【问题描述】：

我正在用 CUDA 编写一些代码（确切地说是霍夫曼算法，但这与案例完全无关）。我有一个文件 Paralellel.cu 有两个函数：一个 (WriteDictionary) 是一个普通函数，第二个 (wrtDict) 是一个在 CUDA GPU 中运行的特殊 CUDA _global_ 函数。以下是这些函数的主体：

//I know body of this function looks kinda not-related 
//   to program main topic, but it's just for tests.
__global__ void wrtDict(Node** nodes, unsigned char* str)
{
    int i = threadIdx.x;

    Node* n = nodes[i];
    char c = n->character;

    str[6 * i] = 1;//c;                         !!!
    str[6 * i + 1] = 2;

    str[6 * i + 2] = 0;
    str[6 * i + 3] = 0;
    str[6 * i + 4] = 0;
    str[6 * i + 5] = 0;
}

我知道前两行似乎毫无意义，因为我在这里没有使用 Node 类的这个对象 n，而只是让它们暂时存在。还有一条超级机密的评论，标有“！！！”。这是WriteDictionary：

void WriteDictionary(NodeList* nodeList, unsigned char* str)
{
    Node** nodes = nodeList->elements;   
    int N = nodeList->getCount();

    Node** cudaNodes;
    unsigned char* cudaStr;

    cudaMalloc((void**)&cudaStr, 6 * N * sizeof(unsigned char));
    cudaMalloc((void**)&cudaNodes, N * sizeof(Node*));

    cudaMemcpy(cudaStr, str, 6 * N * sizeof(char), cudaMemcpyHostToDevice); 
    cudaMemcpy(cudaNodes, nodes, N * sizeof(Node*), cudaMemcpyHostToDevice);

    dim3 block(1);
    dim3 thread(N);

    std::cout << N << "\n";

    wrtDict<<<block,thread>>>(cudaNodes, cudaStr);

    cudaMemcpy(str, cudaStr, 6 * N * sizeof(unsigned char), cudaMemcpyDeviceToHost);


    cudaFree(cudaNodes);
    cudaFree(cudaStr);
}

正如所见，WriteDictionary 函数是 CUDA 和程序其余部分之间的一种代理。我在 Node 指向的普通内存中的某处有一堆我的类 Node 的对象 * 数组 elements 保存在我的对象中节点列表。现在了解 Node 就足够了，它有一个公共字段 char character。 char * str 现在将填充一些测试数据。它包含 6 * N 个为字符分配的内存，其中 N = elements 数组中所有元素的计数。所以我在 CUDA 中为 6 * N chars 和 N Node 指针分配了一个内存空间。然后我将 Node 指针复制到那里，它们仍然指向普通内存。我正在运行该功能。在函数 wrtDict 中，我将字符提取到 char c 变量中，这次不尝试将其放入输出数组 str。

所以，当我编写输出数组 str 的内容时（在 WriteDictionary 函数之外），我得到了完全正确的答案，即：

1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0   1  2  0  0  0  0 
1  2  0  0  0  0

是的，这里我们有 39 个正确的 chars 六进制（以十六进制显示）。但是当我们在 wrtDict 函数中稍微改变我们的超级秘密评论时，像这样：

__global__ void wrtDict(Node** nodes, unsigned char* str)
{
    int i = threadIdx.x;

    Node* n = nodes[i];
    char c = n->character;

    str[6 * i] = c;//1;                         !!!
    str[6 * i + 1] = 2;

    str[6 * i + 2] = 0;
    str[6 * i + 3] = 0;
    str[6 * i + 4] = 0;
    str[6 * i + 5] = 0;
}

我们会看到奇怪的事情。我现在期望每六个中的第一个 char 是数组所指向的 Node 中的一个字符 - 每个都不同。或者，即使它失败了，我预计每六个字符中只有第一个 char 会被弄乱，但其余的保持不变：? 2 0 0 0 0。但不是！当我这样做时，一切都变得一团糟，现在输出数组 str 的内容如下所示：

70 21 67 b7 70 21  67 b7  0  0  0  0 
 0  0  0  0 18 d7  85  8 b8 d7 85  8 
78 d7 85  8 38 d9  85  8 d8 d7 85  8 
f8 d5 85  8 58 d6  85  8 d8 d5 85  8 
78 d6 85  8 b8 d6  85  8 98 d7 85  8 
98 d6 85  8 38 d6  85  8 d8 d6 85  8 
38 d5 85  8 18 d6  85  8 f8 d6 85  8 
58 d9 85  8 f8 d7  85  8 78 d9 85  8 
98 d9 85  8 d8 d4  85  8 b8 d8 85  8 
38 d8 85  8 38 d7  85  8 78 d8 85  8 
f8 d8 85  8 d8 d8  85  8 18 d5 85  8 
61 20 75 6c 74 72  69 63 65 73 20 6d 
6f 6c 65 73 74 69  65 20 73 69 74 20 
61 6d 65 74 20 69  64 20 73 61 70 69 
65 6e 2e 20 4d 61  75 72 69 73 20 73 
61 70 69 65 6e 20  65 73 74 2c 20 64 
69 67 6e 69 73 73  69 6d 20 61 63 20 
70 6f 72 74 61 20  75 74 2c 20 76 75 
6c 70 75 74 61 74  65 20 61 63 20 61 
6e 74 65 2e 20 46

我现在问 - 为什么？是因为我试图从 CUDA GPU 中获取普通内存吗？我收到一个警告，可能就是关于这种情况，说：

Cannot tell what pointer points to, assuming global memory space

我已经用谷歌搜索了这个，只找到了这个，CUDA 它正在到达一个普通的内存，导致无法找到到达哪里，99.99% 的这个警告应该被忽略。所以我忽略了它，认为它会没事的，但事实并非如此——我的情况在这 0.01% 之内吗？

我该如何解决这个问题？我知道我可以将 Nodes（而不是指向它们的指针）复制到 CUDA 中，但我认为复制它们会花费我更多的时间，而不是节省并行化内部对它们所做的事情。我还可以从每个 Node 中提取 character，将它们全部放入一个数组中，然后将其复制到 CUDA，但是 - 与上一条语句中的问题相同。

我完全不知道该做什么，更糟糕的是，我大学的 CUDA 项目的截止日期是今天，apx。晚上 17 点（我只是没有足够的时间来早点，该死的......）。

PS。如果有帮助：我正在使用非常简单的（没有任何开关）命令进行编译：

nvcc -o huff ArchiveManager.cpp IOManager.cpp Node.cpp NodeList.cpp Program.cpp Paraleller.cu

【问题讨论】：

对于您发布的所有（大部分不相关的）信息，缺少很多关键信息。如果您不愿意解释Node 的结构或“超级秘密评论”是什么，我非常怀疑是否有人可以帮助您。我投票结束这个问题，它写得非常糟糕（尽管它很长），而且它的解决方案非常不太可能对其他人有益 - 它是too localised。

标签： c++ memory cuda gpu nvidia

【解决方案1】：

这是一个可怕的问题，请参阅 talonmies 的评论。

检查每个 CUDA API 调用的错误值。内核启动后，您将在cudaMemcpy 上收到启动失败消息
运行cuda-memcheck 帮助调试错误（这基本上是分段错误）
意识到您正在从 GPU 取消引用（未映射的）指向主机内存的指针，您需要复制节点，而不仅仅是指向节点的指针

【讨论】：

扩展一点......你不能只使用指向 GPU 的指针的 memcpy 结构。为什么？因为只有结构会在 GPU 内存中。指针仍将指向 GPU 无法访问的 CPU 内存。您要么需要单独上传结构的所有小部分，更改指针，要么使用不涉及或涉及更少指针的不同数据结构。

【解决方案2】：

您也可以在 cuda-gdb 中运行您的程序。 cuda-gdb 会告诉你你遇到了什么错误。另外，在cuda-gdb的开头，做一个“set cuda memcheck on”，它会在cuda-gdb里面打开memcheck。

在最新的 cuda-gdb 版本（截至今天为 5.0）中，如果您不检查 API 调用的返回码并且这些 API 调用失败，您也会看到警告。

【讨论】：