初始化常量全局数组 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];
  ...
  }

您可以在

编程指南

中阅读有关__constant__变量的信息。 从设备代码（内核代码）的角度来看，

__constant__

变量是只读。 但从主机上，可以使用

cudaMemcpyToSymbol/cudaMemcpyFromSymbol

API 读取或写入它们。

编辑：根据您现在发布的代码，至少有 2 个错误：

1. 您对

   dev_v

   和

   dev_w

   的分配大小不正确。
2. 您没有

   w

   的主机分配。

通过这两个修复，以下代码似乎对我来说可以正常工作：

$ 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 代码时遇到问题，最好使用正确的 cuda 错误检查。
2. 您可以使用

   cuda-memcheck

   运行代码（就像我上面的那样），以快速了解是否遇到任何 CUDA 错误。
3. 我还没有验证数值结果或进行数学计算。 如果这不是你想要的，我想你可以解决它。
4. 除了我认为合理的修复明显错误并使结果可用于教育目的之外，我没有对您的代码进行任何更改。 当然可以讨论首选分配方法，

   printf

   与

   cout

   ，以及你有什么。 我在这个答案中主要关注 CUDA 主题。