copy

Synchronous element copy src → dst between global, shared, and register (local) memory. CUDA currently registers eight variants: five explicit fixed-width variants, ldstmatrix, vec_auto, and fallback. The vec_auto implementation contains separate global/shared and register paths; gmem_smem and reg below name those implementation paths, not selectable dispatch variants.

Variant	Pair	Prio	Lowering
vec_16b / vec_32b / vec_64b / vec_128b / vec_256b	global ↔ shared/local, or shared ↔ local	20	explicit thread-scope transfer of exactly the named width; optional global-load cache controls
vec_auto: copy → vec_auto global/shared path path	global ↔ shared	10	synthesized [outer, threads, vec] partition with direct-PTX vector loads and stores
vec_auto: copy → vec_auto register path path	register ↔ shared/global	10	partition induced by the register layout’s thread axes
copy → ldstmatrix	register ↔ shared	10	warp-collective ldmatrix / stmatrix (m8n8 fragments)
copy → fallback	global / shared / local	0	scalar single-thread copy (last resort)

The detailed pages cover the two vec_auto paths, ldstmatrix, and the fallback, including accepted input, lowering algorithm, emitted TIRx IR, and generated CUDA:

• copy → vec_auto global/shared path
• copy → vec_auto register path
• copy → ldstmatrix
• copy → fallback