Cuda 中存储的文字浮点/整数常量在哪里答案

【问题标题】：Where are literal float/int constants stored in CudaCuda 中存储的文字浮点/整数常量在哪里
【发布时间】：2020-09-05 23:51:16
【问题描述】：

内核中遇到的字面常量是否存储在常量内存中？例如，

y = x * 3.14f; // x,y are floats.

3.14f 是否存储在常量内存中并在乘法之前复制到寄存器中？

【问题讨论】：

标签： cuda

【解决方案1】：

根据我的测试，该常量似乎是在指令流中编码的（即它是与指令一起获取的）：

$ cat t1799.cu
__global__ void k(float *d){

  float y, x = *d;
  y = x * 3.14f;
  *d = y;
}
$ nvcc -c t1799.cu
$ cuobjdump -sass t1799.o

Fatbin elf code:
================
arch = sm_30
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

        code for sm_30
                Function : _Z1kPf
        .headerflags    @"EF_CUDA_SM30 EF_CUDA_PTX_SM(EF_CUDA_SM30)"
                                                                           /* 0x22f04283f2804307 */
        /*0008*/                   MOV R1, c[0x0][0x44];                   /* 0x2800400110005de4 */
        /*0010*/                   MOV R2, c[0x0][0x140];                  /* 0x2800400500009de4 */
        /*0018*/                   MOV R3, c[0x0][0x144];                  /* 0x280040051000dde4 */
        /*0020*/                   LD.E R0, [R2];                          /* 0x8400000000201c85 */
        /*0028*/                   FMUL32I R4, R0, 3.1400001049041748047;  /* 0x310123d70c011c02 */
        /*0030*/                   ST.E [R2], R4;                          /* 0x9400000000211c85 */
        /*0038*/                   EXIT;                                   /* 0x8000000000001de7 */
        /*0040*/                   BRA 0x40;                               /* 0x4003ffffe0001de7 */
        /*0048*/                   NOP;                                    /* 0x4000000000001de4 */
        /*0050*/                   NOP;                                    /* 0x4000000000001de4 */
        /*0058*/                   NOP;                                    /* 0x4000000000001de4 */
        /*0060*/                   NOP;                                    /* 0x4000000000001de4 */
        /*0068*/                   NOP;                                    /* 0x4000000000001de4 */
        /*0070*/                   NOP;                                    /* 0x4000000000001de4 */
        /*0078*/                   NOP;                                    /* 0x4000000000001de4 */
                .................



Fatbin ptx code:
================
arch = sm_30
code version = [6,5]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
$

如果您花一些时间在documentation 上，您可能可以自己进行这种分析

感兴趣的主要说明如下：

        /*0020*/      LD.E R0, [R2]; // float y, x = *d;  R0 is x
        /*0028*/      FMUL32I R4, R0, 3.1400001049041748047; // y = x * 3.14f; R4 is y
        /*0030*/      ST.E [R2], R4;   //       *d = y;

指令语法FMUL32I R4, R0, 3.1400001049041748047; 清楚地表明3.14 常量在指令中被编码为立即操作数：

 FMUL32I
 ||  | Immediate
 ||  32-bit
 |Multiply
 Floating point

代码在内核的前 3 条指令中访问常量内存，但唯一有意义的是：

MOV R2, c[0x0][0x140];

将d 参数的地址加载到寄存器中以进行寄存器间接加载。这些来自常量内存的负载都与3.14 常量无关。

3.14f 是否存储在常量内存中

没有。

并在乘法之前复制到寄存器中？

没有。

【讨论】：