Использование 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;
}

Этот код дает мне NaN за последние 8 элементов результата устройства.

Этот код работает на GTX650 Ti Boost в Linux Mint15. Я использую 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, но не с плавающей точкой.

редактирование: благодаря Роберту Кровелле этот вопрос изначально задавался в отношении 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,

Можете ли вы попробовать внести эти изменения и посмотреть, какие результаты вы получите?

Я немного боролся с этим. Если вы все еще можете воспроизвести ошибку, я буду признателен, если вы сможете:

1. перезагрузите компьютер и попробуйте снова
2. ответьте актуальным обновленным кодом, который вы используете, командой компиляции, версией CUDA, версией CUB и графическим процессором, на котором вы работаете, а также операционной системой системы. (отредактируйте исходный вопрос с помощью этой информации)

Answer 2

Когда я запустил код, я обнаружил, что это не последние 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 он дал четкую схему, упомянутую выше.