c++ - Cuda:将主机数据复制到共享内存阵列

Question

我在主机和设备上定义了一个结构。在主机中，我用值初始化了这个结构的数组。

MyStruct *h_s = (MyStruct *) malloc(objsize*sizeof(MyStruct));
hs[0] = ...

Mystruct *d_s;
cudaMalloc( &d_s, objsize * sizeof(MyStruct));
cudaMemcpy( d_s, h_s, objsize * sizeof(MyStruct), cudaMemcpyHostToDevice );
init<<< gridSize, blockSize >>> ( d_s );

在我的内核中，我有大约 7 个函数应该使用这个数组。其中一些是全局的，一些是简单的设备功能。为了简单和高效，我想使用共享内存阵列。

__shared__ Mystruct *d_s;

__global__ void init(Mystruct *theStructArray){
   //How to allocate memory for d_s
   //How copy theStructArray to d_s
}

所以问题是:如何为共享数组分配内存并使用函数参数设置其值？

编辑:我正在尝试将 smallpt 代码写入 cuda。

struct Sphere {
double rad;       // radius
Vec p, e, c;      // position, emission, color
Refl_t refl;      // reflection type (DIFFuse, SPECular, REFRactive)

Sphere(){
    rad = 16.5;
    p = (Vec(27,16.5,47) + Vec(73,16.5,78))*0.5;
    e = Vec();
    c = Vec(0.75, 0.75, 0.75);
    refl = DIFF;
}

Sphere(double rad_, Vec p_, Vec e_, Vec c_, Refl_t refl_):
rad(rad_), p(p_), e(e_), c(c_), refl(refl_) {}

__device__ double intersect(const Ray &r) const { // returns distance, 0 if nohit
    Vec op = p-r.o; // Solve t^2*d.d + 2*t*(o-p).d + (o-p).(o-p)-R^2 = 0
    double t, eps=1e-4, b=op.dot(r.d), det=b*b-op.dot(op)+rad*rad;
    if (det<0) return 0; else det=sqrt(det);
    return (t=b-det)>eps ? t : ((t=b+det)>eps ? t : 0);
} 

};

Answer 1

如果您了解共享内存的范围和大小限制，那么问题似乎是

1. 如何为共享内存数组动态预留内存
2. 如何在内核中使用动态共享内存

你的内核变成这样:

__shared__ Mystruct *d_s;

__global__ void init(Mystruct *theStructArray){

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

    // load to shared memory array
    // assumes Mystruct has correct copy assignment semantics
    d_s[threadIdx.x] = theStructArray[tid]

    __syncthreads();

    // Each thread has now loaded one value to the block
    // scoped shared array
}

[免责声明:在浏览器中编写的代码，从未编译或测试，并注意注释中有关复制分配的警告]

调用主机代码需要向内核调用添加一个额外的参数来为共享数组保留内存:

MyStruct *h_s = (MyStruct *) malloc(objsize*sizeof(MyStruct));
hs[0] = ...

Mystruct *d_s;
cudaMalloc( &d_s, objsize * sizeof(MyStruct));
cudaMemcpy( d_s, h_s, objsize * sizeof(MyStruct), cudaMemcpyHostToDevice );
init<<< gridSize, blockSize, blockSize * sizeof(MyStruct) >>> ( d_s );

注意 <<< >>> 的第三个参数内核调用的节。这指定了每个 block 保留的内存字节数。硬件对您可以进行的共享内存分配的大小有限制，它们可能会对超出硬件限制的性能产生额外影响。

共享内存是 CUDA 的一个非常有据可查的功能，我会推荐 Mark Harris's blog还有这个 Stack Overflow Question作为 CUDA 中共享内存机制的良好起点。