Аргументы не соответствуют инструкциям 'ld' и 'add'

Question

Я хотел бы знать, почему следующие asm инструкции внутри кода cuda

__global__ void access( double *posArray )
{
  uint32_t tid = threadIdx.x;
  double sink = 0;
  for(uint32_t i = tid; i < L1_SIZE; i += THREADS_NUM) {
    double* ptr = posArray+i;
    asm volatile("{\t\n"
                 ".reg .f32 data;\n\t"
                 "ld.global.ca.f64 data, [%1];\n\t"
                 "add.f64 %0, data, %0;\n\t"
                 "}" : "+d"(sink) : "l"(ptr) : "memory"
    );
  }
  // synchronize all threads
  asm volatile("bar.sync 0;");

  for(uint32_t i = 0; i < L1_SIZE; i += THREADS_NUM) {
    double* ptr = posArray+i;
    // every warp loads all data in l1 cache
    for(uint32_t j = 0; j < THREADS_NUM; j += WARP_SIZE) {
      uint32_t offset = (tid+j)%THREADS_NUM/;
      asm volatile("{\t\n"
                   ".reg .f64 data;\n\t"
                   "ld.global.ca.f64 data, [%1];\n\t"
                   "add.f64 %0, data, %0;\n\t"
                   "}" : "+d"(sink) : "l"(ptr+offset) : "memory"
      );
    }
  }
}

не могут быть скомпилированы с таким выводом:

ptxas /tmp/tmpxft_00003451_00000000-5_l1.ptx, line 63; error   : Arguments mismatch for instruction 'ld'
ptxas /tmp/tmpxft_00003451_00000000-5_l1.ptx, line 64; error   : Arguments mismatch for instruction 'add'
...
...

И эти две строки повторяются.

Answer 1

Эта строка:

             ".reg .f32 data;\n\t"

должна быть изменена на:

             ".reg .f64 data;\n\t"

(что соответствует второму использованию / вызову встроенной сборки.)

Ошибка возникает из-за того, что data является аргументом для следующих 2 инструкций:

             "ld.global.ca.f64 data, [%1];\n\t"
             "add.f64 %0, data, %0;\n\t"

, и обе инструкции работают с 64-битными операндами с плавающей запятой (.f64), а не с 32-битными плавающимиточечные операнды (.f32).