CUDA 分析 - 高共享事务/访问但低本地重播率答案

【问题标题】：CUDA profiling - high shared transactions/access but low local replay rateCUDA 分析 - 高共享事务/访问但低本地重播率
【发布时间】：2019-01-25 02:48:13
【问题描述】：

运行 Visual Profiler 后，引导分析告诉我，我受内存限制，特别是我的共享内存访问对齐/访问不佳 - 基本上我访问共享内存的每一行都标记为每次访问约 2 个事务.

但是，我无法弄清楚 为什么 会出现这种情况（我的共享内存被填充/跨步，因此不应该发生银行冲突），所以我回去检查了共享重放指标 - 这表示只有 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 会使问题消失，我也不明白。

【问题讨论】：

不可能看到完成的代码。请添加一个。 minimal reproducible example你的问题
@talonmies 好吧，我宁愿知道可能一般会发生什么/该模式通常表明什么，但可以肯定的是，我会在早上看到我可以添加什么.
如果您要加载 8 字节数量的 warp-wide（例如 double、int2、float2 等），则每次访问将包含（至少）2 个事务.这并不意味着共享内存访问对齐不佳或有任何问题，这是正常的（在那种特定情况下）。每次访问 2 笔交易也不意味着（必然）存在银行冲突。

标签： optimization cuda shared-memory pycuda

【解决方案1】：

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

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

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

【讨论】：