当 PTX 中的变量“已降级”时，这意味着什么？

编程问答 2022-09-04

问题描述

在我的 CUDA 内核的函数体中，我有几个固定大小的 __shared__ 数组变量。当我查看这些数组之一的已编译 PTX 代码 (SM 7.5) 时，我看到一条评论说：

// my_kernel(t1 p1,t2 p2)::my_variable has been demoted

... 并且此行出现在 PTX 中的 .global 行之间，就在编译内核本身之前。然后，在内核中，我得到：

// demoted variable
.shared .align 4 .b8 my_kernel(t1 p1,t2 p2)::my_variable [1234];

我的问题：

  1. 这样的变量如何“降级”？是不是像我要求的那样定义？
  2. 在什么情况下这些变量会被“降级”？

注意事项：

  • 我使用的是 CUDA 11.2。
  • 我引用了 PTX 中的乱码。实际名称为 _ZZ8blahblah...。
  • 我看到这种“降级”的数组变量要么是二维固定大小的数组，要么具有结构体的元素类型（例如 struct { unsigned short data[2]; }）；也许这在某种程度上是相关的。

解决方法

根据讨论 here 这似乎是基于共享变量范围是否可以限制为单个函数（即单个内核）。即使是非常复杂的具有共享用法的内核函数，我看到的共享变量也会降级。

这是一个简单的例子，它被降级和不降级。

未降级：

$ vi t1.cu
$ cat t1.cu
__shared__ float s[32];
__global__ void k(float *  my_ptr){
  s[threadIdx.x] = threadIdx.x;
  *my_ptr = s[threadIdx.x];
}
__global__ void k1(float *  my_ptr){
  s[threadIdx.x] = 0.0f;
  *my_ptr = s[threadIdx.x];
}
$ nvcc -ptx t1.cu
$ cat t1.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-29373293
// Cuda compilation tools,release 11.2,V11.2.67
// Based on NVVM 7.0.1
//

.version 7.2
.target sm_52
.address_size 64

        // .globl       _Z1kPf
.shared .align 4 .b8 s[128];

.visible .entry _Z1kPf(
        .param .u64 _Z1kPf_param_0
)
{
        .reg .f32       %f<2>;
        .reg .b32       %r<5>;
        .reg .b64       %rd<3>;


        ld.param.u64    %rd1,[_Z1kPf_param_0];
        cvta.to.global.u64      %rd2,%rd1;
        mov.u32         %r1,%tid.x;
        cvt.rn.f32.u32  %f1,%r1;
        shl.b32         %r2,%r1,2;
        mov.u32         %r3,s;
        add.s32         %r4,%r3,%r2;
        st.shared.f32   [%r4],%f1;
        st.global.f32   [%rd2],%f1;
        ret;

}
        // .globl       _Z2k1Pf
.visible .entry _Z2k1Pf(
        .param .u64 _Z2k1Pf_param_0
)
{
        .reg .b32       %r<6>;
        .reg .b64       %rd<3>;


        ld.param.u64    %rd1,[_Z2k1Pf_param_0];
        cvta.to.global.u64      %rd2,%tid.x;
        shl.b32         %r2,%r2;
        mov.u32         %r5,0;
        st.shared.u32   [%r4],%r5;
        st.global.u32   [%rd2],%r5;
        ret;

}

降级：

$ cat t1.cu
__global__ void k(float *  my_ptr){
  __shared__ float s[32];
  s[threadIdx.x] = threadIdx.x;
  *my_ptr = s[threadIdx.x];
}
__global__ void k1(float *  my_ptr){
  __shared__ float s[32];
  s[threadIdx.x] = 0.0f;
  *my_ptr = s[threadIdx.x];
}
$ nvcc -ptx t1.cu
$ cat t1.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-29373293
// Cuda compilation tools,V11.2.67
// Based on NVVM 7.0.1
//

.version 7.2
.target sm_52
.address_size 64

        // .globl       _Z1kPf
// _ZZ1kPfE1s has been demoted
// _ZZ2k1PfE1s has been demoted

.visible .entry _Z1kPf(
        .param .u64 _Z1kPf_param_0
)
{
        .reg .f32       %f<2>;
        .reg .b32       %r<5>;
        .reg .b64       %rd<3>;
        // demoted variable
        .shared .align 4 .b8 _ZZ1kPfE1s[128];

        ld.param.u64    %rd1,_ZZ1kPfE1s;
        add.s32         %r4,%f1;
        ret;

}
        // .globl       _Z2k1Pf
.visible .entry _Z2k1Pf(
        .param .u64 _Z2k1Pf_param_0
)
{
        .reg .b32       %r<6>;
        .reg .b64       %rd<3>;
        // demoted variable
        .shared .align 4 .b8 _ZZ2k1PfE1s[128];

        ld.param.u64    %rd1,_ZZ2k1PfE1s;
        add.s32         %r4,%r5;
        ret;

}

顺便说一句，似乎确实存在 PTX 生成器可以完全删除共享变量的情况，但这与这里的问题没有直接关系。

cudacudanvccptx