参数传递过多问题

　　CUDA对kernel函数中参数大小要求不超过128kb，但是在SRIM程序的并行化中需要传入和传出很多参数，假如只对这些参数简单地罗列，其远远大于128kb，那么我们可以采用结构体将这些参数传入。用结构体从CPU到GPU进行传值时，对于单个变量是很简单的，但如果对数组传值就需要一定的技巧了，以下举例说明。

　　若结构体定义为

　　typedef struct
　　{float E0kev;
　　int *N;
　　…
　　}Test;Kernel函数定义为
　　__global__ void D_test(Test *d_test , int particleNum)
　　{
　　int i=threadIdx.x+blockDim.x*blockIdx.x;
　　if(i<particleNum)
　　{
　　cuPrintf("%d ",(*d_test).N[i];
　　}
　　}

　　结构体在CPU 端的定义及内存分配为：

　　Test *test1=(Test *)malloc(sizeof(Test));//cpu
　　(*test1).N=(int *)malloc(maxLayerNum*sizeof(int));

　　若在GPU端结构体内存分配写为：

　　Test * test2;
　　CUDA_SAFE_CALL(cudaMalloc((void**)&test2,sizeof(Test )));
　　CUDA_SAFE_CALL(cudaMal loc((void**)&(*t est2).N,
　　maxLayerNum*sizeof(int )));

　　则系统会报Segmentation fault错误。

　　在实现时需要定义中间变量，并在GPU端为数组分配内存，代码如下：

　　Test test2;
　　CUDA_ SAFE_CALL (cudaMal l oc(( void* *)&t es t2 .N,
　　maxLayerNum*sizeof(int )));
　　CUDA_SAFE_CALL(cudaMemcpy(test 2.N,(* t est1).N,
　　maxLayerNum*sizeof(int),cudaMemcpyHostToDevice));

　　然后对GPU上的结构体变量定义，并为其传值：

　　Test * d_test ;
　　CUDA_SAFE_CALL(cudaMalloc((void**)&d_test,sizeof(Test )));
　　CUDA_SAFE_CALL(cudaMemcpy(d_test,&test2,sizeof(Test),
　　cudaMemcpyHostToDevice));

　　SRIM并行化模拟结果

图1 H—>C在1000Kev时能损曲线对比

　　图1为H 打C在1000kev时能损曲线对比图，横轴表示靶材料的深度，纵轴表示能损，是由前面总能损MTOT[i]除以总粒子数所得。左图为运行SRIM软件的能损曲线图，右图为用CUDA书写的GPU并行程序能损曲线图，可以看出来两个图的结果基本相同，证明了并行程序的正确性。

表1 P—>C的CPU与GPU时间对比

　　表1列出了质子打碳时入射能分别在10kev、100kev、1000kev，入射粒子数分别为10000、100000、1000000时只统计能损和射程，运行全部程序的时间对比。可以看出，在粒子数比较少的情况下在CPU上运行要明显快于GPU，在粒子数比较多时GPU体现出了它的优越性，这是因为在粒子数比较少时，计算量很小，GPU在启动内核时占用很大一部分时间。并且采用该并行算法GPU的时间增长比较缓慢，那么我们可以看出粒子数越多该并行程序的优越性越明显。

　　本文主要以SRIM为物理背景，讨论了在CUDA架构下的并行编程，通过串行程序与并行程序的时间对比，可以看出GPU并行程序的优越性。CUDA仅是并行计算的基础，为了提高效率我们试图可以考虑MPI+OpenMP+CUDA的三级并行编程模式，即引入节点内的细粒度并行以及节点间的粗粒度并行机制，形成一种更高效的并行策略。

　　（作者单位：1为中国科学院近代物理研究所；2为兰州大学）