PTX ISA 완전판 단권화

1 Virtual ISA 가상 명령어 집합

• PTX는 stable contract — 세대 건너 forward-compat
• SASS는 세대마다 opcode·encoding 변경 ↗ V04 §12
• LLVM IR과 유사하나 GPU-specialized (warp·shared·predicate)

2 Compilation Flow ★ .cu→PTX→SASS

.cu  --(nvcc front)-->  .ptx
.ptx --(ptxas)-->        .cubin (SASS)
.cubin ──pack──>         fatbinary
fatbin + embedded PTX ──>  application

// runtime (JIT) path:
driver.loadModule(ptx) → ptxas(JIT) → SASS

Source: CUDA C PG Compilation 섹션 (13.1)

3 Arch String -arch / -code

4 Forward Compatibility

• fatbin = SASS(여러 arch) + PTX(1 virtual arch) 번들
• 일치하는 SASS 없으면 PTX가 드라이버에서 JIT
• JIT 캐시: ~/.nv/ComputeCache

5 PTX Version ↔ Arch

cf. PTX ISA 8.4 (CUDA 12.4) 기준

6 왜 PTX를 읽는가

1. inline asm로 intrinsic 없는 명령 호출
2. 컴파일러 출력 검증 (의도한 fma가 났는가)
3. CUTLASS 등 라이브러리 커널의 실제 동작 추적
4. register pressure·unroll 결과 확인

7 PTX ↔ SASS 대응

• PTX mad.lo.s32 → SASS IMAD
• PTX ld.global → SASS LDG.E.128
• PTX mma.sync → SASS HMMA.16816.F32
• 상세 매핑 ↗ V04 §13

8 PTX 생성 트리거

# .ptx 파일 직접 보기
nvcc -ptx kernel.cu -arch=sm_80

# PTX + SASS 모두
nvcc -keep kernel.cu -arch=sm_80
# → kernel.ptx, kernel.cubin 생성

# cubin에서 SASS 역)
cuobjdump --dump-sass kernel.cubin

9 핵심 용어 한 줄

• PTX: virtual ISA text (human-readable)
• SASS: real ISA binary (GPU가 실제 실행)
• ptxas: PTX → SASS compiler
• fatbinary: multi-arch 번들
• JIT: driver가 runtime에 PTX 컴파일

1 파일 골격 ★ ver · target · addr · func

// PTX skeleton
.version 7.8
.target  sm_80
.address_size 64

.visible .entry mykernel(
  .param .u64 mykernel_param_0,
  .param .u32 mykernel_param_1
)
{
  .reg .b32 %r<8>;
  .reg .b64 %rd<4>;
  .reg .pred %p<4>;
  // body
  ret;
}

2 Directive . 으로 시작

3 Instruction 형식 ★

// 예
@%p1 add.s32 %r3, %r1, %r2;
     mad.lo.s32 %r5, %r3, %r4, %r5;
     ld.global.nc.f32 %f1, [%rd1];
     st.shared.v4.f32 [%rs0], {%f0,%f1,%f2,%f3};

4 Type Modifier (!) uspbf

5 Register Naming convention

compiler 생성 PTX의 관례. inline asm은 자유.

6 Special Register %tid %ctaid %ntid %nctaid

1 State Space 정의

2 전체 목록 ★ reg · sreg · param · local · global · shared · const · tex

3 물리 매핑 실제 HW

A100 기준. 정확 수치 ↗ V02 §14

4 .param 상세

• kernel 인자는 constant memory bank에 배치
• 최대 4KB (sm_70~), sm_90+는 32KB
• device 함수의 .param은 ABI에 따라 stack 또는 reg

ld.param.u64 %rd1, [mykernel_param_0];
cvta.to.global.u64 %rd2, %rd1;
ld.global.f32 %f1, [%rd2];

5 Generic vs Specific Addr cvta · cvta.to

// shared → generic
cvta.shared.u64 %rd_gen, %rd_sh;
// generic → global
cvta.to.global.u64 %rd_g, %rd_gen;

6 .shared 선언

.shared .align 16 .b8
  smem_buf[16384];

// dynamic (extern shmem)
.extern .shared .align 16
  .b8 smem_dyn[];

• 정적 shmem: __shared__ float buf[N] 대응
• 동적 shmem: extern __shared__ ..., 런치 3번째 인자로 크기
• CTA 내 모든 thread가 같은 주소 공간

