Cuda中存储的文字float / int常量在哪里

问题描述

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

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

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

解决方法

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

$ 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,3.1400001049041748047; // y = x * 3.14f; R4 is y
        /*0030*/      ST.E [R2],R4;   //       *d = y;

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

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

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

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

正在将d参数的地址加载到用于寄存器间接加载的寄存器中。常量存储器中的所有负载都与3.14常量无关。

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

否。

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

否。