Блог

Встречаются ли литеральные константы в ядрах, хранящихся в постоянной памяти? Например.,

 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
$
  

Вероятно, вы можете провести такой анализ самостоятельно, если потратите некоторое время на документацию

Основные инструкции, представляющие интерес, следующие:

         /*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];