最近需要将之前的CUDA优化代码复盘,得到每一种优化下的加速比,因此,每一次优化后需要对比C,CUDA的结果,赶巧不巧,遇上了这么一种情况:
核函数功能 :核函数D_SDF1_SRC1_MultiChan
需要在一个Block
中完成多路的执行,因为可以共用数据,利用共享内存提供更快和更少次数的内存访问;其Block图如图所示:
于是希望在每一行的Thread中取不同的常数,使用threadIdx.y
作为标识:
__global__ void SRC1Kernel_MultiChan(short *d_sr1, float *d_hsrc1, complexType *d_ssrc1, double *d_fc)
{
int bx = blockIdx.x; //unique i
int tx = threadIdx.x; //unique m
int ty = threadIdx.y; //unique fc
int stride = gridDim.x;
int stride_SM = blockDim.x * ty;
__shared__ float add_re[Q * ChanNum];
__shared__ float add_im[Q * ChanNum];
__shared__ double fc_kernel[ChanNum];
if (tx == 0) fc_kernel[ty] = d_fc[ty];
__syncthreads();
for (int i = bx; i < OLENGTH - Q; i += stride) {
int t1 = tx + ((D * i) % I) * Q;
int t2 =