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

好久没更博客了。

最近在学cuda c/c++,正在看一本叫做《cuda并行程序设计》的书，在京东上买的（应该不是盗版吧）风评不错。但是我这么一看啊....发现代码的错误简直就是遍地都是.....然后我给cudabook@gmail.com发邮件，也不知道是没空理我还是这个邮箱已经荒废了...没人理我。所以呢，我决定把目前发现的错误发出来，以免耽误其他人的学习。

就是这本书。

看着不像盗版吧。

（120页代码也有错，我没有改）

P122页 GPU版本的基排序代码应改为：

__device__ void radix_sort(unsigned int *array_,const unsigned int numlists,const unsigned int elements,const size_t tid,unsigned int * sort_tmp){
	for(u32 bit=0;bit<32;bit++){
	u32 mask=(1<<bit);
	u32 tmp0count=0;
	u32 tmp1count=0;
	u32 per_list=elements/numlists;
	for(u32 i=0;i<per_list;i++){
		if((i+tid*per_list)<elements){
			u32 elem=array_[i+tid*per_list];
			if((elem&mask)>0){
				sort_tmp[tmp1count+tid*per_list]=elem;
				tmp1count++;
				}
			else{
				array_[tmp0count+tid*per_list]=elem;
				tmp0count++;
				}
			}		
		}
	
	for(u32 i=0;i<tmp1count;i++){
		array_[tmp0count+i+tid*per_list]=sort_tmp[i+tid*per_list];
		}

	}
	__syncthreads();
}

p124页 find_min函数应改为

u32 my_find_min(u32* src,u32 *index,u32 num_lists,u32  nums_per_list){
	u32 min_val=0xFFFFFFFF;
	u32 min_idx=0;
	for(u32 i=0;i<num_lists;++i){
		if(index[i]<nums_per_list){
			u32 src_idx=i*nums_per_list+index[i];
			u32 data=src[src_idx];
			if(data<min_val){
				min_val=data;
				min_idx=i;
			}
		}
	}
	index[min_idx]++;
	return min_val;
}

p125页copy_data_to_shared函数应改为：

__device__ void copy_to(const u32*data,u32* sort_tmp,u32 num_per_list,u32 tid){
	for(u32 i=0;i<num_per_list;i++){
		sort_tmp[i+tid*num_per_list]=data[i+tid*num_per_list];
	}
	__syncthreads();
}

p128页 并行合并代码应改为：

__device__ void merge_single_thread( u32 * src,u32 *dest,u32 tid,u32 num_lists,u32 num_elements){
	const u32 nums_per_list=num_elements/num_lists;
	u32 list_indexes[256];
	list_indexes[tid]=0;
	__syncthreads();
	if(tid==0){
		for(u32 i=0;i!=num_elements;++i){
			u32 min_val=0xFFFFFFFF;
			u32 min_idx=0;
			for(u32 i=0;i<num_lists;++i){
				if(list_indexes[i]<nums_per_list){
					u32 src_idx=i*nums_per_list+list_indexes[i];
					u32 data=src[src_idx];
					if(data<min_val){
						min_val=data;
						min_idx=i;
					}
				}
			}
		dest[i]=min_val;
		list_indexes[min_idx]++;
		}
	}
}

p132页代码应当改成：

__device__ void merge_multi_threads(u32 * src,u32 *dest,u32 tid,u32 num_lists,u32 num_elements){
	__shared__ u32 list_index[256];
	list_index[tid]=0;
	__syncthreads();
	u32 per_list=num_elements/num_lists;
	for(u32 i=0;i<num_elements;i++){
		__shared__ u32 min_val;
		__shared__ u32 min_tid;
		u32 data;
		if(list_index[tid]<per_list){
			 data=src[list_index[tid]+tid*per_list];
		}
		else
			data=0xFFFFFFFF;
		if(tid==0){
			min_val=0xFFFFFFFF;
			min_tid=0xFFFFFFFF;
		}
		__syncthreads();
		atomicMin(&min_val,data);
		__syncthreads();
		if(min_val==data){
			atomicMin(&min_tid,tid);
		}
		__syncthreads();
		if(min_tid==tid){
			list_index[min_tid]++;
			dest[i]=min_val;
		}
	}


}

注：本人代码是对长度为1024的数组排序； block数量为1，threads数量为256