《cuda并行程序设计》勘误（2）

注：本人代码是对长度为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个线程的线程竞争规模能很好的处理，而此算法用到了多次线程同步，从而拖慢了速度）