1 정수 산술 add sub mul mad div rem

2 Float 산술 ★

3 Rounding Mode (!) rn · rz · rm · rp

• fma/mul/add는 rounding 필수
• .ftz: flush subnormal to zero
• .sat: clamp to [0,1] (f32) 또는 type range

4 FMA의 의미 핵심

5 Logical & Bit and or xor not shl shr

6 setp / selp ★ compare → pred → select

// compare
setp.lt.s32 %p1, %r1, %r2;
// select: %r3 = %p1 ? %r4 : %r5
selp.s32 %r3, %r4, %r5, %p1;

// predicated
@%p1 add.s32 %r3, %r1, %r2;

7 Type 변환 cvt

// f32 → s32 (truncate)
cvt.rzi.s32.f32 %r, %f;
// f16 → f32
cvt.f32.f16 %f, %h;
// u32 → f32 (round-nearest)
cvt.rn.f32.u32 %f, %r;

• float ↔ int: rounding mode 필수
• float wide→narrow: rounding + optional .sat
• packed: cvt.rn.f16x2.f32 (두 f32 → f16x2)

1 ld / st 형식 ★

ld.global.f32 %f1, [%rd1];
ld.global.v4.f32 {%f0,%f1,%f2,%f3}, [%rd1];
ld.global.nc.f32 %f1, [%rd1];  // non-coherent
st.global.wt.f32 [%rd1], %f1;  // write-through
st.shared.v2.b32 [%rs1], {%r0,%r1};

2 Cache Modifier ★ load ca · cg · cs · lu · cv · nc

sm_80+: .nc는 __ldg intrinsic 대응

3 Cache Modifier store

4 Vector Load ★ .v2 .v4

// 128-bit load (4 × f32)
ld.global.v4.f32
  {%f0,%f1,%f2,%f3}, [%rd1];
// 128-bit load (2 × f64)
ld.global.v2.f64
  {%fd0,%fd1}, [%rd1];
// 64-bit load
ld.global.v2.f32
  {%f0,%f1}, [%rd1];

• peak bandwidth에 필수: 128-bit transaction
• 주소 align: v4.f32 → 16B align
• destination regs는 연속 번호 권장

5 Alignment 규칙

misaligned → 여러 transaction으로 분해, BW 손실

6 Volatile & Weak

• ld.volatile.global: 최적화 불가, cache bypass 의도
• ld.weak: relaxed consistency 명시 (PTX 6+)
• 일반 ld는 relaxed (acquire/release 아님)
• ordering 필요 시 ld.acquire / fence 병행 ↗ §12

1 cp.async 정의 ★ L1 bypass · ld→reg 거치지 않음

• 기존: ld.global → reg → st.shared (2-hop)
• 신규: cp.async (1-hop, DMA-like)
• async → prefetch 파이프라인 구성 가능

2 기본 형태

// 16B copy, cache global (L1 bypass)
cp.async.cg.shared.global
  [smem_addr], [gmem_addr], 16;

// 8B copy, cache all
cp.async.ca.shared.global
  [smem_addr], [gmem_addr], 8;

// 4B copy
cp.async.ca.shared.global
  [smem_addr], [gmem_addr], 4;

• byte size: 4 / 8 / 16만 허용
• .ca: L1+L2 · .cg: L2만 (큰 복사에 유리)
• 16B 경우 align 16 필수

3 commit / wait ★

// 현재까지 issue한 cp.async들을
// 하나의 async-group으로 묶음
cp.async.commit_group;

// 가장 오래된 그룹부터 최대 N개만 남기고 대기
cp.async.wait_group N;

// 모든 그룹 완료 대기
cp.async.wait_all;

4 Byte Size × Cache 매트릭스

warp 32 × 16B = 512B/issue, 128B 트랜잭션 4개로 퍼짐

5 Double Buffering Loop ★ 전체 pseudo-code

// STAGES = 2: double buffer
// STAGES = 3~: multi-stage
const S = 3;
.shared .align 16 .b8 sA[S][TILE_BYTES];
.shared .align 16 .b8 sB[S][TILE_BYTES];

// prologue: S-1 stage prefetch
for (s = 0; s < S-1; ++s) {
  cp_async_tile(sA[s], gA[s]);
  cp_async_tile(sB[s], gB[s]);
  cp.async.commit_group;
}

