我有以下CUDA内核,它似乎非常"难以"优化:
__global__ void DataLayoutTransformKernel(cuDoubleComplex* d_origx, cuDoubleComplex* d_origx_remap, int n, int filter_size, int ai ) { for(int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < filter_size; idx+=blockDim.x * gridDim.x) { int index = (idx * ai) & (n-1); d_origx_remap[idx] = d_origx[index]; } } //Parameters were defined before int permute[loops] = {29165143,3831769,17603771,9301169,32350975, ...} int n = 33554432; int filter_size = 1783157; for(int i=0; i>>((cuDoubleComplex*) d_origx,(cuDoubleComplex*)d_origx_remap+i*filter_size, n, filter_size, permute[i]); }
内核的目的是重新排序d_origx[]
从不规则到常规(d_origx_remap
)的数据布局.内核使用不同的访问步幅(ai
)启动多次.
这里的挑战是引用数组中的不规则内存访问模式d_origx[index]
.我的想法是使用共享内存.但是对于这种情况,使用共享内存来合并全局内存访问似乎非常困难.
有没有人有关于如何优化这个内核的建议?