GPGPU 총정리 (shared memory , global memory , texture memory) - (5)

Shared Memory Bank Conflict

기본 개념에 대한 것은 이미 알고 있다고 가정하고 . (이번 범위 참고) 

bank 와 bank conlfict 에 대해서 자세히 알아보자. 

 

 

 

shared memory 가 하나의 큰 메모리 처럼 보이지만 , 내부적으로 여러 bank 로 나뉘어져있다.

CUDA 5.0 이상에서는 shared memory 가 32개의 bank 로 구성되고 , 연속된 32-bit word 가 bank 에 매핑이된다. 

-> 이때 각 bank 는 독립적인 포트를 가져가서 동시에 read/write 가능 

 

warp 가 32thread 인데 , 이 32개 thread 가 동시에 shared memory 에 접근하기 때문에 , bank 32개를 두면 , 32thread 가 각자 다른 bank 를 건든다고하면 1cycle 에 처리가 가능해진다. ! 

 

예시:

arr[0] → bank 0

arr[1] → bank 1

arr[2] → bank 2

... arr[31] → bank 31

arr[32] → bank 0 ← 다시 순환!

arr[33] → bank 1 ...

 

 

왜 conflict 가 생기는가 ? 

 

GPU 는 warp 단위로 명령을 실행한다. ( 한 warp 는 32개의 thread ) 

한 warp 의 32개 thread 가 shared memory 를 다음과 같이 접근할 수 있는 경우의 수가 생긴다. 

 

 

32 thread가 32개 서로 다른 bank 접근	No conflict — 1 cycle
N개 thread가 같은 bank의 서로 다른 word 접근	N-way conflict — N cycle 동안 직렬화
여러 thread가 같은 bank의 같은 word를 read	Broadcast — conflict 아님, 1 cycle
여러 thread가 같은 bank의 같은 word에 write	Conflict는 아니지만, 어느 thread의 값이 쓰일지 undefined

 

주의 : read 인 경우에는 conflict 가 아니다.

-> broadcast 된다. 

 

but,

하지만 write는 다르다. 여러 thread가 같은 shared memory address에 write하면 어떤 thread의 값이 최종적으로 남을지 undefined다.

s[0] = threadIdx.x;  // 여러 thread가 같은 위치에 write: 결과 undefined

 

 

 

** 팁 

Case 3: stride = 3 (홀수)

tid:    0   1   2   3   4   5   ...
idx:    0   3   6   9  12  15   ...
bank:   0   3   6   9  12  15   ...

3과 32는 서로소(gcd=1)이므로 32개 thread가 모두 다른 bank를 침. No conflict. 일반적으로 stride가 홀수면 conflict가 없다 

 

( 마찬가지 논리로 1,3,17,33 stride 인 경우에 conflict 가 0 )

 

 

 

 

Global Memory 

기본 개념은 역시 이전에 설명했었고 필요한 부분만 설명해보면..

- global memory 는 off-chip 에 존재하고 , 모든 thread block 에서 모든 thread 가 접근 가능하다. kernel launch 가 끝나고 같은 application 안에서는 계속 유지된다. ( shared memory 는 lbock 내부에서만 볼 수 있고 , block 이 끝나면 사라짐 ) 

 

(비교) 

Shared Memory : 작다 / 빠르다 / 같은 block 내부 공유 / block 끝나면 사라짐
Global Memory : 크다 / 느리다 / 모든 block 접근 가능 / kernel 이후에도 유지 

 

 

메모리 접근을 한번 보면..

 

Thread
  ↓
Register
  ↓
L1 Cache / Shared Memory  (per SM)
  ↓
L2 Cache                  (GPU 전체 공유)
  ↓
Global Memory / VRAM

 

이런 식으로 접근한다. 특히 cache miss 가 많이 난 경우에는 수백 clock cycle 이 걸릴 수 있다. 또한 global memory 로 오가는 데이터가 항상 L2 cache 통과하고 , L1 같은 경우는 경우에 따라서 bypass 될 수 있다. ( 추후에 살펴보겠다 ) 

 

