Shared memory, fragment 의 차이 이해

1. Shared memory = 그냥 byte array

shared memoy 자체는 그 어떤 행렬, layout 의미도 없음

__shared__ float smem[BLOCK_SIZE][BLOCK_SIZE];

이 형태도 사실 float 배열, row, col-major, tile layout 같은 건 메모리 레이아웃개념에 불과

단지 바이트가 연속된 저장 공간

 

2. 이전에 row/col 을 쓴 것처럼 보인 이유

이전의 GEMM/커널 생성 시 

int row = threadIdx.y;
int col = threadIdx.x;

shared[row][col] = global[A_row * K + A_col];

 GPU 는 단지 linear 메모리로 취급

row/col 개념은 임의로 만든 관습적 coordinate system

 

shared memory 자체는 row-major, col-major 도 아님

shared memory 는

• logical row 없음
• logical column 없음
• tile 개념 없음
• stride 개념 없음

모두 다 임의의 인덱스로 정의한 것

shared 가 하는 일은 

주소 = base + (row * stride + col) * sizeof(float)
bank = 주소 / 4 % 32

 

 

Fragment 에서 row-major, col-major 를 명시해야 하는 이유

단순 배열이 아닌, 행렬 타일 16, 16 의 의미 구조 자체를 갖고 있는 데이터 구조이기 때문

fragment = Tensor Core 가 이해하는 행렬 조각