CUDA 分析 - 高共享 transactions/access 但本地重放率低

Question

在运行 Visual Profiler 之后，引导分析告诉我我受内存限制，尤其是我的共享内存访问很差 aligned/accessed - 基本上我访问共享内存的每一行被标记为每次访问 ~2 个事务。

但是，我想不通为什么会这样（我的共享内存是padded/strided所以不应该有银行冲突），所以我回去检查了共享重播指标 - 表明只有 0.004% 的共享访问被重播。

那么，这是怎么回事，我应该注意什么来加速我的内核？

编辑：最小复制：

import numpy as np
import pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
import pycuda.gpuarray as gp

mod = SourceModule("""

（拆分代码块以获得 Python 和 CUDA/C++ 着色）

typedef unsigned char ubyte;

__global__ void identity(ubyte *arr, int stride) 
{
    const int dim2 = 16;
    const int dim1 = 64;
    const int dim0 = 33;
    int shrstrd1 = dim2;
    int shrstrd0 = dim1 * dim2;
    __shared__ ubyte shrarr[dim0 * dim1 * dim2];

    auto shrget = [shrstrd0, shrstrd1, &shrarr](int i, int j, int k) -> int{ 
        return shrarr[i * shrstrd0 + j * shrstrd1 + k]; 
    };

    auto shrset = [shrstrd0, shrstrd1, &shrarr](int i, int j, int k, ubyte val) -> void {
        shrarr[i * shrstrd0 + j * shrstrd1 + k] = val;
    };

    int in_x = threadIdx.x;
    int in_y = threadIdx.y;

    shrset(in_y, in_x, 0, arr[in_y * stride + in_x]);
    arr[in_y * stride + in_x] = shrget(in_y, in_x, 0);
}
""",

（同上）

options=['-std=c++11'])

#Equivalent to identity<<<1, dim3(32, 32, 1)>>>(arr, 64);
identity = mod.get_function("identity")
identity(gp.zeros((64, 64), np.ubyte), np.int32(64), block=(32, 32, 1))

每次访问 2 个事务，共享重放开销 0.083。将 dim2 减少到 8 会使问题消失，我也不明白。

Answer 1

部分回答：我对共享内存库的工作方式存在根本性的误解（即，它们是 banks，每个大约有一千个字节库），所以没有意识到它们是循环的，所以过多的填充意味着 32 行元素最终可能会不止一次地使用每个 bank。

不过，据推测，这种冲突并不是每次都出现 - 相反，它出现了，哦，从数字来看，一个区块大约出现 85 次。

我会把它留在这里一天，希望得到更完整的解释，然后关闭并接受这个答案。

CUDA 分析 - 高共享 transactions/access 但本地重放率低

CUDA profiling - high shared transactions/access but low local replay rate

optimization

cuda

shared-memory

pycuda