《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个线程的线程竞争规模能很好的处理,而此算法用到了多次线程同步,从而拖慢了速度)
阅读更多
版权声明:本文为博主原创文章,未经博主允许不得转载。 https://blog.csdn.net/qq_27582707/article/details/58016306
文章标签: C++ cuda
个人分类: C++ cuda
想对作者说点什么? 我来说一句

CUDA并行程序设计-GPU编程指南

2017年11月12日 12.64MB 下载

没有更多推荐了,返回首页

加入CSDN,享受更精准的内容推荐,与500万程序员共同成长!
关闭
关闭