Shared memory 최적화
__syncthreads() 과다 사용 문제
__syncthreads()는 필요하지만 비용이 크다. 불필요하게 여러 번 호출하면 성능이 떨어진다.
문제 상황
데이터를 업데이트한 후에도 __syncthreads()를 한 번 더 호출하는 경우가 있다.
__global__ void kernelAdjDiff(float* b, const float* a, unsigned num) {
__shared__ float s_data[1024];
register unsigned tx = threadIdx.x;
register unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < num) {
s_data[tx] = a[i];
}
__syncthreads();
if (tx > 0) {
b[i] = s_data[tx] - s_data[tx - 1];
} else if (i > 0) {
b[i] = s_data[tx] - a[i - 1];
} else {
b[i] = s_data[tx] - 0.0f;
}
__syncthreads(); // 이것이 필요한가?
}
마지막 __syncthreads()는 데이터 업데이트 이후에도 나올 것을 가정하고 호출한다. 하지만 계산 후에 shared memory를 더 이상 사용하지 않는다면 불필요하다.
성능 비교
adjdiff-shared.cu 587 usec (__syncthreads() 1번)
adjdiff-overused.cu 629 usec (__syncthreads() 2번, 속도 저하!)
불필요한 동기화 하나가 42 usec의 오버헤드를 만들었다. 이는 약 7%의 성능 감소다.
속도 저하가 나쁜 것이다. __syncthreads()를 해줘도 나빠질 것은 없지 않냐고 생각할 수 있다. 하지만 실제로 측정해보면 확실히 느려진다.
왜 오버헤드가 발생하는가
"그냥 기다리기만 하는데 왜 시간이 걸리나?"라고 생각할 수 있다. 하지만 barrier synchronization은 단순 대기가 아니다.
Barrier Synchronization이란?
Block 내 모든 스레드가 특정 지점에 도착할 때까지 기다리는 동기화 메커니즘이다. 마치 단체 관광에서 "다음 장소로 가기 전에 모두 모이세요"와 같다.
하드웨어가 해야 할 일
1024개 스레드를 가진 블록이 __syncthreads()를 실행한다고 하자.
1단계: 도착 체크
각 스레드가 barrier에 도착했는지 추적한다.
Warp 0 (Thread 0-31): 도착 완료
Warp 1 (Thread 32-63): 도착 완료
Warp 2 (Thread 64-95): 아직 실행 중...
...
SM(Streaming Multiprocessor)은 32개 warp의 도착 상태를 확인해야 한다. 이 과정에서 하드웨어 카운터를 업데이트하고 비교한다.
2단계: 대기 상태 유지
도착한 warp는 실행을 멈추고 대기 상태로 들어간다. Warp scheduler는 이들을 "대기 중" 목록에 넣는다.
실행 가능: Warp 2, Warp 5, Warp 10
대기 중: Warp 0, Warp 1, Warp 3, ...
Warp scheduler가 대기 중인 warp를 관리하는 비용이 발생한다.
3단계: 모두 도착 확인
마지막 스레드가 도착하면 하드웨어는 "모두 도착했다"는 것을 확인한다. 32개 warp가 모두 도착했는지 체크한다.
4단계: 재개 신호
모든 warp에게 "이제 계속하세요" 신호를 보낸다. 대기 중이던 32개 warp를 다시 "실행 가능" 상태로 변경한다.
5단계: 실행 재개
Warp scheduler가 warp들을 다시 스케줄링한다.
시간이 걸리는 이유
위 과정이 모두 하드웨어 오버헤드다.
- 도착 카운터 업데이트: 각 warp마다 atomic 연산
- 대기 상태 관리: Warp scheduler의 상태 변경
- 전체 확인: 32개 warp 상태를 모두 체크
- 재개 신호: 모든 warp에 broadcast
- 재스케줄링: 32개 warp를 다시 큐에 넣기
1GHz GPU라면 42 usec = 약 42,000 clock cycles다. 1024개 스레드를 관리하려면 이 정도 비용이 든다.
실제 영향
벤치마크 결과를 보면 알 수 있다.
587 usec (1번 동기화)
629 usec (2번 동기화)
---
차이: 42 usec (7%)
동기화 1번당 약 20 usec 정도의 오버헤드다. 커널 실행 시간이 짧을수록 비율이 커진다.
만약 커널이 10 usec만 실행된다면?
10 usec (원래 실행 시간)
+ 20 usec (__syncthreads() 1번)
= 30 usec (3배 느려짐!)
불필요한 동기화를 제거하면 성능이 크게 개선된다.
Shared Memory 동적 할당
Compile time에 shared memory 크기를 알 수 없는 경우가 있다. 이때 동적 할당을 사용한다.
왜 동적 할당이 필요한가
정적 할당의 한계를 먼저 보자.
__global__ void kernel() {
__shared__ float s_data[1024]; // 크기 고정
}
블록 크기가 512라면? 512개만 사용하고 나머지 512개는 낭비된다.
블록 크기를 사용자가 선택하게 하려면? 컴파일 시점에 크기를 모른다.
// main 함수에서 명령행 인자로 블록 크기 받음
blocksize = procArg(argv[0], argv[2], 32, 1024); // 32~1024 사이
이런 경우 커널 코드에 크기를 하드코딩할 수 없다.
선언 방법
커널에서 extern __shared__로 선언한다.
__global__ void kernelAdjDiff(float* b, const float* a, int num) {
extern __shared__ float s_data[]; // 크기 없음!
register unsigned tx = threadIdx.x;
register unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < num) {
s_data[tx] = a[i];
}
__syncthreads();
if (tx > 0) {
b[i] = s_data[tx] - s_data[tx - 1];
} else if (i > 0) {
b[i] = s_data[tx] - a[i - 1];
} else {
b[i] = s_data[tx] - 0.0f;
}
}
extern __shared__ float s_data[]의 의미:
extern: 크기가 외부에서 결정됨__shared__: shared memory에 할당float s_data[]: float 배열, 크기는 나중에 지정
컴파일러는 "이 커널이 shared memory를 쓴다"는 것만 안다. 얼마나 쓸지는 모른다.
커널 실행 시 크기 지정
커널 호출 시 세 번째 파라미터로 바이트 단위 크기를 전달한다.
unsigned blocksize = 1024;
int main(const int argc, const char* argv[]) {
blocksize = procArg(argv[0], argv[2], 32, 1024); // 런타임에 결정
dim3 dimBlock(blocksize, 1, 1);
dim3 dimGrid(div_up(num, dimBlock.x), 1, 1);
kernelAdjDiff<<<dimGrid, dimBlock, blocksize * sizeof(float)>>>(dev_vecB, dev_vecA, num);
// ^^^^^^^^^^^^^^^^^^^^^^^^^
// 세 번째 파라미터!
cudaDeviceSynchronize();
}
세 번째 파라미터 상세 설명
<<<dimGrid, dimBlock, sizeInByte>>>
- 첫 번째: Grid 구성 (블록 개수)
- 두 번째: Block 구성 (스레드 개수)
- 세 번째: Shared memory 크기 (바이트 단위)
blocksize * sizeof(float)의 의미:
blocksize = 1024 (스레드 개수)
sizeof(float) = 4 (float는 4바이트)
-----------------------------------
1024 * 4 = 4096 bytes (4KB)
스레드마다 float 하나씩 저장하려면 4KB가 필요하다.
만약 blocksize가 512라면?
512 * 4 = 2048 bytes (2KB)
필요한 만큼만 할당한다.
Compile Time vs Runtime
정적 할당 (Compile Time):
__shared__ float s_data[1024]; // 컴파일 시 크기 결정
- 장점: 빠름
- 단점: 유연성 없음
동적 할당 (Runtime):
extern __shared__ float s_data[];
<<<grid, block, 1024*sizeof(float)>>>(...)
- 장점: 유연함 (실행 시 크기 결정)
- 단점: 약간 느림 (overhead)
동적 할당의 약간 느린 점 때문에 성능이 미세하게 떨어질 수 있다.
성능 비교
adjdiff-shared.cu 587 usec
adjdiff-shared2.cu 594 usec (dynamic allocation)
Dynamic allocation을 사용해도 7 usec 정도만 차이가 난다. 이는 약 1%의 차이로 무시할 수 있는 수준이다.
Device Query로 Shared Memory 크기 확인
GPU마다 사용 가능한 shared memory 크기가 다르다. Device query로 확인하여 적절한 크기를 설정할 수 있다.
주요 함수
cudaGetDeviceCount
__host__ __device__ cudaError_t cudaGetDeviceCount(int* count);
compute-capable device의 개수를 반환한다. count는 사용 가능한 device 개수다. count = 0이면 CUDA를 사용할 수 없다. count = 1은 보통의 경우다. count = 2 이상이면 여러 개의 CUDA-capable device가 있다.
cudaDriverGetVersion과 cudaRuntimeGetVersion
__host__ cudaError_t cudaDriverGetVersion(int* driverVersion);
__host__ __device__ cudaError_t cudaRuntimeGetVersion(int* runtimeVersion);
driver와 runtime 버전을 반환한다.
cudaGetDeviceProperties
__host__ cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop, int device);
compute-device에 관한 정보를 반환한다. prop은 지정된 device의 속성이다. device는 속성을 가져올 device 번호다 (0이 기본 CUDA device).
cudaDeviceProp 구조체는 다음을 포함한다.
struct cudaDeviceProp {
char name[256];
cudaUUID_t uuid;
size_t totalGlobalMem;
size_t sharedMemPerBlock;
int maxThreadsPerBlock;
int regsPerBlock;
...
};
sharedMemPerBlock이 블록당 shared memory 크기다. maxThreadsPerBlock이 블록당 최대 스레드 수다.
사용 예제
int main(const int argc, const char* argv[]) {
int deviceCount = 0;
cudaGetDeviceCount(&deviceCount);
cudaGetDeviceProperties(&deviceProp, 0); // 0은 기본 CUDA device
cudaDriverGetVersion(&driverVersion);
cudaRuntimeGetVersion(&runtimeVersion);
// device query 후 설정
blocksize = deviceProp.maxThreadsPerBlock;
kernelAdjDiff<<<dimGrid, dimBlock, blocksize * sizeof(float)>>>(dev_vecB, dev_vecA, num);
}
deviceProp.maxThreadsPerBlock을 사용하여 블록 크기를 자동으로 설정한다.
실행 결과는 다음과 같다.
CUDA device "GeForce RTX 2070": driver ver 11.2, runtime ver 11.3
max num of threads per block = 1024
max dim size of a thread block (x,y,z) = (1024, 1024, 64)
이 GPU는 블록당 48KB shared memory와 1024개의 스레드를 지원한다. 스레드당 몇 바이트를 사용할 수 있는가? 48KB / 1024 threads = 48 bytes per thread다.
성능 결과
adjdiff-shared.cu 587 usec
adjdiff-shared2.cu 594 usec (dynamic alloc)
adjdiff-devQuery.cu 595 usec (dev query)
Device query를 추가해도 거의 차이가 없다. 595 usec로 여전히 빠르다.
모든 버전 성능 비교
Intel Core i5-3570과 GeForce RTX 2070으로 테스트한 결과다.
host version: 31,810 usec
CUDA global mem: 473 usec (67배 향상)
CUDA shared mem: 587 usec
CUDA dynamic allocation: 594 usec
CUDA sync twice: 629 usec
CUDA dev query: 595 usec
Global memory를 사용하는 것이 가장 빠르다. 이 예제에서는 각 데이터를 한 번만 읽기 때문이다. Shared memory로 복사하는 오버헤드가 더 크다.
하지만 복잡한 경우에는 shared memory를 쓰면 확실히 빠르다. 최신 CUDA device들은 L1/L2 cache가 추가되고 성능이 매우 좋다. Device마다 테스트가 필요하다.
2년 전 구형 아키텍처에서는 결과가 달랐다.
host version: 65973.269601 msec
CUDA global mem: 69.164193 msec
CUDA shared mem: 47.299384 msec (가장 빠름!)
CUDA dynamic allocation: 51.761590 msec
CUDA sync twice: 61.578443 msec
구형 GPU에서는 shared memory가 확실히 빠르다. 47 msec vs 69 msec로 약 1.5배 차이가 난다.
교훈은 다음과 같다. Global memory를 사용하는 쪽이 더 빠른 경우도 있지만, 상황에 따라서는 shared memory를 쓰면 확실히 빠르다. __syncthreads() 함수를 과다 사용하지 말아야 한다.
Kernel에서 Pointer 사용
Shared memory를 활용하다 보면 포인터를 사용해야 하는 경우가 생긴다. 동적 할당된 extern __shared__ 배열도 결국 포인터로 접근하고, 복잡한 자료구조를 shared memory에 구성하려면 포인터가 필수다.
하지만 CUDA에서 포인터 사용은 일반 C/C++와 다른 규칙을 따른다. Qualifier의 의미가 다르고, 같은 warp 내에서 포인터를 잘못 사용하면 system crash까지 발생할 수 있다. Shared memory를 안전하고 효율적으로 사용하려면 포인터 동작 방식을 정확히 이해해야 한다.
Qualifier란?
변수나 함수 선언 앞에 붙는 키워드로, 메모리 위치나 속성을 지정한다. CUDA에서는 shared, device, constant, global 등이 있다. 예를 들어 shared int x는 "x가 shared memory에 위치한다"는 의미다. C/C++ 표준의 const, static도 qualifier다.
기본 사용법
__device__ int my_global_variable;
__constant__ int my_constant_variable = 13;
__global__ void foo(void) {
__shared__ int my_shared_variable;
int* ptr_to_global = &my_global_variable;
const int* ptr_to_constant = &my_constant_variable;
int* ptr_to_shared = &my_shared_variable;
*ptr_to_global = *ptr_to_shared;
}
위 코드에서 ptr_to_global, ptr_to_constant, ptr_to_shared는 모두 레지스터 변수다. 각 포인터가 가리키는 메모리 영역이 다르다.
Pointer의 타입
Pointer는 memory space에 typed되지 않는다. 이게 무슨 뜻인지 구체적으로 살펴보자.
__shared__ int* ptr;
이 선언에서 ptr이 가리키는 곳은 어디인가?
많은 사람들이 "shared memory를 가리키는 포인터"라고 생각한다. 하지만 틀렸다.
정답은 "ptr 자체가 shared memory에 위치한 포인터 변수"다.
헷갈리는 이유
C/C++에서는 이렇게 생각한다.
const int* ptr; // "const int를 가리키는 포인터"
qualifier가 "무엇을 가리키는지"를 나타낸다고 생각한다.
하지만 CUDA는 다르다.
__shared__ int* ptr; // "shared memory에 저장된 포인터 변수"
qualifier가 "포인터 자체가 어디에 위치하는지"를 나타낸다.
메모리 레이아웃으로 이해하기
구체적인 예를 보자.
__global__ void foo() {
__shared__ int shared_var; // (1)
__shared__ int* ptr; // (2)
ptr = &shared_var; // (3)
}
메모리 맵은 아래와 같다.
Shared Memory (블록당 48KB)
[shared_var] <- 4 bytes (int)
[ptr] <- 8 bytes (포인터, 64bit 시스템)
Global Memory
(없음)
(1) shared_var는 shared memory에 있는 int 변수다.
(2) ptr는 shared memory에 있는 포인터 변수다.
(3) ptr에 shared_var의 주소를 저장한다.
ptr 변수 자체가 shared memory에 위치한다. 하지만 ptr이 가리키는 대상은 어디든 될 수 있다.
다양한 조합
포인터 변수의 위치와 가리키는 대상은 독립적이다.
__device__ int global_var;
__shared__ int shared_var;
__global__ void foo() {
// Case 1: shared에 있는 포인터가 global을 가리킴
__shared__ int* ptr1;
ptr1 = &global_var;
// Case 2: shared에 있는 포인터가 shared를 가리킴
__shared__ int* ptr2;
ptr2 = &shared_var;
// Case 3: register에 있는 포인터가 shared를 가리킴
int* ptr3; // qualifier 없음 = register
ptr3 = &shared_var;
}
메모리 맵은 아래와 같다.
Shared Memory
[shared_var] <- int
[ptr1] <- 포인터 변수 (global_var 주소 저장)
[ptr2] <- 포인터 변수 (shared_var 주소 저장)
Register (스레드별)
[ptr3] <- 포인터 변수 (shared_var 주소 저장)
Global Memory
[global_var] <- int
ptr1, ptr2는 shared memory에 위치한 포인터 변수다. 하지만 각각 다른 메모리를 가리킨다.
Pointer 사용 시 주의점
같은 warp 내에서 pointer가 다른 memory 영역을 access하면 CUDA system crash 또는 성능 저하가 발생한다.
__device__ int my_global_variable;
__global__ void foo(int *input) {
__shared__ int my_shared_variable;
int* ptr = nullptr;
if (input[threadIdx.x] % 2) {
ptr = &my_global_variable;
} else {
ptr = &my_shared_variable;
}
// where does ptr point?
}
문제가 발생하는 이유
위 코드를 구체적으로 분석해보자.
Warp는 32개 스레드가 동시에 실행되는 단위다. Warp 0번을 예로 들면 Thread 0~31이다.
Thread 0: input[0] % 2 = 0 → ptr = &my_shared_variable
Thread 1: input[1] % 2 = 1 → ptr = &my_global_variable
Thread 2: input[2] % 2 = 0 → ptr = &my_shared_variable
Thread 3: input[3] % 2 = 1 → ptr = &my_global_variable
...
같은 warp 내에서 Thread 0과 Thread 2는 shared memory를 가리키고, Thread 1과 Thread 3는 global memory를 가리킨다.
GPU의 메모리 접근 메커니즘
GPU는 warp 단위로 메모리 접근을 처리한다. 모든 스레드가 같은 종류의 메모리를 접근한다고 가정한다.
- 모두 global memory 접근 → 하나의 global memory 트랜잭션
- 모두 shared memory 접근 → 하나의 shared memory 트랜잭션
하지만 섞여 있으면?
- 성능 저하: GPU가 여러 번 나누어 접근해야 한다
- 먼저 global memory 접근하는 스레드들 처리
- 그 다음 shared memory 접근하는 스레드들 처리
- 원래 1번이면 될 일을 2번 이상 해야 함
- System crash: 하드웨어가 혼란스러워서 예측 불가능한 동작
- 어떤 GPU는 잘못된 메모리 주소를 읽을 수 있음
- 프로그램이 죽거나 이상한 값이 나올 수 있음
Warp Divergence와의 차이
Warp divergence (if문으로 스레드가 다른 경로 실행)는 느려지지만 안전하다.
if (threadIdx.x % 2) {
a = b + c; // 짝수 스레드
} else {
d = e + f; // 홀수 스레드
}
이건 괜찮다. 느리긴 하지만 crash는 안 난다.
하지만 메모리 종류가 섞이면 훨씬 심각하다. Crash까지 날 수 있다.
안전한 예시
모든 스레드가 같은 메모리 종류를 가리키면 안전하다.
__shared__ int shared_array[1024];
__global__ void safe_example(int* input) {
int* ptr = &shared_array[threadIdx.x]; // 모두 shared memory
*ptr = input[threadIdx.x];
}
모든 스레드가 shared memory의 다른 위치를 가리킨다. 안전하다.
Pointer 사용 권장사항
Simple, regular access patterns로 pointer 사용을 제한한다. 단순하고 규칙적인 접근 패턴에서만 포인터를 사용한다.
Pointers to pointers는 피한다. GPU는 pointer chase를 잘 처리하지 못한다. Linked lists는 성능이 나쁘다.
Compiler warning에 주의한다. 컴파일러는 pointer가 어디를 가리키는지 알 수 없어서 global memory space로 가정한다. Warning: Cannot tell what pointer points to, assuming global memory space가 나오면 나중에 system crash 가능성이 있다.
Kernel Function Parameter의 Memory Space
Kernel function 호출 시 파라미터 처리 방식을 이해해야 한다.
기본적으로 Call-by-Value
Kernel function call은 call-by-value다. struct를 넘기면 그대로 copy된다.
Pointer는 CUDA Global Memory로 가정
Pointer 파라미터는 모두 CUDA global memory space로 가정한다.
__global__ void foo(int *input) {
// input은 global memory를 가리킨다고 가정
}
위 커널은 input이 device의 global memory에 있다고 가정한다.
배열 파라미터의 함정
int dim[3] 처럼 배열을 넘길 때 주의가 필요하다.
C/C++ 배열-포인터 Decay
C/C++에서 함수 파라미터로 배열을 전달하면 자동으로 포인터로 변환된다.
Array-to-Pointer Decay란?
함수 인자로 배열을 전달할 때, 배열이 자동으로 포인터로 변환되는 C/C++의 규칙이다. int arr[10]을 전달하면 int* arr이 된다.
// 이 두 선언은 동일하다
void foo(int dim[3]);
void foo(int* dim);
컴파일러는 두 번째 형태로 변환한다. 배열 크기 정보는 사라진다.
CUDA에서의 문제
CUDA kernel도 같은 규칙을 따른다.
__global__ void kernel(int dim[3]) {
// 실제로는 int* dim으로 처리됨
int x = dim[0];
}
컴파일러가 보는 것은 int* dim이다. 그리고 CUDA는 모든 포인터를 global memory로 가정한다.
잘못된 사용: Host 배열 전달
다음 코드는 작동하지 않는다.
int main() {
int host_dim[3] = {10, 20, 30}; // Host 메모리
kernel<<<1, 1>>>(host_dim); // 위험!
cudaDeviceSynchronize();
}
실행 과정:
- Host에서
host_dim배열 선언 (CPU 메모리) - Kernel 호출 시 배열이 포인터로 decay
- GPU가 받는 것: CPU 메모리의 주소
0x7fff1234abcd - Kernel에서
dim[0]접근 시도 - GPU가 device global memory의 주소
0x7fff1234abcd를 찾으려 함 - 그런 주소 없음 → Segmentation fault 또는 이상한 값
Host와 Device의 주소 공간 분리
CPU와 GPU는 완전히 독립된 메모리 공간을 가진다.
Host Memory (CPU):
주소: 0x7fff1234abcd
[host_dim[0]] = 10
[host_dim[1]] = 20
[host_dim[2]] = 30
Device Memory (GPU):
주소: 0x7fff1234abcd <- 이 주소는 GPU에 존재하지 않음!
Host의 주소 0x7fff1234abcd를 GPU가 접근하면 실패한다. GPU는 자신의 메모리 공간에서만 접근할 수 있다.
마치 다른 컴퓨터에 있는 메모리 주소를 읽으려는 것과 같다.
올바른 사용: Device 메모리 할당
Device 메모리를 명시적으로 할당하고 복사해야 한다.
int main() {
// Host 배열
int host_dim[3] = {10, 20, 30};
// Device 메모리 할당
int* dev_dim = nullptr;
cudaMalloc(&dev_dim, 3 * sizeof(int));
// Host → Device 복사
cudaMemcpy(dev_dim, host_dim, 3 * sizeof(int), cudaMemcpyHostToDevice);
// Kernel 실행 (Device 포인터 전달)
kernel<<<1, 1>>>(dev_dim); // 안전!
cudaDeviceSynchronize();
// 정리
cudaFree(dev_dim);
}
실행 과정은 아래와 같다.
cudaMalloc으로 GPU 메모리 할당 → 주소0xd000012340반환cudaMemcpy로 데이터 복사 → GPU 메모리에{10, 20, 30}저장- Kernel에 GPU 주소
0xd000012340전달 - Kernel에서
dim[0]접근 → GPU 메모리의0xd000012340읽기 → 성공!
Host Memory:
[host_dim] = {10, 20, 30}
Device Memory:
주소: 0xd000012340
[dev_dim[0]] = 10 <- cudaMemcpy로 복사됨
[dev_dim[1]] = 20
[dev_dim[2]] = 30
GPU가 자신의 메모리 공간에 있는 주소를 접근하므로 안전하다.
핵심 규칙
Host 배열을 직접 전달하면 안 된다.
int arr[10];
kernel<<<1, 1>>>(arr); // 잘못됨
Device 메모리를 할당하고 복사해야 한다.
int* dev_arr;
cudaMalloc(&dev_arr, 10 * sizeof(int));
cudaMemcpy(dev_arr, arr, 10 * sizeof(int), cudaMemcpyHostToDevice);
kernel<<<1, 1>>>(dev_arr); // 올바름
cudaFree(dev_arr);
배열 파라미터는 자동으로 포인터로 변환되고, CUDA는 그 포인터가 global memory를 가리킨다고 가정한다. Host 메모리 주소를 전달하면 GPU가 접근할 수 없다.
관련 문서
- CUDA Runtime API - cudaGetDeviceProperties - 디바이스 속성 조회
- CUDA C Programming Guide - Dynamic Shared Memory - 동적 Shared Memory 할당
- CUDA C Programming Guide - Pointer Restrictions - 포인터 사용 제약사항
- CUDA C Best Practices Guide - Shared Memory - Shared Memory 최적화
- Bank Conflicts in Shared Memory - Shared memory 성능 최적화를 더 공부하려면, Bank Conflict 주제도 참고해보면 좋다.