如果您打算在具有最新版本的 CUDA 的计算能力 2.x 或 3,x 设备上运行此代码,那么您的内核代码几乎是正确的。new
Fermi 和 Kepler 硬件上的 CUDA 4.x 和 5.0 支持C++运算符。请注意,使用new
或malloc
在设备上的运行时堆上分配的内存。它具有创建上下文的生命周期,但您目前无法直接从 CUDA 主机 API(通过cudaMemcpy
或类似方式)访问它。
我把你的结构和内核变成了一个简单的示例代码,你可以自己尝试看看它是如何工作的:
#include <cstdio>
struct myStruct {
float *data;
};
__device__
void fill(float * x, unsigned int n)
{
for(int i=0; i<n; i++) x[i] = (float)i;
}
__global__
void kernel(myStruct *input, const unsigned int imax)
{
for(unsigned int i=0,N=1; i<imax; i++, N*=2) {
float * p = new float[N];
fill(p, N);
input[i].data = p;
}
}
__global__
void kernel2(myStruct *input, float *output, const unsigned int imax)
{
for(unsigned int i=0,N=1; i<imax; i++, N*=2) {
output[i] = input[i].data[N-1];
}
}
inline void gpuAssert(cudaError_t code, char * file, int line, bool Abort=true)
{
if (code != 0) {
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
if (Abort) exit(code);
}
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
int main(void)
{
const unsigned int nvals = 16;
struct myStruct * _s;
float * _f, * f;
gpuErrchk( cudaMalloc((void **)&_s, sizeof(struct myStruct) * size_t(nvals)) );
size_t sz = sizeof(float) * size_t(nvals);
gpuErrchk( cudaMalloc((void **)&_f, sz) );
f = new float[nvals];
kernel<<<1,1>>>(_s, nvals);
gpuErrchk( cudaPeekAtLastError() );
kernel2<<<1,1>>>(_s, _f, nvals);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaMemcpy(f, _f, sz, cudaMemcpyDeviceToHost) );
gpuErrchk( cudaDeviceReset() );
for(int i=0; i<nvals; i++) {
fprintf(stdout, "%d %f\n", i, f[i]);
}
return 0;
}
需要注意的几点:
- 此代码只能在 Fermi 或 Kepler GPU 上使用 CUDA 4.x 或 5.0 编译和运行
- 您必须将 GPU 的正确架构传递给 nvcc 进行编译(例如,我曾经
nvcc -arch=sm_30 -Xptxas="-v" -o dynstruct dynstruct.cu
在 linux 上为 GTX 670 编译)
- 示例代码使用“收集”内核将数据从运行时堆中的结构复制到主机 API 可以访问的分配,以便打印结果。这是针对我之前提到的
cudaMemcpy
无法直接从运行时堆内存中的地址复制的限制的一种解决方法。我希望这可能会在 CUDA 5.0 中得到解决,但最新的候选版本仍然有这个限制。