初始化常量全局數組CUDA C

Question

我遇到問題了！我需要在cuda c中初始化一個常量全局數組。要初始化數組，我需要使用for！我需要這樣做，因爲我必須在某些內核中使用這個數組，並且我的教授告訴我只能在設備中定義爲常量可見。

我該怎麼做？

我想要做這樣的事情：

#include <stdio.h> 
#include <math.h> 
#define N 8 

__constant__ double H[N*N]; 

__global__ void prodotto(double *v, double *w){ 

     int k=threadIdx.x+blockDim.x*blockIdx.x; 

     w[k]=0; 
     for(int i=0;i<N;i++) w[k]=w[k]+H[k*N+i]*v[i]; 
} 

int main(){ 

     double v[8]={1, 1, 1, 1, 1, 1, 1, 1}; 
     double *dev_v, *dev_w, *w; 
     double *host_H; 
     host_H=(double*)malloc((N*N)*sizeof(double)); 
     cudaMalloc((void**)&dev_v,sizeof(double)); 
     cudaMalloc((void**)&dev_w,sizeof(double)); 

     for(int k=0;k<N;k++){ 
      host_H[2*N*k+2*k]=1/1.414; 
      host_H[2*N*k+2*k+1]=1/1.414; 
      host_H[(2*k+1)*N+2*k]=1/1.414; 
      host_H[(2*k+1)+2*k+1]=-1/1.414; 

     } 

     cudaMemcpyToSymbol(H, host_H, (N*N)*sizeof(double)); 
     cudaMemcpy(dev_v, v, N*sizeof(double), cudaMemcpyHostToDevice); 
     cudaMemcpy(dev_w, w, N*sizeof(double), cudaMemcpyHostToDevice); 

     prodotto<<<1,N>>>(dev_v, dev_w); 

     cudaMemcpy(v, dev_v, N*sizeof(double), cudaMemcpyDeviceToHost); 
     cudaMemcpy(w, dev_w, N*sizeof(double), cudaMemcpyDeviceToHost); 


     for(int i=0;i<N;i++) printf("\n%f %f", v[i], w[i]); 

     return 0; 
    } 

但產量爲零的數組...我想輸出數組來填充矩陣H的產品（這裏看作是一個數組）和數組v。 謝謝!!!!!

Answer 1

像這樣的東西應該工作：

#define DSIZE 32 
__constant__ int mydata[DSIZE]; 

int main(){ 
    ... 
    int *h_mydata; 
    h_mydata = new int[DSIZE]; 
    for (int i = 0; i < DSIZE; i++) 
    h_mydata[i] = ....; // initialize however you wish 
    cudaMemcpyToSymbol(mydata, h_mydata, DSIZE*sizeof(int)); 
    ... 
} 

並不困難。然後，您可以在內核直接使用__constant__數據：

__global__ void mykernel(...){ 
    ... 
    int myval = mydata[threadIdx.x]; 
    ... 
    } 

您可以閱讀programming guide約__constant__變量。從設備代碼（內核代碼）的角度來看，__constant__變量爲只讀。但是從主機，他們可以使用cudaMemcpyToSymbol/cudaMemcpyFromSymbolAPI來讀取或寫入。

編輯：立足現在你已經發布的代碼，至少有2個錯誤：

1. 您分配尺寸dev_v和dev_w是不正確的。
2. 您沒有爲w分配主機。

下面的代碼似乎爲我與2個修復正常工作：

$ cat t579.cu 
#include <stdio.h> 
#include <math.h> 
#define N 8 

__constant__ double H[N*N]; 

__global__ void prodotto(double *v, double *w){ 

     int k=threadIdx.x+blockDim.x*blockIdx.x; 

     w[k]=0; 
     for(int i=0;i<N;i++) w[k]=w[k]+H[k*N+i]*v[i]; 
} 

int main(){ 

     double v[N]={1, 1, 1, 1, 1, 1, 1, 1}; 
     double *dev_v, *dev_w, *w; 
     double *host_H; 
     host_H=(double*)malloc((N*N)*sizeof(double)); 
     w  =(double*)malloc( (N)*sizeof(double)); 
     cudaMalloc((void**)&dev_v,N*sizeof(double)); 
     cudaMalloc((void**)&dev_w,N*sizeof(double)); 

     for(int k=0;k<N;k++){ 
      host_H[2*N*k+2*k]=1/1.414; 
      host_H[2*N*k+2*k+1]=1/1.414; 
      host_H[(2*k+1)*N+2*k]=1/1.414; 
      host_H[(2*k+1)+2*k+1]=-1/1.414; 

     } 

     cudaMemcpyToSymbol(H, host_H, (N*N)*sizeof(double)); 
     cudaMemcpy(dev_v, v, N*sizeof(double), cudaMemcpyHostToDevice); 
     cudaMemcpy(dev_w, w, N*sizeof(double), cudaMemcpyHostToDevice); 

     prodotto<<<1,N>>>(dev_v, dev_w); 

     cudaMemcpy(v, dev_v, N*sizeof(double), cudaMemcpyDeviceToHost); 
     cudaMemcpy(w, dev_w, N*sizeof(double), cudaMemcpyDeviceToHost); 


     for(int i=0;i<N;i++) printf("\n%f %f", v[i], w[i]); 
     printf("\n"); 
     return 0; 
    } 
$ nvcc -arch=sm_20 -o t579 t579.cu 
$ cuda-memcheck ./t579 
========= CUDA-MEMCHECK 

1.000000 0.000000 
1.000000 -0.707214 
1.000000 -0.707214 
1.000000 -1.414427 
1.000000 1.414427 
1.000000 0.707214 
1.000000 1.414427 
1.000000 0.707214 
========= ERROR SUMMARY: 0 errors 
$ 

的幾個注意事項：

1. 您在使用CUDA代碼的麻煩任何時候，它的使用proper cuda error checking的良好做法。
2. 您可以使用cuda-memcheck（就像我上面所述）運行您的代碼以快速瞭解是否遇到任何CUDA錯誤。
3. 我還沒有驗證數字結果或通過數學工作。如果這不是你想要的，我認爲你可以把它整理出來。
4. 我沒有對你的代碼進行任何修改，除了對我來說修復明顯的錯誤並使結果適合教育目的顯得有意義之外。當然可以討論優先分配方法，printf與cout，以及你有什麼。在這個答案中，我主要關注CUDA主題。