0

我正在分析以下 CUDA 内核

__global__ void fftshift_2D(double2 *data, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    if (i < N1 && j < N2) {
        double a = pow(-1.0, (i+j)&1);

        data[j*blockDim.x*gridDim.x+i].x *= a;
        data[j*blockDim.x*gridDim.x+i].y *= a;
    }
 }

它基本上将二维双精度复数数据矩阵乘以标量双精度变量。

可以看出，我正在执行合并的全局内存访问，我想通过 NVIDIA Visual Profiler 通过检查全局内存负载和存储效率来验证这一点。令人惊讶的是，这样的效率都恰好是 50%，与合并内存访问的预期 100% 相去甚远。这与复数实部和虚部的交错存储有关吗？如果是这样，我可以利用任何技巧来恢复 100% 的效率吗？

先感谢您。

附加信息

BLOCK_SIZE_x=16
BLOCK_SIZE_y=16

dim3 dimBlock2(BLOCK_SIZE_x,BLOCK_SIZE_y);
dim3 dimGrid2(N2/BLOCK_SIZE_x + (N2%BLOCK_SIZE_x == 0 ? 0:1),N1/BLOCK_SIZE_y + (N1%BLOCK_SIZE_y == 0 ? 0:1));

N1 和 N2 可以是任意偶数。

该卡是 NVIDIA GT 540M。

4

2 回答 2

5

看看这篇关于各种内存访问模式效率的 NVIDIA 博客文章。您遇到了跨步内存访问问题。

由于每个组件都是独立使用的，因此您可以将double2数组视为普通的普通double数组（就像Robert Crovella 建议的那样）。

__global__ void fftshift_2D(double *data, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    if (i < N1 * 2 && j < N2) {
        double a = pow(-1.0, (i / 2 + j)&1);
        data[j*blockDim.x*gridDim.x+i] *= a;
    }
}

但是，如果您需要在单个线程中同时访问 x 和 y 组件，您可以尝试：

使用 2 个单独的数组。一个带有 x 分量 一个带有 y 分量。像那样：

__global__ void fftshift_2D(double *dataX, double *dataY, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    if (i < N1 && j < N2) {
        double a = pow(-1.0, (i+j)&1);

        dataX[j*blockDim.x*gridDim.x+i] *= a;
        dataY[j*blockDim.x*gridDim.x+i] *= a;
    }
}

或者保持数据布局不变，但将其加载到共享内存中并从共享内存中重新洗牌。那看起来或多或少是这样的：

__global__ void fftshift_2D(double2 *data, int N1, int N2)
{
    __shared__ double buff[BLOCK_SIZE*2];
    double2 *buff2 = (double2 *) buff;
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
    double ptr = (double *) &data[j*blockDim.x*gridDim.x + blockDim.x * blockIdx.x];

    // TODO add guarding with N1 & N2
    buff[threadIdx.x] = ptr[threadIdx.x];
    buff[blockDim.x + threadIdx.x] = ptr[blockDim.x + threadIdx.x];
    __syncthreads();

    double a = pow(-1.0, (i+j)&1);
    buff2[threadIdx.x].x *= a 
    buff2[threadIdx.x].y *= a 

    __syncthreads();
    ptr[threadIdx.x] = buff[threadIdx.x];
    ptr[blockDim.x + threadIdx.x] = buff[blockDim.x + threadIdx.x];
}
于 2013-01-10T08:36:00.237 回答
4

是的，因为您有一个结构数据存储格式的数组，并且您仅使用以下行引用所有其他元素：

    data[j*blockDim.x*gridDim.x+i].x *= a;

那么作为结果发生的全局加载和全局存储将各自只有 50% 的利用率。请注意，我认为缓存在这里应该有所帮助，因为您在下一行引用了备用元素。但是加载/存储效率仍然是 50%。

我相信你可以使用某种方法来解决这个问题（对于这个特定的例子）*data：

double *mydata = (double *)data;
...
mydata[2*(j*blockDim.x*gridDim.x)+i] *= a;

请注意，我并不是想确切地展示如何获得相同的覆盖范围，只是说明这个想法。上面的代码大致是需要的，但您需要调整代码以确保正确处理要相乘的所有元素。

于 2013-01-09T22:18:52.727 回答