鴻鵠國際提供SRIM與GPU的軟硬體解決方案。 我們提供整體的針對
國際上通用的計算帶電粒子在靶材料中入射和能量傳遞的計算機程序 提供相對應的諮詢與技術整合。 並且提供相關人員的培訓與人力的調度。 SRIM介紹 SRIM（the stopping and range of ions into matter）系列軟件基於Monte Carlo方法,是目前國際上通用的計算帶電粒子在靶材料中入射和能量傳遞的計算機程序。其原名為TRIM（the Transport of Ions in Matter）1985年由JF Ziegler JP Biersack和U.littmark編寫，主要計算入射粒子在物質中的能損，射程等。以後在此基礎上不斷升級完善，並於1991年更名為SRIM。 SRIM對入射原子在靶中的運動分析採用連續慢化假設： 1.入射原子與材料靶原子核的碰撞採用兩體彈性碰撞來描述, 這一部分主要導致入射原子運動軌蹟的偏折, 能量損失來自彈性碰撞的能量損失部分。 2.兩次碰撞之間的距離以及碰撞後的參數通過隨機抽樣得到。 3.在兩次碰撞之間，認為入射原子與材料中的電子作用，連續均勻地損失能量。 粒子的位置、能量損失以及次級粒子的各種參數都在整個跟踪過程中存儲下來,最後得到各種所需物理量的期望值和相應的統計誤差。
SRIM並行化中的一些技術 那麼代碼框架可寫為如下： … #define yn (100) … __global__ void mcc(參數){ int tid=threadIdx.x+blockDim.x*blockIdx.x; int i=tid; … while(i<particleNum&&tid<blockDim.x*gridDim.x) { … temp=(i%blockDim.x*gridDim.x)*yn+m-1;//m>=0&&m<yn MTOT[temp]=MTOT[temp]+E_loss; … i+= blockDim.x*gridDim .x; } } int main(){ int blockdim=16; int threaddim=16; int n_HN_total= yn* blockdim* threaddim; … mcc<<<blockdim,threaddim>>>( d_MTOT，…參數); CUDA_SAFE_CALL(cudaThreadSynchronize ()); CUDA_ SAFE_CAL L (cudaMemcp y( h_MTOT,d_MTOT, n_HN_total*sizeof(float),cudaMemcpyDeviceToHost)); for(int u=0;u<thread_total;u++){ for(int v=0;v<yn ;v++){ MTOT[v]=MTOT[v]+h_MTOT[u*yn+v]; … } } … } 這種實現方法的好處是：1.無論線程數增加到多大，該代碼都是可行的；2.其占用的內存是固定的不會隨著線程數的增加而增大。 參數傳遞過多問題 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為蘭州大學） |
分子動力學 >