즉 thread 입장에서 global memory load는 이런 식이다.

float x = d_in[idx];

 

겉으로는 한 줄이지만 내부적으로는:

1. thread가 global load instruction 발행
2. SM이 L1 cache 확인
3. L1 miss면 L2 cache 확인
4. L2 miss면 VRAM에서 가져옴
5. 가져온 데이터가 L2/L1을 거쳐 register로 들어옴

 

그래서 global memory 접근은 계산보다 훨씬 비쌀 수 있다. ( 실제로 실험을 해봤을때도 결국 계산 문제가아니라 메모리 접근 문제가 대부분이였다. ) 

 

 

 

선언 방식에는 두가지가 있는데 

 

1. Dynamic ( 런타임 할당 ) 

- cudamalloc() -> cudaMemcpy() -> cudafree()

 

2. static ( __device__ 단독 사용 )

- host 에서 쓰려면 전용 API 가 필요하다. 

cudaMemcpyToSymbol(d_val, &h_val, sizeof(float));

cudaMemcpyFromSymbol(&h_val, d_val, sizeof(float));

 

 

중요한 차이: dynamic은 컴파일러가 크기를 모르고 런타임에 결정, static은 컴파일 타임에 박힘. 그래서 __device__ 변수의 주소는 그 자체로는 host 포인터로 못 씀 — cudaMemcpyToSymbol을 거쳐야 함.

 

 

 

 

 

Global memory 에 두가지 캐시가 존재한다. 

L1 Cache

• 위치: on-chip, 각 SM에 하나씩
• 논리적 단위: 128-byte cache line
• sector 구조: 한 cache line은 32-byte sector 4개로 나뉨 (128 = 32 × 4)
• Shared memory와 같은 물리 메모리를 분할해서 씀

L2 Cache

• 위치: on-chip이지만 SM 바깥에 있고, device 전체가 공유
• 물리적 transaction 단위: VRAM과의 통신은 무조건 32-byte sector 단위
• 역할: 여러 SM에서 같은 주소 요청이 오면 합쳐서 VRAM 부담을 줄임 (aggregator)

 

핵심: VRAM ↔ L2는 32-byte sector 단위.

SM ↔ L1은 128-byte cache line 단위(4 sectors).

그래서 "한 번에 받아 와도 손해 안 보는 단위"가 32 byte (=float 8개), 128 byte (=float 32개).

 

 

 

 

 

(Warp 가 global memory load 를 요청시 ) 

 

1. Warp가 load 명령 발행
        ↓
2. L1 check (on-chip, SM 내부)
   Hit  → register로 ~20-30 cycles
   Miss → 32-byte sector 요청을 crossbar로 보냄
        ↓
3. L2 check (on-chip, device 공용)
   Hit  → L1으로 ~150-200 cycles
   Miss → memory controller 호출
        ↓
4. VRAM fetch
   요청된 32-byte sector만 GDDR/HBM에서 가져옴 (~400-800+ cycles)
   L2에 채워짐
        ↓
5. L1 fill: L2의 32-byte sector가 L1의 해당 sector로 옮겨짐

 

data가 global memory와 오가는 길은 반드시 L2를 거침. L1은 선택적으로 bypass 가능(streaming data 같은 경우). 즉 L2는 모든 global memory access의 관문.

 

 

Coalesced Global Memory Access 

 

 

Coalesced Memory Access 란 

Warp의 32 thread가 동시에 global memory에 접근할 때, 그 요청을 가능한 한 적은 32-byte sector 요청으로 묶어 처리하는 것이다. 

 

32byte 단위 ? 

VRAM (hbm ) 은 burst transfer 용으로 설계된다. 즉 , 한 byte 만 달라 라는 요청이 불가능하고 한 번 주소르 보내면 반드시 일정 크기를 한번에 받아온다 

