CUDA求序列最大值

淩亂°似流年 2022-04-13 02:41 313阅读 0赞

CUDA并行计算中，求最大值和求和原理一样，都属于规约算法。我是求复数模长的最大值。

核函数如下：

__global__ void array_max(const Complex* __restrict__ in, float* __restrict__ blockMax, const int in_len)
    {
    	extern __shared__ float sha_partialMax2[];
    	const int tid = threadIdx.x;
    	const int gTid = blockIdx.x * (blockDim.x * 2) + tid;
    
    	// 复制全局数据到共享内存
    
    	Complex in0, in1;
    	in0 = in[gTid];
    	sha_partialMax2[tid] = in0.x*in0.x + in0.y*in0.y;
    
    	in1 = in[gTid + blockDim.x];
    	sha_partialMax2[tid + blockDim.x] = in1.x*in1.x + in1.y*in1.y;
    
    	// block内前半部分与后半部分对应比较大小	
    	if (blockDim.x > 512)
    	{
    		__syncthreads();
    		if (sha_partialMax2[tid] < sha_partialMax2[tid + 1024])
    			sha_partialMax2[tid] = sha_partialMax2[tid + 1024];
    	}
    	if (blockDim.x > 256 && tid < 512)
    	{
    		__syncthreads();
    		if (sha_partialMax2[tid] < sha_partialMax2[tid + 512])
    			sha_partialMax2[tid] = sha_partialMax2[tid + 512];
    	}
    	if (tid < 256)
    	{
    		__syncthreads();
    		if (sha_partialMax2[tid] < sha_partialMax2[tid + 256])
    			sha_partialMax2[tid] = sha_partialMax2[tid + 256];
    	}
    	if (tid < 128)
    	{
    		__syncthreads();
    		if (sha_partialMax2[tid] < sha_partialMax2[tid + 128])
    			sha_partialMax2[tid] = sha_partialMax2[tid + 128];
    	}
    	if (tid < 64)
    	{
    		__syncthreads();
    		if (sha_partialMax2[tid] < sha_partialMax2[tid + 64])
    			sha_partialMax2[tid] = sha_partialMax2[tid + 64];
    	}
    
    	if (tid < 32)//步长stride小于等于warp
    	{
    		__syncthreads();
    		volatile float* vol_sh_max = sha_partialMax2;//volatile 确保每次对共享内存的修改都即时修改
    		float register val_t;
    		for (int stride = 32; stride > 0; stride >>= 1)
    		{
    			val_t = vol_sh_max[tid + stride];
    			if (vol_sh_max[tid] < val_t)
    				vol_sh_max[tid] = val_t;
    		}
    	}
    
    	if (tid == 0)
    	{
    		blockMax[blockIdx.x] = sqrtf(sha_partialMax2[0]);
    	}
    }

当活动的线程数少于一个线程束（Warp）时，指令的执行遵循SIMD（单指令多数据流）模式，也就是说在活动线程数少于32个时，我么不需要进行同步控制，但需要volatile关键字修饰，确保每次修改都即时生效。当线程数少于一个warp时也无需判断tid是否小于步长，我们让多余的线程做无效的操作，避免每次判断增加降低效率。