【发布时间】：2014-03-26 16:24:31
【问题描述】：
作为我论文工作的一部分，我正在从事一个 CUDA 项目（修改别人的代码、添加功能等）。作为 CUDA 的新手，这对我来说是一个真正的挑战。我正在使用 计算能力 1.3 卡，4 x Tesla C1060。遗憾的是，我遇到了平台的一些限制。
我需要将几个新结构传递给设备，我认为这些结构已正确复制。但是，当尝试在我的内核调用中将指针传递给设备上的结构时，我达到了 256 字节的限制（如question 中所述）。
我的代码是这样的：
// main.cu
static void RunGPU(HostThreadState *hstate)
{
SimState *HostMem = &(hstate->host_sim_state);
SimState DeviceMem;
TetrahedronStructGPU *h_root = &(hstate->root);
TetrahedronStructGPU *d_root;
TriangleFacesGPU *h_faces = &(hstate->faces);
TriangleFacesGPU *d_faces;
GPUThreadStates tstates;
unsigned int n_threads = hstate->n_tblks * NUM_THREADS_PER_BLOCK;
unsigned int n_tetras = hstate->n_tetras; // 9600
unsigned int n_faces = hstate->n_faces; // 38400
InitGPUStates(HostMem, h_root, h_faces, &DeviceMem, &tstates, hstate->sim,
d_root, d_faces, n_threads, n_tetras, n_faces );
cudaThreadSynchronize();
...
kernel<<<dimGrid, dimBlock, k_smem_sz>>>(DeviceMem, tstates, /*OK, these 2*/
d_root, d_faces);
// Limit of 256 bytes adding d_root and/or d_faces
cudaThreadSynchronize();
...
}
InitGPUStates 函数在另一个源文件中：
// kernel.cu
int InitGPUStates(SimState* HostMem, TetrahedronStructGPU* h_root,
TriangleFacesGPU* h_faces,
SimState* DeviceMem, GPUThreadStates *tstates,
SimulationStruct* sim,
TetrahedronStructGPU* d_root, TriangleFacesGPU* d_faces,
int n_threads, int n_tetras, int n_faces)
{
unsigned int size;
// Allocate and copy RootTetrahedron (d_root) on device
size = n_tetras * sizeof(TetrahedronStructGPU); // Too big
checkCudaErrors(cudaMalloc((void**)&d_root, size));
checkCudaErrors(cudaMemcpy(d_root, h_root, size, cudaMemcpyHostToDevice));
// Allocate and copy Faces (d_faces) on device
size = n_faces * sizeof(TriangleFacesGPU); // Too big
checkCudaErrors(cudaMalloc((void**)&d_faces, size));
checkCudaErrors(cudaMemcpy(d_faces, h_faces, size, cudaMemcpyHostToDevice));
...
}
我知道我只需要传递指向设备内存位置的指针。如何获取设备中的地址？这种指针传递是否正确完成？
两个新结构是：
// header.h
typedef struct {
int idx;
int vertices[4];
float Nx, Ny, Nz, d;
} TriangleFacesGPU;
typedef struct {
int idx, region;
int vertices[4], faces[4], adjTetras[4];
float n, mua, mus, g;
} TetrahedronStructGPU;
// other structures
typedef struct {
BOOLEAN *is_active;
BOOLEAN *dead;
BOOLEAN *FstBackReflectionFlag;
int *NextTetrahedron;
UINT32 *NumForwardScatters;
UINT32 *NumBackwardScatters;
UINT32 *NumBackwardsSpecularReflections;
UINT32 *NumBiases;
UINT32 *p_layer;
GFLOAT *p_x, *p_y, *p_z;
GFLOAT *p_ux, *p_uy, *p_uz;
GFLOAT *p_w;
GFLOAT *Rspecular;
GFLOAT *LocationFstBias;
GFLOAT *OpticalPath;
GFLOAT *MaxDepth;
GFLOAT *MaxLikelihoodRatioIncrease;
GFLOAT *LikelihoodRatioIncreaseFstBias;
GFLOAT *LikelihoodRatio;
GFLOAT *LikelihoodRatioAfterFstBias;
GFLOAT *s, *sleft;
TetrahedronStructGPU *tetrahedron;
TriangleFacesGPU *faces;
} GPUThreadStates;
typedef struct {
UINT32 *n_p_left;
UINT64 *x;
UINT32 *a;
UINT64 *Rd_ra;
UINT64 *A_rz;
UINT64 *Tt_ra;
} SimState;
kernel的定义是
__global__ void kernel(SimState d_state, GPUThreadStates tstates,
TetrahedronStructGPU *d_root,
TriangleFacesGPU *d_faces);
我将努力将SimState d_state 更改为指针传递SimState *d_state。以及GPUThreadStates tstates 到GPUThreadStates *tstates。
【问题讨论】：
-
kernel的声明是什么？例如，您似乎将tstates按值 传递给kernel。如果sizeof(GPUThreadStates)很大，您可以通过指针而不是值传递该结构来释放一些喘息空间。问题是，d_root和d_faces是 已经 指针。因此，如果仅添加这两个指针就超出了参数空间，则需要缩小要传递的其他内容的大小，例如DeviceMem(sizeof(SimState)) 和tstates(sizeof(GPUThreadStates)) .这也会影响引用这些实体的内核代码。 -
@RobertCrovella 你是对的。我不确定我是否正确传递了指针。内核定义
__global__ void MCMLKernel(SimState d_state, GPUThreadStates tstates, TetrahedronStructGPU *d_root, TriangleFacesGPU *d_faces)以及d_state和tstates都是按值传递的，不是吗？ -
是的，它们似乎是，尽管您实际上没有显示
GPUThreadStates和SimState的定义。如果它们的大小很大，阻止您添加d_root（指针）和d_faces（指针），那么您将不得不关注这些。 -
@RobertCrovella 再次感谢。我正在研究这些，
GPUThreadStates和SimState也很大。我在上面添加了这些定义。 -
@RobertCrovella 我发布了我所做的修改，作为更好格式化的答案。我有错误
code=11(cudaErrorInvalidValue) "cudaMalloc((void**)&DeviceMem->n_photons_left, size)"。我将衷心感谢您的帮助！谢谢！
标签： c++ pointers memory-management cuda