使用 CUB::DeviceScan

Question

我正在尝试在 CUDA 中进行独家减和。我正在使用 CUB 库并决定尝试 CUB::DeviceReduce。但是，我的结果是 NaN，我不知道为什么。

代码是：

#include <cub/cub.cuh>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
using std::cout;
using std::endl;
#define DSIZE 512

void dev_cumsum( const float *dev_inData, float *dev_outData ) {
    int n = 512;
    void* dev_temp_storage = NULL;
    size_t temp_storage_bytes = 0;
    cub::DeviceScan::ExclusiveSum(dev_temp_storage,temp_storage_bytes,const_cast<float*>(dev_inData),dev_outData,n);
    cudaMalloc(&dev_temp_storage,temp_storage_bytes);
    cub::DeviceScan::ExclusiveSum(dev_temp_storage,temp_storage_bytes,const_cast<float*>(dev_inData),dev_outData,n);
}

int main(){
    float h_data[512];
    float* d_data;
    float* d_result;
    float h_result[512];
    cudaMalloc(&d_data, DSIZE*sizeof(float));
    cudaMalloc(&d_result, DSIZE*sizeof(float));
    h_data[0] = rand()%10;
    h_result[0] = 0;
    for (int i=1; i<DSIZE; i++) {
        h_data[i] = rand()%10;
        h_result[i] = h_data[i-1]+h_result[i-1];
    }
    cudaMemcpy(d_data, h_data, DSIZE*sizeof(float), cudaMemcpyHostToDevice);
    dev_cumsum(d_data, d_result);
    printf("CPU result = %f\n", h_result[511]);
    cudaMemcpy(h_result, d_result, DSIZE*sizeof(float), cudaMemcpyDeviceToHost);
    printf("GPU result = %f\n", h_result[511]);
    for( int i = 0; i < DSIZE; i++ ) {cout << h_result[i] << " ";}
    cout << endl;
    return 0;
}

此代码为设备结果的最后 8 个元素提供 NaN。

此代码在 Linux Mint15 的 GTX650 Ti Boost 上运行。我正在使用 NSight，控制台输出编译命令是：

Invoking: NVCC Compiler
/usr/local/cuda-5.5/bin/nvcc -G -g -O0 -gencode arch=compute_30,code=sm_30 -odir "" -M -o "main.d" "../main.cu"
/usr/local/cuda-5.5/bin/nvcc --device-c -G -O0 -g -gencode arch=compute_30,code=compute_30 -gencode arch=compute_30,code=sm_30  -x cu -o  "main.o" "../main.cu"

Cuda 版本是 5.5 CUB 版本 1.0.2

这是在另一台装有 Cuda 6、OSX10.9.2、CUB 1.2.3 并运行 GT750M 的计算机上测试的，并重现了最后 8 个数字为 NaN 的错误

编辑：代码在 int 和 double 下正常工作，但不是 float。

编辑：感谢 Robert Crovella，这个问题最初是针对 DeviceReduce 提出的。该代码有效，它抛出了 NaN，因为早期使用 DeviceScan 的代码将 NaN 作为输入提供给它。问题已修改以适应

Answer 1

当我运行代码时，我发现不是最后 8 个值设置为 NaN，而是事实上，自最后一个 72 的整数倍数以来的所有值都设置为 NaN。在您的示例中，有 512 个值：这意味着前 504 个（7 * 72）是正确的，接下来的 8 个值是 NaN。

这种行为似乎一直持续到 568 (8 * 72) 个值，此后它似乎可以正常工作。

我用来测试的代码在这里：http://pastebin.com/kXVvuKAN

我使用以下命令编译了代码：

nvcc --relocatable-device-code=true -gencode arch=compute_30,code=compute_30 -G -o main main.cu

注意：如果我不使用 -G 参数，结果会更加随机。但是，使用 -G 命令，它给出了上面提到的清晰模式。

Answer 2

编辑： cub 1.3.0 最近发布了，我相信它包含针对此问题的修复程序。

---

我会对您的代码进行一些更改，我认为这些更改是错误的，但我不知道它们是否会影响您所看到的内容。在下面的代码部分中，您使用的是h_result[0]，但没有对其进行初始化，因此请添加我用注释标记的行：

h_data[0] = rand()%10;
h_result[0] = 0;    // ADD THIS LINE
for (int i=1; i<DSIZE; i++) {
    h_data[i] = rand()%10;
    h_result[i] = h_data[i-1]+h_result[i-1];
}

（显然这不应该影响您的 GPU 结果。）另外，您最终的cudaMemcpy 操作不太正确：

cudaMemcpy(&h_result, d_result, DSIZE*sizeof(float), cudaMemcpyDeviceToHost);
           ^
           delete this ampersand

由于h_result 已经是您的公式中的一个指针，我们不需要将它的地址传递给cudaMemcpy。

您可以尝试进行这些更改，看看您会得到什么样的结果？

我一直在努力解决这个问题。如果您仍然可以重现该错误，我将不胜感激：

1. 重启你的机器，然后再试一次
2. 回复您正在运行的实际更新代码、您正在使用的编译命令、CUDA 版本、CUB 版本和您正在运行的 GPU，以及系统操作系统。 （使用此信息编辑您的原始问题）