分析 CUDA 代码:合并内存读取时出现意外指令计数

Posted

技术标签:

【中文标题】分析 CUDA 代码:合并内存读取时出现意外指令计数【英文标题】:Profiling CUDA code: Unexpected instruction counts on coalesced memory reads 【发布时间】:2013-08-27 11:03:27 【问题描述】:

我正在为小输入数据(= 512 个元素)分析一个非常转储的排序算法。我正在调用一个从结构数组中读取合并的内核。

结构如下所示:

struct __align__(8) Elements 

     float weight;
     int value;
;

nvprof 为 L1 未命中/命中和 gdl 指令提供以下指令计数:

                  Invocations    Avg       Min       Max  Event Name
        Kernel: sort(Elements*)
                      500         0         0         0  gld_inst_8bit
                      500         0         0         0  gld_inst_16bit
                      500      1024      1024      1024  gld_inst_32bit
                      500         0         0         0  gld_inst_64bit
                      500         0         0         0  gld_inst_128bit
                      500       120       120       120  l1_global_load_hit
                      500       120       120       120  l1_global_load_miss
                      500         0         0         0  uncached_global_load_tr.

如果我改变结构的布局如下:

struct __align__(8) Elements 

     float weight;
     float value;
;

分析输出如下所示:

                  Invocations    Avg       Min       Max  Event Name
Device 0
        Kernel: sort(Elements*)
                      500         0         0         0  gld_inst_8bit
                      500         0         0         0  gld_inst_16bit
                      500         0         0         0  gld_inst_32bit
                      500       512       512       512  gld_inst_64bit
                      500         0         0         0  gld_inst_128bit
                      500         0         0         0  l1_global_load_hit
                      500       120       120       120  l1_global_load_miss
                      500         0         0         0  uncached_global_load_tr.

对执行时间没有任何影响,但我不明白为什么 GPU 在代码的第一个变体上执行 32 位加载指令而在第二个变体上执行 64 位指令。

内核使用 1 个块和 512 个线程调用(因此 l1_global_load_x 计数器可能不正确)。一切都在配备 CUDA 5.0 的 GeForce 480 上进行。

编辑: 排序内核(有点缩短):

__global__ void sort(Elements* nearest)

    ThreadIndex idx = index();

    __shared__ Elements temp[MAX_ELEMENTS];
    __shared__ int index_cache[MAX_ELEMENTS];

    temp[idx.x] = nearest[idx.x];

    WeightedElements elem = temp[idx.x];
    __syncthreads();

    int c = 0;

    // some index crunching 

    nearest[idx.x] = temp[c];

【问题讨论】:

可以添加你的内核代码吗? GPU 在第一种情况下执行 32 位加载,因为编译器在这种情况下生成 32 位加载指令(在第二种情况下生成 64 位加载。)我猜你的问题是“为什么当我的结构依次具有两个float 类型时编译器会生成64 位加载,但是当我的结构有float 后跟int 时会加载两个32 位? 【参考方案1】:

造成这种情况的基本原因归结于编译器的代码生成。 PTX 汇编器具有用于浮点和整数的不同虚拟寄存器状态空间,并且(我认为)不可能将 64 位加载到不同状态空间的两个寄存器中。因此编译器在混合整数/浮点结构中发出两个 32 位加载,但在浮点/浮点结构情况下可以将 64 位向量加载到两个寄存器中。

这可以通过考虑以下代码模型来说明:

struct __align__(8) ElementsB 

    float weight;
    float value;
;

struct __align__(8) ElementsA 

    float weight;
    int value;
;

template<typename T>
__global__ void kernel(const T* __restrict__ in, T* __restrict__ out, bool flag)

    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    T ival = in[idx];
    if (flag) 
        out[idx] = ival;
    



template __global__ void kernel<ElementsA>(const ElementsA *, ElementsA *, bool);
template __global__ void kernel<ElementsB>(const ElementsB *, ElementsB *, bool);

这里我们有您提到的两种结构,以及为这两种类型实例化的简单模板内核。如果我们查看编译器为 sm_20(CUDA 5.0 发行版编译器)发出的 PTX,差异是显而易见的。对于ElementsA 实例:

    ld.param.u32    %r4, [_Z6kernelI9ElementsAEvPKT_PS1_b_param_0];
    ld.param.u32    %r5, [_Z6kernelI9ElementsAEvPKT_PS1_b_param_1];
    ld.param.u8     %rc1, [_Z6kernelI9ElementsAEvPKT_PS1_b_param_2];
    cvta.to.global.u32      %r1, %r5;
    cvta.to.global.u32      %r6, %r4;
    .loc 2 16 1
    mov.u32         %r7, %ntid.x;
    mov.u32         %r8, %ctaid.x;
    mov.u32         %r9, %tid.x;
    mad.lo.s32      %r2, %r7, %r8, %r9;
    .loc 2 18 1
    shl.b32         %r10, %r2, 3;
    add.s32         %r11, %r6, %r10;
    ld.global.u32   %r3, [%r11+4];  // 32 bit integer load
    ld.global.f32   %f1, [%r11];  // 32 bit floating point load

(为强调而添加了cmets)

对于Element B 实例:

    ld.param.u32    %r3, [_Z6kernelI9ElementsBEvPKT_PS1_b_param_0];
    ld.param.u32    %r4, [_Z6kernelI9ElementsBEvPKT_PS1_b_param_1];
    ld.param.u8     %rc1, [_Z6kernelI9ElementsBEvPKT_PS1_b_param_2];
    cvta.to.global.u32      %r1, %r4;
    cvta.to.global.u32      %r5, %r3;
    .loc 2 16 1
    mov.u32         %r6, %ntid.x;
    mov.u32         %r7, %ctaid.x;
    mov.u32         %r8, %tid.x;
    mad.lo.s32      %r2, %r6, %r7, %r8;
    .loc 2 18 1
    shl.b32         %r9, %r2, 3;
    add.s32         %r10, %r5, %r9;
    ld.global.v2.f32        %f9, %f10, [%r10];  // 64 bit float2 load

两者之间没有性能损失的原因是底层硬件使用 128 字节获取来进行合并的扭曲级别负载,并且在这两种情况下,事务都会导致同一对 128 字节获取。

【讨论】:

以上是关于分析 CUDA 代码:合并内存读取时出现意外指令计数的主要内容,如果未能解决你的问题,请参考以下文章

使用队列和侦听器登录到控制台时出现意外行为

使用 Numba 进行矩阵乘法时出现 CUDA 内存不足错误

从标准输入读取参数时出现意外行为

从套接字读取时出现意外行为

直接打印到文本显存时出现意外输出

在 CUDA 的 __device__ 函数中使用动态分配时出现“未知错误”