Is it correct to inspect the destionation regsiter value of ldmatrix PTX

Accelerated Computing CUDA CUDA Programming and Performance
1 

__global__ void kernel(int8_t *src_A, int8_t *src_B, void *dst) {

// the smem
  __shared__ int8_t src_a[128];
  __shared__ int8_t src_b[128];

/* some code has been omitted */

  uint64_t smem_src_a = reinterpret_cast<uint64_t>(src_a);
  uint64_t smem_src_b = reinterpret_cast<uint64_t>(src_b);
    
  int32_t ret;
  int32_t r_a, r_b;

  asm volatile (
    "ldmatrix.sync.aligned.m8n8.x1.shared::cta.b16 {%0}, [%1];"
    : "=r"(r_a)
    : "l"(smem_src_a)
  );
  asm volatile (
    "ldmatrix.sync.aligned.m8n8.x1.shared::cta.b16 {%0}, [%1];"
    : "=r"(r_b)
    : "l"(smem_src_b)
  );

printf("%d", r_a); // this code cause illegal memory access, but why?

  return;
}

as i’m understand, ldmatrix gonna write the corresponding fragment of matrix value from shared mem(which has been initialized in other emitted code) to destination register. why it cannot be read?

2

The error is most likely caused by the code before printf.
If printf is commented out, the code before is probably optimized away, which is why it appears that without printf there is no error.

3

thanks for you point out and it solved. i error caused becasue shared memory pointer didnt converted. i have referred to this doc: CUDA C++ Programming Guide