// main loop
for (k = 0; k < K_TILES; ++k) {
  // issue next stage (overlap)
  if (k + S - 1 < K_TILES) {
    cp_async_tile(sA[(k+S-1) % S],
                  gA[k+S-1]);
    cp_async_tile(sB[(k+S-1) % S],
                  gB[k+S-1]);
  }
  cp.async.commit_group;

  // wait so [k] is ready
  cp.async.wait_group S-2;
  bar.sync 0;  // __syncthreads

  // compute on stage k
  ldmatrix(rA, sA[k % S]);
  ldmatrix(rB, sB[k % S]);
  mma.sync(rD, rA, rB, rD);
}

// epilogue
cp.async.wait_all;
bar.sync 0;

1 mma.sync 의미 ★

2 Ampere 기본 형태

mma.sync.aligned.m16n8k16.row.col
  .f32.f16.f16.f32
  {%f0,%f1,%f2,%f3},   // D (f32 × 4)
  {%r0,%r1,%r2,%r3},   // A (f16 × 8)
  {%r4,%r5},           // B (f16 × 4)
  {%f4,%f5,%f6,%f7};   // C (f32 × 4)

A100 throughput: 1 mma / cycle / SM (FP16 m16n8k16)

3 shape × dtype 완전 표 ★ Ampere sm_80

.row.col: A row-major, B col-major (standard). 다른 조합은 데이터 재배치 필요.

4 Fragment Layout m16n8k16

A: 16 × 16 (row-major, f16)
  thread t 의 소유:
    row = (t / 4)          (0..7)
    col = (t % 4) * 2 + j  (j=0,1)
    plus  row+8 행의 같은 col
  → t당 8 × f16 = 4 × .b32

B: 16 × 8 (col-major, f16)
  row = (t / 4)
  col = (t % 4) * 2 + j
  → t당 4 × f16 = 2 × .b32

D: 16 × 8 (row-major, f32)
  row = (t / 4), col = (t%4)*2 + j
  plus row+8
  → t당 4 × f32

5 중요 규칙

• .sync = warp 내 암묵적 barrier (activemask 기준)
• .aligned = 모든 32 thread가 같은 instruction 도달 필수
• divergent 상태에서 호출 금지
• accumulator D==C 허용 (in-place 누적)
• fragment는 thread-private register; shmem 아님 (Hopper WGMMA에서 바뀜)

1 ldmatrix 정의 ★

2 기본 사용

// x4: 4 fragment 한 번에
ldmatrix.sync.aligned.x4.m8n8.shared.b16
  {%r0,%r1,%r2,%r3}, [%rs_base];

// x4 + transpose (B matrix용)
ldmatrix.sync.aligned.x4.trans
  .m8n8.shared.b16
  {%r0,%r1,%r2,%r3}, [%rs_base];

// x2 (A: 16x16 = 2 x 8x8)
ldmatrix.sync.aligned.x2.m8n8.shared.b16
  {%r0,%r1}, [%rs_base];

3 Thread → Row Address 매핑

x1 (1 fragment, 8 rows):
  thread 0~7   : row 0~7 base 제공
  thread 8~31  : 주소 무시

x2 (2 fragment, 16 rows):
  thread 0~7   : frag0 row 0~7
  thread 8~15  : frag1 row 0~7
  thread 16~31 : 무시

x4 (4 fragment, 32 rows):
  thread 0~7   : frag0
  thread 8~15  : frag1
  thread 16~23 : frag2
  thread 24~31 : frag3

4 각 thread 부담

5 mma와의 정합 ★

• ldmatrix.x4 → A(16×16 f16) 4개 fragment
• ldmatrix.x2.trans → B(16×8 f16) 2개 fragment
• 로드 결과는 그대로 mma.sync에 입력
• 별도 shuffle / permute 불필요

6 shmem Swizzle과의 결합

// XOR swizzle (8-way)
col_swz = col ^ ((row & 7) << 3);
// 쓸 때와 읽을 때 같은 swizzle
// → ldmatrix가 conflict 없이 읽음

CUTLASS의 SmemLayoutSwizzle 참고. 상세 ↗ V06 §9.

7 Ampere .x4 패턴

