Peut-on utiliser du code C++ dans un programme CUDA C ? Lorsque j’inclus le fichier d’en-tête iostream, le compilateur confond le ou les appels du noyau.



par

CUDA ne supporte pas la STL. Vector ne fonctionne tout simplement pas dans un noyau CUDA. Ce qui s'en rapproche le plus doit être la bibliothèque Thrust. For example,

  1. thrust::device_vector vect; 

but it doesn’t work in device kernel code. Its just an STL-like structure with a GPU backed storage so that you can do the math on it through GPU-computation using simple codes like this:

  1. thrust::transform( 
  2. vect.begin(),  
  3. vect.end(),  
  4. vect.begin(),  
  5. [=] __device__ (float x) {return x + 1;} 
  6. ); 

So its still a CUDA C/C++ program but not an explicit device kernel code. If you meant to run them in device kernel code(not host), then you have to implement one yourself because there are many use cases and every one of them will be highly fluctuating in performance depending on hardware and algorithm.

For example, a parallel-vector in a kernel could be like:

  1. class fake_heap 
  2. { 
  3. public: 
  4. __device__ fake_heap(float * heap) 
  5. { 
  6. h=heap; 
  7. ctr=0; 
  8. }  
  9. int ctr; 
  10. float * h;  
  11. }; 
  12.  
  13. class vector 
  14. { 
  15. public: 
  16. __device__ vector(fake_heap * heapInput) 
  17. { 
  18. heap = heapInput; 
  19. index = 0; 
  20. } 
  21.  
  22. __device__ float& operator[](int idx) 
  23. { 
  24. return heap->h[idx]; 
  25. } 
  26.  
  27. __device__ void push_back(float data) 
  28. { 
  29. index = atomicAdd(heap->h,1); 
  30. heap->h[index]=data; 
  31. __syncthreads(); 
  32. } 
  33.  
  34. int index; 
  35. fake_heap * heap; 
  36. }; 
  37.  
  38. extern "C" 
  39. __global__ void parallel_vector_test(int * __restrict__ data) 
  40. { 
  41. const int i=threadIdx.x + blockIdx.x * blockDim.x; 
  42. __shared__ float memory[256]; 
  43. fake_heap heap(memory); 
  44. vector vect(&heap); // a vector that spans N threads 
  45. vect.push_back(i); // all threads push into same vector 
  46. data[i]=vect[i]; 
  47. }  

output of data array is (launched with 32-threads, pushed into same same vector of 32 elements):

  1. 0,8,1,9,2,10,3,11,4,12,5,13,6,14, 
  2. 7,15,16,24,17,25,18,26,19,27,20, 
  3. 28,21,29,22,30,23,31 

and not everyone wants it this way. Cela peut être n'importe quoi. Quelqu'un peut vouloir que les valeurs soient ordonnées de la même manière que les valeurs de thread-id, quelqu'un peut avoir besoin que la taille du vecteur ne soit pas limitée à quelques kilo-octets et reste sur la mémoire globale et plus probablement quelqu'un peut utiliser des structures vectorielles per-thread à la place, pour avoir plusieurs vecteurs indépendants.

Mais, comme vous pouvez le voir dans les définitions de "classe", vous pouvez utiliser "C++". (bien qu'un peu rogné dans les caractéristiques) dans le code du noyau. Regardez l'opérateur d'index surchargé. Ce n'est certainement pas du C99. À partir de ces capacités de type C++, vous pouvez construire des conteneurs avancés pour effectuer votre travail comme vous le souhaitez. For example, you can do this in a CUDA device function:

  1. float * mem = new float[256]; 
  2. delete mem; 

but this will be in local memory and possibly cached by L1 and is per-thread allocation/deletion but is too slow since there is (at least)256 stride-length between neighboring threads accessing their own mem buffer with same index. You may need a parallel allocator/indexer implementation for this to work fast.

When you include a header, it confuses its calling source

  1. __device__ 
  2. __global__ 
  3. __host__ 

these do not exist in definitions of functions in headers that are irrelevant to CUDA. You can try some CUDA compiler options to make it assume all functions in the header to be

  1. __device__ 

so that it may be called from device kernel.