.cg Cache modifier with async copy size of less than 16

Accelerated Computing CUDA CUDA NVCC Compiler
1

Hi all,

I was just wondering if it’s documented anywhere that when using the cp.async instruction on Ampere targets, the copy size must be 16 bytes when using the cg cache modifier. PTX documentation appears to omit this, providing only examples with a copy size of 16 (I have only checked the cp.async section).

Surprisingly, the .ca cache modifier does not suffer from such limitations.

Thus the following code snippet does not compile (but wll compile if you switch to using .ca):

__device__ void load_smem(const int64_t gmem, const int32_t smem) {
    asm volatile("cp.async.cg.shared.global [%0], [%1], 4;\n" :: "r"(smem), "l"(gmem) : "memory");
}

Thus, my question is: Is this expected behavior or a compiler bug by any chance?

Thanks a Lot !

2

It seems to be documented. The first code block there are not examples, that is “syntax”. The 16 there indicates that cp-size is not applicable to that case, and the parameter in that position must be 16. The “examples” is a code block that comes later in that section.

cp.async

Initiates an asynchronous copy operation from one state space to another.

Syntax

cp.async.ca.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], cp-size{, src-size}{, cache-policy} ;
cp.async.cg.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], 16{, src-size}{, cache-policy} ;
cp.async.ca.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], cp-size{, ignore-src}{, cache-policy} ;
cp.async.cg.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size} [dst], [src], 16{, ignore-src}{, cache-policy} ;

.level::cache_hint = { .L2::cache_hint }
.level::prefetch_size = { .L2::64B, .L2::128B, .L2::256B }
cp-size = { 4, 8, 16 }