// A 16x16 f16 → fragment 4개
// B 16x8 f16 → fragment 2개 (trans)
ldmatrix.x4     {a0,a1,a2,a3}, [sA];
ldmatrix.x2.trans {b0,b1},     [sB];
mma.sync.m16n8k16.row.col
  .f32.f16.f16.f32
  {d0,d1,d2,d3},
  {a0,a1,a2,a3},
  {b0,b1},
  {d0,d1,d2,d3};

1 shfl.sync 4가지 ★ idx · up · down · bfly

shfl.sync.down.b32
  %r_out | %p_valid,
  %r_in, 0x10, 0x1f, 0xffffffff;
// delta=16, clamp=31, mask=full

2 인자 해석

3 vote.sync

vote.sync.ballot.b32
  %r_mask, %p_cond, 0xffffffff;
// == __ballot_sync(0xff..., cond)

4 activemask

activemask.b32 %r_mask;
// 현재 실행 중인 lane bitmap
// == __activemask()

• divergent branch 내부에서 현재 active lane 파악
• 다음 shfl.sync의 membermask로 사용
• 주의: convergence 보장 아님

5 match.sync Volta+

// 같은 값 가진 lane mask
match.any.sync.b32 %r_m, %r_v, %mask;
match.all.sync.b32 %r_m | %p,
                   %r_v, %mask;

6 membermask의 역할 ★

• 0xffffffff: 전체 32 lane
• divergent branch 내부: activemask 결과 사용
• 참여 lane에서 값 unspecified → UB
• 동기화 contract를 HW가 검증 (warning/error)

7 Warp Reduction 예

// sum reduction
for (d = 16; d > 0; d >>= 1) {
  shfl.sync.bfly.b32 %t|%p,
    %sum, %d, 0x1f, 0xffffffff;
  @%p add.f32 %sum, %sum, %t;
}
// lane 0 에 총합

5 단계 = log₂32, 총 32개 값 합산

1 Predication ★ @p / @!p

setp.lt.s32 %p1, %r1, %r2;
@%p1  add.s32 %r3, %r1, %r2;
@!%p1 sub.s32 %r3, %r1, %r2;

2 언제 predication이 유리한가

ptxas가 heuristic으로 선택. -maxrregcount·__builtin_expect 영향.

3 Branch & Jump

4 Divergence & Reconverge

// 권장 패턴 (divergent region)
@%p bra TAKEN;
    // not-taken path
    bra DONE;
TAKEN:
    // taken path
DONE:
bar.warp.sync 0xffffffff;
// reconverge 명시

5 Early Exit 패턴

// boundary guard: thread 단위 early exit
setp.ge.s32 %p, %tid, %N;
@%p exit;

// 전 lane이 exit 하면 warp 제거
// 일부만 exit → 남은 lane만 진행

• exit: thread 종료, warp 수축
• ret: 함수 복귀
• kernel entry의 ret = kernel 종료

6 Uniform Branch 선언

• bra.uni: 모든 lane이 같은 target임을 컴파일러가 보장
• HW가 warp 전체를 한 번에 점프
• unify 못 하면 실행은 맞지만 스케줄링 손실
• @p bra.uni: 같은 predicate 값을 모두 가짐 전제

7 Branch Metadata

.branchtargets L0, L1, L2;
bra.sync LB, %r_idx;
// indirect branch

1 atom 형식 ★

atom.global.add.u32 %r_old,
  [%rd_addr], %r_val;

atom.global.cas.b32 %r_old,
  [%rd_addr], %r_cmp, %r_new;

atom.shared.max.s32 %r_old,
  [%rs_addr], %r_val;

2 지원 op 매트릭스

3 atom vs red

atom.global.add.u32 %r_old,
  [%rd], %r_val;
// vs
red.global.add.u32 [%rd], %r_val;
// no dst

• global histogram: red 선호
• lock / counter 시퀀싱: atom (old 필요)

4 Scope Qualifier ★ cta · gpu · sys

scope가 좁을수록 latency / trafic ↓

5 CAS 패턴 ★

// lock-free update template
LOOP:
  ld.global.u32 %old, [%rd];
  // compute new from old
  ...
  atom.global.cas.b32 %cur,
    [%rd], %old, %new;
  setp.ne.b32 %p, %cur, %old;
@%p bra LOOP;

• CAS 실패 → 재시도
• contention ↑ → throughput ↓
• atomic FP32 add 없는 타입 구현에 사용

6 Ordering Qualifier sm_70+

atom.acq_rel.global.add.u32
  %r_old, [%rd], %r_val;

