注:本人代碼是對長度為1024的數組排序; block數量為1,threads數量為256
p135頁代碼應改為:
__device__ void merge_array(u32 *src,u32 *dest,u32 num_lists,u32 num_elements,u32 tid){
u32 reduction_shift=3;
u32 reduction_size=8;
u32 per_list=num_elements/num_lists;
u32 data=src[tid*per_list];
u32 s_idx=tid>>reduction_shift;
u32 num_reductions=num_lists/reduction_size;
__shared__ u32 list_index[256];
list_index[tid]=0;
__shared__ u32 min_val[32];
__shared__ u32 min_tid;
for(u32 i=0;i<num_elements;i++){
if(tid<num_lists){
min_val[s_idx]=0xFFFFFFFF;
min_tid=0xFFFFFFFF;
}
__syncthreads();
atomicMin(&min_val[s_idx],data);
if(num_reductions>0){
__syncthreads();
if(tid<num_reductions)
atomicMin(&min_val[0],min_val[tid]);
__syncthreads();
}
if(min_val[0]==data)
atomicMin(&min_tid,tid);
__syncthreads();
if(tid==min_tid){
list_index[tid]++;
dest[i]=data;
if(list_index[tid]<per_list)
data=src[tid*per_list+list_index[tid]];
else
data=0xFFFFFFFF;
}
__syncthreads();
}
}
(在gtx970上此代碼沒有純使用atomicMin()(即不采用規約)的算法快,餘竊以為是因為在新的裝置上256個線程的線程競争規模能很好的處理,而此算法用到了多次線程同步,進而拖慢了速度)