1

我试图在 GPU 上实现 FDTD 方程。我最初实现了使用全局内存的内核。内存合并不是那么好。因此，我实现了另一个使用共享内存来加载值的内核。我正在研究1024x1024.

代码如下

__global__ void update_Hx(float *Hx, float *Ez, float *coef1, float* coef2){
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
    __shared__ float  Ez_shared[BLOCKSIZE_HX][BLOCKSIZE_HY + 1];
    /*int top = offset + x_index_dim;*/
    if(threadIdx.y == (blockDim.y - 1)){
        Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
        Ez_shared[threadIdx.x][threadIdx.y + 1] = Ez[offset + x_index_dim];
   }
    else{
        Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];
    }
}

常数BLOCKSIZE_HX=16和BLOCKSIZE_HY= 16。

当我运行视觉分析器时，它仍然说内存没有合并。

编辑：我使用 GT 520 显卡，cuda 计算能力为 2.1。我的全局 L2 事务/访问 =7.5即有245 760用于 32768执行该行的 L2 事务Ez_shared[threadIdx.x][threadIdx.y] = Ez[offset];

Global memory load efficiency是50%。

Global memory load efficiency=100 * gld_requested_throughput/ gld_throughput

尽管我的线程正在查看 16 个连续值，但我无法弄清楚为什么会有如此多的内存访问。有人可以指出我做错了什么吗？

编辑：感谢所有帮助。

4

1 回答 1

1

您的内存访问模式是这里的问题。您只能获得 50% 的效率（对于 L1 和 L2），因为您正在访问 16 个浮点数的连续区域，即 64 个字节，但 L1 事务大小为 128 个字节。这意味着对于请求的每 64 个字节，必须将 128 个字节加载到 L1（因此也加载到 L2）。

您也有共享内存库冲突的问题，但这目前不会对您的全局内存加载效率产生负面影响。

您可以通过多种方式解决负载效率问题。最简单的方法是将 x 维度块大小更改为 32。如果这不是一个选项，您可以更改全局内存数据布局，以便每两个连续的 blockIdx.y（[0, 1]、[2,3] 等） ) 值将映射到一个连续的内存块。如果即使这样也不是一个选项，并且您必须只加载一次全局数据，您可以使用非缓存全局内存加载来绕过 L1 - 这将有所帮助，因为 L2 使用 32 字节事务，因此您的 64 字节将在两个 L2 中加载无开销的事务。

于 2013-02-12T10:59:24.887 回答