Utilisation de CUB :: DeviceScan

Question

J'essaye de faire une réduction de somme exclusive dans CUDA. J'utilise la bibliothèque CUB et j'ai décidé d'essayer le CUB :: DeviceReduce. Cependant, mon résultat est NaN, et je n'arrive pas à comprendre pourquoi.

code

est:

#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; 
} 

Ce code me donne NaN pour les 8 derniers éléments du résultat de l'appareil.

Ce code fonctionne sur une GTX650 Ti Boost sous Linux Mint15. J'utilise nSight et la sortie de la console compiler commande est:

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" 

la version 5.5 est Cuda version CUB 1.0.2

Cela a été testé sur un autre ordinateur avec Cuda 6, OSX10.9.2, CUB 1.2 .3 et l'exécution d'un GT750M, et reproduit l'erreur des 8 derniers numéros étant NaN

edit: Le code fonctionne correctement avec int et double, mais pas flottant.

edit: Merci à Robert Crovella, cette question a été posée à l'origine en ce qui concerne DeviceReduce. Ce code fonctionnait, il lançait NaN parce que le code précédent utilisant DeviceScan lui fournissait NaN comme entrée. La question est révisée pour convenir à

Answer 1

EDIT:cub 1.3.0 a récemment été publié, et je crois qu'il comprend un correctif pour ce problème.

---

Il y a quelques changements que je voudrais apporter à votre code, que je considère comme des erreurs, mais je ne sais pas si elles affectent ce que vous voyez. Dans la section de code suivant, vous utilisez h_result[0] sans l'initialiser, il faut donc ajouter la ligne que j'ai marqué avec un commentaire: (. Il est clair que l'on ne devrait pas être Influencer le résultat de votre GPU)

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]; 
} 

En outre, votre dernière cudaMemcpy opération n'est pas tout à fait raison:

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

Depuis h_result est déjà un pointeur dans votre formulation, on n'a pas besoin de passer l'adresse de celui-ci à cudaMemcpy.

Pouvez-vous essayer de faire ces changements et voir quel genre de résultats vous obtenez?

J'ai eu du mal avec ça. Si vous pouvez toujours reproduire l'erreur, je vous en serais reconnaissant si vous pouvez:

1. redémarrer votre machine et essayez à nouveau
2. répondre en arrière avec le code mis à jour réel que vous utilisez, la commande de compilation que vous utilisez, la version CUDA, la version CUB et le GPU que vous utilisez, ainsi que le système d'exploitation. (Éditez votre question originale avec cette information)

Answer 2

Quand j'ai couru le code, j'ai trouvé que ce n'est pas les 8 dernières valeurs qui sont mises à NaN mais c'est, en fait, toutes les valeurs depuis le dernier nombre entier de 72 qui sont définis sur NaN. Dans votre exemple, il y a 512 valeurs: cela signifie que les premiers 504 (7 * 72) étaient corrects et que les 8 valeurs suivantes étaient NaN.

Ce comportement semble se poursuivre jusqu'à 568 (8 * 72) valeurs et par la suite il semble fonctionner correctement.

Le code que je l'habitude de tester c'est ici: http://pastebin.com/kXVvuKAN

Je compilé le code avec la commande suivante:

nvcc --relocatable-device code = true arc -gencode = compute_30, le code = compute_30 -G -o main.cu principal

REMARQUE: Si je n'ai pas utilisé le paramètre -G, les résultats sont plus aléatoires. Cependant, avec la commande -G, il a donné le motif clair mentionné ci-dessus.