본문 바로가기

C++

 (22)
타일 내부 외부 경계 지점 살펴보기 - 행렬 곱의 전혀 다른 표현 형태 등장 ( QK, softmax, PV 연산의 타일 형태로의 최적화된 표현으로 인해 ) 1. 타일 내부 연산 : u_part[d] 만들기 (tile-local, parallelish)float u_part[Dv];#pragma unrollfor (int d = 0; d 타일의 한 row i 에 대해각 lane 이 자기 column j 에 해당하는 V[j][d] 를 들고e * V 를 만들어서 u_part[d] 에 담아둠d 루프는value dimension 방향으로 벡터 곱을 만드는 것각 column 이 기여하는 16 차원 벡터 조각을 만드는 과정범위 : 타일 내부아직 streaming 아님병렬성 : column 방향은 warp lane 으로 병렬, d 는 lane 내부 unroll 2. 타일 내부 reduction : u[d] 만들기 ( tile-local, warp reduce )float..
Warp All-Reduce (max / sum) 코드 분석 - warp 내 lane 들의 동시 동작, __shfl_ ) // ---------- warp reductions ----------__inline__ __device__ float warp_allreduce_max(float v) { unsigned mask = 0xffffffffu; v = fmaxf(v, __shfl_xor_sync(mask, v, 16)); v = fmaxf(v, __shfl_xor_sync(mask, v, 8)); v = fmaxf(v, __shfl_xor_sync(mask, v, 4)); v = fmaxf(v, __shfl_xor_sync(mask, v, 2)); v = fmaxf(v, __shfl_xor_sync(mask, v, 1)); return v;}__inline__ __device__ fl..
C++ lambda device code 용 local function define - CUDA 커널 안에서의 람다, 그냥 코드 블록이 펼쳐진 형태로 사용됨 코드 내용 // --- helper lambdas: loader warp copies one tile into a given buffer --- auto load_tile_KV = [&](int t, int buf) { int col_start = t * N_TILE; // K: 16x16 half = 256 half // global K layout: K_b[k * seq_len + (col_start + n)] // shared shK[buf][k][n] contiguous with ld=N_TILE for (int idx = lane; idx 1. 문법의 확인auto load_tile_KV = [&](int t, int bu..
FlashAttention-like v5 (Warp Specialization + Ping-Pong, no cp.async) 코드 분석 이 커널은 2-warp 블록에서 warp 역할을 분리한다warp 0 = Loader : global - shared 로 K/V 타일을 미리 로드warp 1 = Compute: shared K 타일로 WMMA 수행 후, streaming softmax + PV 누적타일 루프 내부에서 __syncthreads() 를 쓰지 않고, ready[2] 플래그로 생산자/소비자 동기화를 구성한다. 1. Warp 역할 분리 / 인덱싱int tid = threadIdx.x;int lane = tid & (WARP_SIZE - 1);int warp = tid >> 5; // 0,1lane 은 warp 내부 0 ... 31warp 는 블록 내 warp ID ( 2개만 존재 )이후 조건 분기로 역할을 강제한다warp = 0 ..

티스토리툴바