커널 내부에서의 sharedmemory 선언 이유, 커널은 그리드 전체에 대응되는 느낌임!

커널이 스레드 하나에 대응되는 것이 아님,

__global__ void gemm_kernel(...) {
    // ...
}

이 함수 하나가 grid 전체 가 실행하는 코드

런치 시, blockIdx.x / threadIdx.x 값만 다를 뿐, 같은 커널 코드를 수백/수천 개 스레드가 동시에 실행하는 구조

커널 안에 적는 건 스레드 하나의 코드가 아닌, 블록/스레드 전체가 공유하는 프로그램 템플릿

 

shared memory 의미 : 블록당 1개씩 있는 on-chip scratchpad

어떤 배열을 __shared__ float buf[1024] 라고 선언하면

grid 에서 블록이 100개 있으면, buf 도 100개 생김 ( 블록마다 1개씩 ) 

같은 블록 안의 모든 스레드가 공유, 같은 blockIdx 를 가진 스레드들은 buf[i] 를 같이 보고, sync 해서 같이 씀, 블록 바깥에서는 절대접근 불가

__global__ void kernel(...) {
    __shared__ float smem[1024]; // "이 커널을 실행하는 각 블록마다 하나씩 주세요"
    ...
}

block 0  → smem[1024] 하나
block 1  → smem[1024] 또 하나
block 2  → smem[1024] 또 하나
...

 

커널 내부에 선언하는 이유는, 커널이 실행될 때 각 블록마다 필요하다는 의미

커널 밖에다가 만들면 의미가 달라짐, 대부분 커널마다 자기 전용 shared 작업 공간을 갖고 싶어함