1 membar vs fence ★

fence는 memory consistency 모델 공식화(PTX 6+).

2 bar.sync / barrier

16개 named barrier: 0~15. CTA 내 역할 분담 가능.

3 의미적 구분 ★

4 Producer-Consumer 예

// Producer
st.global.u32 [buf], %data;
fence.release.gpu;
st.global.u32 [flag], 1;

// Consumer
LOOP: ld.global.u32 %f, [flag];
      setp.eq.u32 %p, %f, 0;
@%p   bra LOOP;
fence.acquire.gpu;
ld.global.u32 %data, [buf];

relaxed ld로 spin. acquire/release로 순서 보장.

5 scope 선택 결정

6 축약 매핑

• __syncthreads() ≡ bar.sync 0
• __threadfence() ≡ membar.gl ≈ fence.sc.gpu
• __threadfence_block() ≡ membar.cta
• __threadfence_system() ≡ membar.sys
• __syncwarp(mask) ≡ bar.warp.sync mask

1 PTX Memory Model ★

• default ld/st: .relaxed
• ordering 필요 시 ld.acquire · st.release 또는 fence
• scope = 가시성 범위 (cta / gpu / sys / cluster)

2 Ordering 5종

3 Scope × Ordering 완전 매트릭스 ★

scope 확장 시 latency ↑ (cta < gpu < sys). 꼭 필요한 최소 scope 선택.

4 Acquire / Release 짝

5 Sequential Consistency (sc)

• acq_rel만으로 표현 불가한 3-thread IRIW 같은 패턴
• Dekker's algorithm류 동기화에 필요

6 PTX ↔ CUDA C++ 매핑

1 축소 맥락

• DL은 ld.global.nc + 수동 interpolation으로 대체
• volumetric rendering / CT 재구성 등에는 유용
• PTX는 지원 유지, 새 기능은 없음

2 Texture PTX 명령

3 기본 사용

// 2D bilinear sample
tex.2d.v4.f32.f32
  {%f0,%f1,%f2,%f3},
  [%rd_tex, {%f_u, %f_v}];

// gather (R channel, 4 texel)
tld4.r.2d.v4.f32.f32
  {%f0,%f1,%f2,%f3},
  [%rd_tex, {%f_u,%f_v}];

4 Surface (읽기 쓰기 모두)

surface = 2D/3D array 이미지. texture와 달리 write 가능.

5 Sampler 기능

• bilinear / trilinear / anisotropic 보간 (HW)
• clamp / wrap / mirror address 모드
• normalize coord [0,1]
• automatic border / out-of-range 처리
• CUDA array에 국한 (일반 linear memory 제한)

6 Bindless Texture modern

// CUDA C: cudaTextureObject_t
tex.2d.v4.f32.f32
  {%f0,%f1,%f2,%f3},
  [%rd_texobj, {%f_u,%f_v}];
// register에 handle 보유 → 함수 전달 OK

• CUDA 5+ bindless: handle을 reg / param으로 전달
• legacy: texture reference = 전역 심볼

1 asm() 형식 ★

asm volatile(
  "mov.u32 %0, %%laneid;"
  : "=r"(lane));

2 Constraint 문자 (!) r · l · f · d · h · n

3 operand 배치

// ldmatrix 인라인
asm volatile(
  "ldmatrix.sync.aligned.x4.m8n8"
  ".shared.b16 "
  "{%0,%1,%2,%3}, [%4];"
  : "=r"(x0), "=r"(x1),
    "=r"(x2), "=r"(x3)
  : "r"(smem_addr));

• %0~%3 → output regs (순서대로)
• %4 → input reg
• PTX의 %%는 실제 % 이스케이프

4 volatile의 효과

• side-effect 있는 op (I/O, atomic) → volatile 필수
• pure compute → volatile 불필요 (컴파일러 최적화 허용)
• "memory" clobber: 메모리 부작용 있음을 알림

5 shmem 주소 변환

// generic → .shared smem_addr
uint32_t smem_addr;
asm("cvta.to.shared.u64 %0, %1;"
  : "=l"(tmp64) : "l"(ptr));
asm("cvt.u32.u64 %0, %1;"
  : "=r"(smem_addr) : "l"(tmp64));

ldmatrix / cp.async에 shared addr은 32-bit 필요.

6 mma 인라인 ★

