【问题标题】:CUDA: 2D array indexing giving unexpected resultsCUDA:二维数组索引给出了意想不到的结果
【发布时间】:2013-05-30 03:28:13
【问题描述】:

我开始学习 CUDA,我想编写一个简单的程序,将一些数据复制到 GPU,对其进行修改,然后将其传回。我已经用谷歌搜索并试图找到我的错误。我很确定问题出在我的内核中,但我不完全确定出了什么问题。

这是我的内核:

__global__ void doStuff(float* data, float* result)
{
    if (threadIdx.x < 9) // take the first 9 threads
    {
        int index = threadIdx.x;
        result[index] = (float) index;
    }
}

这是我main的相关部分:

#include <stdlib.h>
#include <stdio.h>

int main(void)
{
    /*
        Setup
    */
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};

    float* data_array;
    float* result_array;

    size_t data_array_pitch, result_array_pitch;
    int width_in_bytes = 3 * sizeof(float);
    int height = 3;

    /*
        Initialize GPU arrays
    */
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);

    /*
        Copy data to GPU
    */
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);

    dim3 threads_per_block(16, 16);
    dim3 num_blocks(1,1);

    /*
        Do stuff
    */
    doStuff<<<num_blocks, threads_per_blocks>>>(data_array, result_array);

    /*
        Get the results
    */
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);

    for (int i = 1; i <= 9; ++i)
    {
        printf("%f ", simple[i-1]);
        if(!(i%3))
            printf("\n");
    }

    return 0;
}

当我运行它时,第一行得到0.000000 1.000000 2.00000,其他两行得到垃圾。

【问题讨论】:

  • 如果您对所有 cuda API 调用和内核调用执行 cuda error checking,您会收到任何错误吗?当您使用 cuda-memcheck 运行代码时会发生什么?
  • 一切都返回了cudaSuccess
  • 访问数组中的元素时是否需要考虑音高?我现在正在查看 NVIDIA 指南的第 30 页。

标签: c++ cuda


【解决方案1】:

如果您刚开始学习 cuda,我不确定我是否会专注于 2D 数组。

如果您在问题中手动键入代码也很好奇,因为您定义了一个 threads_per_block 变量,但随后您在内核调用中使用了 threads_per_blocks

不管怎样,你的代码有几个问题:

  1. 使用 2D 数组时,几乎总是需要传递间距 参数(以某种方式)到内核。 cudaMallocPitch 在每行的末尾分配带有额外填充的数组,这样 下一行从一个很好对齐的边界开始。这通常会 导致分配粒度为 128 或 256 字节。所以你的第一个 行有 3 个有效数据实体,后面有足够的空白空间来填充 最多,比如说 256 个字节(等于你的音高变量)。所以我们必须改变内核调用和内核本身来解决这个问题。
  2. 您的内核本质上是一维内核(例如,它不理解或使用threadIdx.y)。因此,启动 2D 网格没有意义。虽然在这种情况下它不会造成任何伤害,但它会产生冗余,这在其他代码中可能会造成混淆和麻烦。

这是一个更新的代码,显示了一些更改,这些更改将为您提供预期的结果,基于上述 cmets:

#include <stdio.h>


__global__ void doStuff(float* data, float* result, size_t dpitch, size_t rpitch, int width)
{
    if (threadIdx.x < 9) // take the first 9 threads
    {
        int index = threadIdx.x;
        result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
    }
}

int main(void)
{
    /*
        Setup
    */
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};

    float* data_array;
    float* result_array;

    size_t data_array_pitch, result_array_pitch;
    int height = 3;
    int width = 3;
    int width_in_bytes = width * sizeof(float);

    /*
        Initialize GPU arrays
    */
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);

    /*
        Copy data to GPU
    */
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);

    dim3 threads_per_block(16);
    dim3 num_blocks(1,1);

    /*
        Do stuff
    */
    doStuff<<<num_blocks, threads_per_block>>>(data_array, result_array, data_array_pitch, result_array_pitch, width);

    /*
        Get the results
    */
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);

    for (int i = 1; i <= 9; ++i)
    {
        printf("%f ", simple[i-1]);
        if(!(i%3))
            printf("\n");
    }
    return 0;
}

您可能还会发现this question 有趣的阅读。

编辑:回答 cmets 中的问题:

result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
              1               2                      3

要计算正确的元素索引到倾斜数组中,我们必须:

  1. 根据线程索引计算(虚拟)行索引。为此,我们将线程索引除以每个(非间距)行的宽度(以元素为单位,而不是字节)。
  2. 将行索引乘以每个间距行的宽度。每个 pitched 行的宽度由 pitched 参数给出,以字节为单位。为了将这个音调 byte 参数转换为音调 element 参数,我们除以每个元素的大小。然后通过将数量乘以在步骤 1 中计算的行索引,我们现在已经索引到了正确的行。
  3. 通过将线程索引的余数(模除)除以宽度(以元素为单位)计算线程索引的(虚拟)列索引。获得列索引(以元素为单位)后,我们将其添加到第 2 步计算的正确行开始索引中,以识别该线程将负责的元素。

对于一个相对简单的操作来说,以上是相当大的努力,这就是为什么我建议首先关注基本 cuda 概念而不是倾斜数组的一个例子。例如,在处理倾斜阵列之前,我将计算如何处理 1 维和 2D 线程块,以及 1 和 2D 网格。在某些情况下,间距数组是访问 2D 数组(或 3D 数组)的有用的性能增强器,但在处理 CUDA 中的多维数组时,它们绝不是必需的。

【讨论】:

  • 嗯...是的,我确实是手动输入的。对于那个很抱歉。您能更详细地解释一下result[...] 行吗?
  • 感谢您的解释!
【解决方案2】:

其实也可以通过换行来完成

int width_in_bytes = 3 * sizeof(float);

作者:

int width_in_bytes = sizeof(float)*9;

因为这是告诉 cudaMemcpy2D 从 src 复制多少字节到 dst 的参数,在第一个代码中您要求复制 3 个浮点数,但是您要复制的数组的长度为 9,所以您需要的宽度是9 个浮点数的大小。

尽管此解决方案有效,但您的代码仍然存在一些低效率;例如,如果您真的希望块的前 9 个线程执行某些操作,则在“if”中您应该使用 and(&&) 添加以下条件

threadIdx.y==0

【讨论】:

    猜你喜欢
    • 1970-01-01
    • 2019-05-19
    • 2017-01-05
    • 2021-10-14
    • 1970-01-01
    • 1970-01-01
    • 2014-06-22
    • 2021-12-17
    • 2021-04-22
    相关资源
    最近更新 更多