여기서 해당 단위가 32byte 인 것이다. 

- L2<-> VRAM : 32byte 단위로 통신 

 

(cf. 32byte = float 8개 ,int 8개 ) - 이게 GPU 메모리 시스템의 최소 단위.

 

 

Warp의 메모리 요청이 실제로 처리되는 과정

다음 코드가 있다고 하자:
cudafloat val = d_in[some_index];
이걸 warp(32 thread)가 동시에 실행하면:
Step 1: 32개의 주소가 모임
        thread 0: addr_0
        thread 1: addr_1
        ...
        thread 31: addr_31

Step 2: 하드웨어가 32개 주소를 보고
        "이 주소들은 어떤 32-byte sector에 속하지?"
        를 계산함

Step 3: 필요한 sector들을 L1/L2에서 요청
        (이미 캐시에 있으면 hit, 없으면 miss → VRAM)

Step 4: 받은 sector 데이터에서
        각 thread가 자기 몫(4 byte)을 골라감
핵심은 Step 2. 32 thread의 주소가 같은 sector에 많이 묶이면 sector 요청 수가 줄어. 흩어져 있으면 늘어남.

 

다음과 같은 step 을 통해서 접근을 하게 된다. 

 

 

강의자료의 Best Case 를 먼저 살펴보자. 

이 경우는 그냥 연속 접근을 한 경우이다. thread 32개 x 4(float 가정) => 128 byte 이므로 , L1 cache miss 상황에서 128-byte cache line 하나를 할당하고 , 네개의 32-byte sector 를 L2 에서 요청하면 된다 !! 

 

 

 

최약의 경우 stride 가 8 ( 32byte)  여서 thread 별로 하나씩 접근하려고 할때 다른 sector 를 요구하는 경우에 

32thread 접근시 32 sector 가 필요하고 -> 1024 byte fetch 해야한다 ( 128 byte 사용 ) 

 

 

 "warp 내 thread들이 반드시 연속된 메모리에 접근해야 coalesced"라고 외우는데 이는 틀린말이다. 

 

진짜 조건: warp 내 32 thread의 주소가 같은 32-byte sector들 안에서 흩어져 있기만 하면 됨. 순서 섞여도 OK. 

-> 어찌됐든 , 내가 가지고온 메모리 안에서 처리할 수 있으면 순서와 관계없이 ok 인것이다. 

 

 

 

 

이제 global memory coalescing 예시를 살펴보자. 

 

 

1. GPU_0  

- 표준 패턴이다 , 연속된 메모리에 접근하는 패턴으로 보면된다. 

다만 이 패턴안에서도 신경써야할 부분이 있다. 

 

blockdim 64 ~ 1024 에서는 좋은 성능을 보여주고 거의 대부분 일정하다 

이는 thread block 크기 자체는 큰 영향이 없는 것을 보여준다  ( memory - bound 가 대부분의 영향을 미침 ) 

 

다만 , blockdim 32 ~ 1 에서는 급격히 느려지는 것을 확인할 수 있다. 

 

왜그럴까 ? 여러가지 이유가 있을 수 있겠으나 . 

warp 하나가 32 thread 인데 block 에 16 thread 밖에 없으면 warp 의 절반을 사용하지 않게 된다.  

또 blockdim 을 작게잡으면 thread block 이 매우 많아지는데 , sm 당 배정할 수있는 최대 Thread block 의 최대 개수가 정해져 있기 때문에 모든 데이터가 SM 에 올라가지 않을 수 있다. 

 

2.GPU_1 

 

GPU_1 Kernel 을 보면 인덱스의 수식이 바뀌어져 있다. 

같은 warp 에서 인덱스에 접근할때 

 

thread 0 → i = 0 * 65536 + 5 = 5

thread 1 → i = 1 * 65536 + 5 = 65541 ← 65536 떨어짐! 

 

