PTX 中的变量“已降级”是啥意思？

Posted 2023-03-23

【中文标题】PTX 中的变量“已降级”是啥意思？

【英文标题】：What does it mean when a variable "has been demoted" in the PTX?PTX 中的变量“已降级”是什么意思？

【发布时间】：2021-10-03 04:13:11

【问题描述】：

在我的 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];

我的问题：

这样的变量是通过什么方式“降级”的？不是按照我的要求定义的吗？ 在什么情况下这些变量会被“降级”？

注意事项：

我使用的是 CUDA 11.2。 我从我的 PTX 中引用了去错线。实际名称是_ZZ8blahblah...。 我看到这个“降级”的数组变量要么是二维固定大小的数组，要么是具有结构的元素类型（例如struct unsigned short data[2]; ）；也许这有某种关系。

【问题讨论】：

@talonmies: 1. 我假设这具有标准含义；但根据你的建议，我会生成一个 MRE。 2. “全局->共享->全局”循环是什么意思？我从一个共享变量开始。

@RobertCrovella：这些变量可能无法优化到寄存器中（它们太大并且在编译时不知道访问索引）；但让我看看那个 MRE。

【参考方案1】：

根据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, %rd1;
        mov.u32         %r1, %tid.x;
        shl.b32         %r2, %r1, 2;
        mov.u32         %r3, s;
        add.s32         %r4, %r3, %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, release 11.2, 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, [_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, _ZZ1kPfE1s;
        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>;
        // demoted variable
        .shared .align 4 .b8 _ZZ2k1PfE1s[128];

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

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

【讨论】：

嗯，在我的例子中，这些变量是先验限制在单个函数中的——因为它们是在内核函数的主体中定义的。所以我的情况本质上是第二个k1。但是那么......那里发生的“降级”是什么？函数范围的共享变量出现在 PTX 和 CUDA 函数中。

降级是将变量的范围从隐式全局范围转换为该函数的局部范围。在这种情况下，该变量显然已经在本地范围内，但前端处理中也可能发生促销活动。您可能希望阅读我链接的整个主题。