asm volatile(
  "mma.sync.aligned.m16n8k16"
  ".row.col.f32.f16.f16.f32 "
  "{%0,%1,%2,%3}, "
  "{%4,%5,%6,%7}, "
  "{%8,%9}, "
  "{%10,%11,%12,%13};"
  : "=f"(d0), "=f"(d1),
    "=f"(d2), "=f"(d3)
  : "r"(a0), "r"(a1),
    "r"(a2), "r"(a3),
    "r"(b0), "r"(b1),
    "f"(c0), "f"(c1),
    "f"(c2), "f"(c3));

f16 fragment는 .b32에 2개 pack → "r" constraint 사용.

7 cp.async 인라인

asm volatile(
  "cp.async.cg.shared.global "
  "[%0], [%1], 16;"
  :: "r"(smem_addr),
     "l"(gmem_ptr));
asm volatile(
  "cp.async.commit_group;");
asm volatile(
  "cp.async.wait_group %0;"
  :: "n"(STAGES-2));

wait_group의 count는 immediate (n constraint).

1 nvcc 옵션 ★ ptx · keep · src-in-ptx · Xptxas

2 전형적 flow

# .cu → .ptx 만 보기
nvcc -O3 -arch=sm_80 \
     -src-in-ptx -lineinfo \
     -ptx kernel.cu -o kernel.ptx

# reg/smem 사용량
nvcc -arch=sm_80 -Xptxas -v \
     -c kernel.cu

# 전 중간 파일 유지
nvcc -arch=sm_80 -keep kernel.cu

3 -Xptxas -v 해석

// 예시 출력
ptxas info    : Compiling entry function
  '_Z6kerneli' for 'sm_80'
ptxas info    : Function properties for _Z6kerneli
  0 bytes stack frame,
  0 bytes spill stores,
  0 bytes spill loads
ptxas info    : Used 32 registers,
  4096 bytes smem,
  376 bytes cmem[0]

• registers: 32 → occupancy 상한 결정
• smem: shared memory 사용량
• spill stores/loads: >0이면 local mem 사용 → 감속 위험
• stack frame: device 함수 호출 스택

4 Name Mangling

kernel<float, 256>  →
_Z6kernelIfLi256EEvPKT_Pi

• c++filt -p _Z...로 역변환
• extern "C" 선언 시 mangling 제거

5 PTX 읽는 순서 루틴

1. .entry 찾기 → target kernel 확인
2. .param 선언 → argument layout
3. .reg 선언 → 사용된 register 종류/수
4. kernel body의 hot loop 위치 탐색
5. loop 내 mma / cp.async / ld.global 비율
6. spill signal: .local 언급 / st.local

6 pragma & hint의 영향

1 자주 쓰는 20 instruction ★

2 Ampere 핵심 3종 (!) async · tile · mma

3 State Space 요약

4 Type Suffix 한 줄

• .u/.s: unsigned/signed
• .b: opaque bit
• .f16/.bf16/.tf32/.f32/.f64: 부동
• .f16x2: packed half pair
• .pred: 1-bit predicate

5 Modifier 한 줄

• .sync: warp-wide barrier 포함
• .aligned: 모든 lane 동일 도달 보장
• .nc: non-coherent RO cache
• .cg: L1 bypass, L2만
• .volatile: 최적화 차단

6 Memory Model 한 줄

7 Out-of-scope (이 권에서 배제)

• SASS 전반 · opcode 대응 ↗ V04 §12~16
• Hopper TMA / WGMMA / mbarrier ↗ V04 §4~10
• Cluster / DSM ↗ V04 §2~3
• 실전 inline asm 작성 팁 / 디버깅 수법
• 특정 커널 튜닝 사례 ↗ V18
• 실측 benchmark 수치 (vendor 공식 외)

8 다음 권으로

1. Hopper PTX 추가분: ↗ V04
2. CUTLASS에서 PTX 실사용: ↗ V06
3. Attention kernel의 PTX: ↗ V07
4. FP8 / 양자화 연산 PTX: ↗ V10
5. SASS 해석: ↗ V04 §12

9 기억 포인트 요약

• PTX = virtual ISA, SASS = 실제
• Ampere 3종: cp.async · ldmatrix · mma
• memory model: scope × ordering 직교
• inline asm: constraint + volatile + %% 이스케이프
• 검증은 -ptx + -Xptxas -v