위처럼 매우 많이 떨어져서 접근하는 것을 확인할 수 있다. 

 

 

위 gpu_0 번과 비교해보면 성능이 매우 떨어지는걸 확인할 수 있다. 

 

 

 

 

그림으로 살펴보면 다음과 같다. 

 

 

 

2. GPU_2 

 

gpu_2 는 2D block 의 모습이다.

 

2D 규칙 

: warp 는 linear ID 순으로 32 thread 씩 묶인다. ( linear ID = threadIdx.y * blockdim.x + threadidx.x ) 

 

 

blockdim 이 (32,32) 인 경우를 살펴보자. 

warp 0 = thread (0,0) (1,0) ,,, (31,0) 

warp 1 = thread (0,1) (1,1) ,,, (31,1) 

이처럼 같은 row 의 column 들이 연속으로 접근하게 된다. 이렇게되면 locality 가 올라가니까 자연스럽게 속도가 빨라지게 된다. 

 

반면 (8,128) 인 경우에는 하나의 warp 에서 줄바꿈(?) 을 4번 하게 되므로 locality 가 떨어질 수 밖에 없다. 

이중에서 최악은 (1,1024) 인 경우인데 , 각 thread 마다 새로운 메모리에 접근하게 된다. 

 

 

 

 

 

 

 

 

Constant Memory 

Constant memory 란 ? 

- Device(vram) 메모리 안에 있는 Read-only 영역으로 , 모든 thread 가 grid 전체에서 공유하며 , warp 내 32 thread 가 같은 주소를 읽을때 매우 빠른 broadcast 를 지원하는 메모리이다. 

- 별도의 constant cachce 가 각 SM 에 따로 있다. 

- 총 용량 : 64k ( gpu 다 합쳐서 ) 

( 주의 : cache 가 SM 에 있는거지 memory 자체는 off-chip 에 존재한다 ) 

 

언제쓰이나 ? 

- constant cache 는 L1 data cache 와 별도이다. 따라서 L1 을 bypass 해서 접근할 수  있게 된다. 

 

 

 

핵심 특성 : Broadcast 

일반 global memory와의 차이

Global memory에서 warp 32 thread가 같은 주소를 읽으면:

• L1/L2를 통해 한 sector를 가져옴 → 32 thread에 분배
• 그래도 sector 단위 fetch라서 32 byte 전부 가져옴

Constant memory에서 warp 32 thread가 같은 주소를 읽으면:

• 단일 read가 broadcast로 32 thread 모두에게 동시에 전달됨
• 1 cycle만에 끝 (cache hit 시)
• 즉 register access만큼 빠름

 

 

선언 방식 : __constant__ 

 

__constant__ 변수를 파일 전역에 선언해야한다. 

 

특성 : 

 

• Lifetime: CUDA context 수명 (application 끝날 때까지)
• Scope: grid 내 모든 thread + host에서도 접근 가능
• Distinct object per device: device마다 별도 인스턴스
• No constant address: 컴파일러가 알아서 주소 잡음 — 우리는 심볼 이름으로만 접근

 

 

추가로

static allocation ( 크기가 컴파일 시에 정해져야한다 ) 

 

__constant__ float cGaussian[64];     // ✓ OK
__constant__ float cArr[];            // x  dynamic 안 됨

 

 

Host에서 값을 설정하는 방법

__constant__ 변수는 직접 cudaMemcpy로 못 씀. 전용 API를 써야 함:

 

cudaGetSymbolAddress(),

cudaGetSymbolSize(),

cudaMemcpyToSymbol(),

and cudaMemcpyFromSymbol()

를 사용해야한다. 

 

 

코드

 

__global__ : ( host 에서 호출하고 device 에서 실행 )

__global__ void kernel(float* d_arr) {       // 인자로 받음
    int i = threadIdx.x;
    float v = d_arr[i];                       // 그냥 사용
}

