使用CUB::DeviceScan

Question

我正试着在CUDA中做一个独占的sum缩减。我正在使用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

这是在另一台运行CUDA6，OSX10.9.2，CUB1.2.3并运行GT750M的计算机上测试的，并重现了最后8个数字为NaN的错误

编辑:代码可以正确使用int和double，但不能使用float。

编辑:感谢Robert Crovella，这个问题最初是关于DeviceReduce的。这段代码起作用了，它抛出了NaN，因为早期使用DeviceScan的代码将NaN作为输入提供给它。问题已修改为适合

Answer 1

编辑： 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。

你能试着做这些改变，看看会得到什么样的结果吗？

我在这个问题上有点纠结。如果你仍然可以重现这个错误，我将不胜感激：

CUDA重新启动您的计算机，并使用您正在运行的实际更新的代码、您正在使用的编译命令、版本、CUB版本、您正在运行的

1. 以及系统操作系统重新尝试again
2. respond。(使用此信息编辑您的原始问题)

Answer 2

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

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

我用来测试它的代码是：http://pastebin.com/kXVvuKAN

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

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

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