CUDA中可以用什麼來替代Vector類？

01-22

__global__ void setInsideColorKernel(float3 *v0,float3 *v1,float3 *v2,byte *color,int3 *comparevertices){
int k=threadIdx.x; int size=128; int ix=0; int iy=0; int iz=0; int offz=0; int offy=0; int index =0; long indexOutput =0; int r=0; int g=0; int b=0; int insidenum=0; int diffx=0; int diffy=0; int diffz=0; double diff01x,diff01y,diff01z; diff01x=v1[k].x-v0[k].x; diff01y=v1[k].y-v0[k].y; diff01z=v1[k].z-v0[k].z; double diff01=sqrt(diff01x*diff01x+diff01y*diff01y+diff01z*diff01z); double diff02x,diff02y,diff02z; diff02x=v2[k].x-v0[k].x; diff02y=v2[k].y-v0[k].y; diff02z=v2[k].z-v0[k].z; double diff02=sqrt(diff02x*diff02x+diff02y*diff02y+diff02z*diff02z); double a=1/(diff01); double c=1/(diff02); float3 insidepoint;
for(double i=0;i&<=1;i+=a) { for(double j=0;j+i&<=1;j+=c) { insidepoint.x=i*(diff01x)+j*(diff02x)+v0[k].x; insidepoint.y=i*(diff01y)+j*(diff02y)+v0[k].y; insidepoint.z=i*(diff01z)+j*(diff02z)+v0[k].z; for(int m=0;m&<48;m++) { if((insidepoint.x&>=comparevertices[m].x-size/2) (insidepoint.x&<=comparevertices[m].x+size/2) (insidepoint.y&>=comparevertices[m].y-size/2) (insidepoint.y&<=comparevertices[m].y+size/2) (insidepoint.z&>=comparevertices[m].z-size/2) (insidepoint.z&<=comparevertices[m].z+size/2)) { diffx=comparevertices[m].x-insidepoint.x; diffy=comparevertices[m].y-insidepoint.y; diffz=comparevertices[m].z-insidepoint.z; ix=(size/2)-diffx; iy=(size/2)-diffy; iz=(size/2)-diffz; if(iz==0) offz = size * size * iz; else offz = size * size * (iz-1); if(iy==0) offy = size * iy; else offy = size * (iy-1); index=ix + offy + offz; indexOutput=index * 3; r=(int)color[indexOutput]; g=(int)color[indexOutput+1]; b=(int)color[indexOutput+2]; } } } } }
由於無法在__global__函數中直接將r,g,b寫入vector容器中（據我所知CUDA暫不支持vector類，或者如果有方法可以將vector類插入進來的話還請大神指點），需要將每個循環產生的r,g,b三個值回傳到CPU上並儲存到vector容器中。請問有什麼方法可以在CUDA中替代vector類先做一個類似容器的功能對循環產生的r,g,b值進行暫存再傳回CPU中。

如＠呂文龍所說，你可以使用Thrust中的device_ptr[1]。這是個模板類所以你可以直接放float3進去。需要使用Raw Pointer的時候調用raw_pointer_cast就可以。但是如果你搜索之後會發現，幾乎沒有人使用這種辦法。原因是這樣會構造一個AoS(Array of Structures)。你目前使用的辦法是SoA(Structure of Arrays)。兩種結構在GPU上的效率對比有很多討論，比較簡潔明了的可以參考[2]中的第一個答案。

我給你一些其他性能優化的建議：

1）如果你的項目對inverse sqrt的精度要求不高（比如，可以容忍千分之一的誤差），那麼建議使用Quake3里的快速inv_sqrt[3]。可以考慮寫成一個__inline__ __device__函數。

2) 關於每個線程，每個block處理多少數據的問題。你目前的代碼默認一個線程處理一個三角形（如果我理解無誤）。GPU上的SM數量是有限的，一個SM上可以同時active的block數量也是有限的（在Fermi應該是8個，Kepler下是16個）。要完全保證GPU的occupancy，需要對每個線程處理的數據數量，每個block的線程數量，每個線程使用的寄存器（本地變數）多少以及總的block數量進行適當的分配。特別的，每個線程處理數據數量應該遵照[4]中的原則。祝你好運。

[1] thrust: thrust::device_ptr&< T

[2] How to structure data for optimal speed in a CUDA app

[3] Fast inverse square root

[4] CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops

thrust