int main() {
    float* d_arr;
    cudaMalloc(&d_arr, N * sizeof(float));                          // 1. 메모리 잡고
    cudaMemcpy(d_arr, h_arr, N*sizeof(float), cudaMemcpyHostToDevice); // 2. 값 복사
    
    kernel<<<grid, block>>>(d_arr);            // 3. 포인터 인자로 전달
    
    cudaFree(d_arr);                            // 4. 해제
}

여기서 device 메모리를 동적으로 cudamalloc 으로 잡고 , 포인터를 커널 인자로 넘긴다. 

 

 

Constant memory 쓰는 흐름

__constant__ float c_kernel[64];               // ← 파일 스코프 전역 선언!

__global__ void kernel(float* d_arr) {
    int i = threadIdx.x;
    float v = d_arr[i] * c_kernel[3];          // ← 인자로 안 받고 그냥 씀
}

int main() {
    float h_kernel[64] = { ... };
    cudaMemcpyToSymbol(c_kernel, h_kernel, sizeof(float)*64);  // ← 심볼에 직접 복사
    
    kernel<<<grid, block>>>(d_arr);            // ← c_kernel 안 넘김
}

 

__constant__ 를 전역번수처럼 선언 -> 컴파일러가 device 의 constant 영역에 잡고 

커널 인자로 안넘긴다 ( 동적으로 할당 불가능 ) 

Host 에서 값 채울시에 cudaMemcpyToSymbol로 심볼이름을 직접 주는 형식으로 

 

 

Texture and Surface Memory

 

원래 GPU 는 그래픽스 처리를 위해 만들어졌다 ( 지금은 AI 에 많이 활용되고 있지만. ) 

기본적으로 그래픽스에서는 texture image 를 자주 읽는다. 이때는 1D 배열 접근 보다 2D locality 가 중요하다 ( 이미지니까 ) 

-> 매번 Polygon 에 텍스처를 입힐떄 많은 작업을 해야하는데 , 이걸 SW 에서 처리하면 너무 비싸진다. 따라서 gpu 에서 이걸 hw 로 처리하는 전용 회로를 CUDA 가 컴퓨팅에서 쓸 수 있게 노출한것 ( = 이게 texture memory ) 

 

그래서 global memory 를 읽는 기존 경로와 다르다고 보면된다. 

 

 

texture / surface 를 이해하기 앞서서 device memory 의 두가지 저장방식을 알아보자. 

 

1. Linear memory 

- 이건 보통 cudaMalloc() 으로 device memory 를 잡을때 패턴이다. 

- 포인터로 서로 참조가 가능한 특징이 있따. 

 

2. CUDA Arrays 

- texture fecthing 에 최적화된 memory 이다. 

- 일반 포인터 처럼 직접 접근하는 배열 x , cuda 내부에서 알아서 texture fecthing 에 맞게 최적화 되어 있다. 

(너무 자세히 다루지는 않는다 ) 

 

Texture memory 

 

Texture memory 

- 2D ,3D 이미지 데이터를 효율적으로 읽기 위해 만들어 졌다. 

- Texture 메모리는 기본적으로 Read-only 이다. 

-  API 로 따로 다뤄야 한다. 

 

 

Example of Using Texture Memory 

#include <cuda_runtime.h>
#include <cstdlib>
#include <cstring>

// ============================================================
// Device code (page 47)
// ============================================================
__global__ void transformKernel(float* output,
                                cudaTextureObject_t texObj,
                                int width, int height,
                                float theta)
{
    // Calculate normalized texture coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float u = x / (float)width;
    float v = y / (float)height;

    // Transform coordinates
    u -= 0.5f;
    v -= 0.5f;
    float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;

    // Read from texture and write to global memory
    output[y * width + x] = tex2D<float>(texObj, tu, tv);
}

// ============================================================
// Host code (page 48~50)
// ============================================================
int main()
{
    const int height = 1024;
    const int width  = 1024;
    float angle = 0.5;

    // -------- ① Allocate and set some host data --------
    float *h_data = (float *)std::malloc(sizeof(float) * width * height);
    for (int i = 0; i < height * width; ++i)
        h_data[i] = i;

    // -------- ② Allocate CUDA array in device memory --------
    cudaChannelFormatDesc channelDesc =
        cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray_t cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);

    // -------- ③ Copy host → CUDA array --------
    // Set pitch of the source (width in bytes, no padding)
    const size_t spitch = width * sizeof(float);
    cudaMemcpy2DToArray(cuArray, 0, 0, h_data, spitch,
                        width * sizeof(float), height,
                        cudaMemcpyHostToDevice);

    // -------- ④ Specify texture (Resource descriptor) --------
    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = cuArray;

    // -------- ⑤ Specify texture object parameters (Texture descriptor) --------
    struct cudaTextureDesc texDesc;
    memset(&texDesc, 0, sizeof(texDesc));
    texDesc.addressMode[0]   = cudaAddressModeWrap;
    texDesc.addressMode[1]   = cudaAddressModeWrap;
    texDesc.filterMode       = cudaFilterModeLinear;
    texDesc.readMode         = cudaReadModeElementType;
    texDesc.normalizedCoords = 1;

    // -------- ⑥ Create texture object --------
    cudaTextureObject_t texObj = 0;
    cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

    // -------- ⑦ Allocate result buffer in device memory --------
    float *output;
    cudaMalloc(&output, width * height * sizeof(float));

    // -------- ⑧ Invoke kernel --------
    dim3 threadsperBlock(16, 16);
    dim3 numBlocks((width  + threadsperBlock.x - 1) / threadsperBlock.x,
                   (height + threadsperBlock.y - 1) / threadsperBlock.y);
    transformKernel<<<numBlocks, threadsperBlock>>>(output, texObj,
                                                     width, height, angle);

    // -------- ⑨ Copy device → host --------
    cudaMemcpy(h_data, output, width * height * sizeof(float),
               cudaMemcpyDeviceToHost);

    // -------- ⑩ Cleanup --------
    cudaDestroyTextureObject(texObj);
    cudaFreeArray(cuArray);
    cudaFree(output);
    free(h_data);

    return 0;
}

 

 

전반적인 흐름

Host(CPU)             Device(GPU)
──────             ─────────────────────────────────
h_data[]   ───③───>   cuArray (CUDA array, opaque) 
                          │
                          │ resDesc(=ArrayType) + texDesc(Wrap/Linear/...)
                          ▼
                       texObj (handle)
                          │
                          │ kernel: tex2D<float>(texObj, tu, tv)
                          ▼
                       output[] (linear device memory)
                          │
h_data[] <─────⑨────── output[]

데이터는 두 번 형태가 바뀜:

1. Linear (h_data) → CUDA array (cuArray): 텍스처 fetch가 최적화되는 opaque layout
2. CUDA array → Linear (output): 커널은 texture 읽고 일반 메모리에 write (texture는 read-only)

 

<단계별 설명>

1. 데이터 준비

 

2. CUDA array 할당 

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaArray_t cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);

 

- (32 , 0 , 0 , 0  , Float ) : x = 32bit , y/z/w = 0 -> 단일 채널로 가정 ( if rgb -> 32 ,32 ,32 , float ) 

- cudamallocArray : cudamalloc 이 아니라 CUDA array 전용이다. ( 내부는 모름 ) 

 

 

3. Host -> CUDA array 복사 

- sptich : source 한행이 차지하는 byte 수 

- 일반 cudaMemcpy가 아니라 2D 전용 cudaMemcpy2DToArray — pitch가 다른 2D 영역도 다룰 수 있게. 

 

 

** 이쪽 설명 보강해두기 **

( 노트에 정리해둔것 다시 옮겨두기 )