[기술] GPU와 CUDA (8) - 공유 메모리

공유 메모리 공유 메모리 사용 방법은 크게 세 가지 케이스로 구분한다. L1 캐시: 자주 사용되는 데이터를 직접 분류, 관리하기 어려운 경우 사용자 관리 캐시 1: 개발자가 커널 내 알고리즘의 데이터 접근 패턴을 파악 후, 직접 제어 사용자 관리 캐시 2: 자주 사용하는 데이터의 전역 메모리 접근을 줄이기 위함 스레드 간 공유 메모리와 L1 캐시 활용 방법 공유 메모리 (Shared Memory) 역할 및 특징: 공유 메모리는 각 블록 내 모든 스레드가 접근할 수 있는 고속 메모리 공간이다. 스레드 간 데이터를 공유하기 위해 중요한 메커니즘으로 활용된다. 동기화의 필요성: 여러 스레드가 공유 메모리에서 데이터를 읽고 쓰는 과정에서 데이터의 일관성을 유지하기 위해 동기화가 필요하다. 이를 위해 __syncthreads() 함수가 사용되며, 이 함수는 모든 스레드가 특정 지점에 도달할 때까지 대기하게 한다. L1 캐시 역할 및 특징: L1 캐시는 GPU가 자동으로 관리하는 고속 메모리로, 최근 접근한 데이터를 저장해 이후 접근 시 메모리 접근 시간을 줄인다. 캐시 히트와 미스: 캐시 히트: 필요한 데이터가 캐시에 있을 경우, 신속하게 데이터를 가져올 수 있어 성능이 향상된다. 캐시 미스: 필요한 데이터가 캐시에 없을 경우, 메모리에서 데이터를 가져와야 하므로 레이턴시가 증가하고 성능이 저하된다. 최적화 고려사항: 공간적 및 시간적 지역성을 고려한 데이터 접근 패턴을 설계하면, 캐시 히트율을 높여 성능을 최적화할 수 있다. SIMT 구조와 메모리 접근 패턴 SIMT 구조: GPU는 SIMT (Single Instruction, Multiple Threads) 구조로, 워프(warp) 단위로 스레드들이 동시에 실행된다. 메모리 접근 최적화: 효율적인 접근: 워프 내의 스레드들이 동일한 데이터 블록에 접근할 경우, 메모리 트랜잭션 수가 최소화되어 성능이 향상된다. 비효율적 접근: 스레드들이 여러 데이터 블록에 분산되어 접근하면, 메모리 트랜잭션 수가 증가해 성능이 저하될 수 있다. 사용자 관리 캐시 (User-managed Cache) 역할 및 특징: 개발자가 공유 메모리나 레지스터를 활용해 데이터를 수동으로 관리하는 것을 의미한다. 이는 자동으로 관리되는 L1 캐시와 달리, 알고리즘의 특성에 맞게 데이터를 직접 제어하는 방식이다. 활용 예시: 예를 들어, 행렬 곱셈에서 공유 메모리를 활용하여 데이터를 재사용함으로써 메모리 대역폭을 효율적으로 사용할 수 있다. 이를 통해 커널의 성능을 크게 최적화할 수 있다. 공유 메모리 사용 예제 1: 작은 행렬 하나의 블록만 사용는 행렬 곱셉 커널은 아래와 같다. 각 스레드의 인덱스는 행렬 C에서 동일 인덱스의 원소를 연산한다. int index는 행렬 C의 원소들을 1차원 형태 인덱스로 표현한 변수이다. ...

[기술] GPU와 CUDA (7) - CUDA 기반 행렬 곱셈

CUDA 기반 행렬 곱셈 행렬 연산은 CUDA 연산에 가장 어울리는 문제이다. 따라서 행렬 곱셈을 GPU에서 수행하는 방법을 이해한다. 스레드 레이아웃 설정 대규모 행렬 곱을 위한 CUDA 프로그램을 위해 스레드 레이아웃을 먼저 결정해야 한다. 어떻게 레이아웃 기준을 잡아야 할까? 두 가지 경우를 생각할 수 있다. 데이터를 읽는 행렬 A, B 기준 결과가 저장되는 행렬 C 기준 C = AB 이 행렬 연산에서 C의 (row, col) 값 계산을 위해서는 A(row, k) B (k, col) 원소들을 불러와야 한다. ...

[기술] GPU와 CUDA (6) - CUDA 실행 모델

GPU 아키텍처 SM: 스트리밍 멀티프로세서 하나의 GPU는 SM이라는 물리적 구조를 여러 개 포함한다. SM은 여러 CUDA 코어를 가진 연산 장치다. Fermi 아키텍처는 하나의 SM에 32개의 CUDA 코어를 가지고 있다. SM에는 CUDA 코어말고 레지스터, 공유 메모리, L1 캐시 등이 포함된다. CUDA 코어 CUDA 코어는 GPU의 가장 기본이 되는 프로세싱 유닛이다. 코어 안에는 FP 연산기, INT 연산기 등이 있으며, CUDA 프로그램의 동작 단위가 스레드이므로, 스레드 1개에 CUDA 코어 1개가 할당된다. CUDA 스레드 계층과 GPU 하드웨어 간단한 요약 1 스레드 = 1 코어 32개 스레드가 모여 1개의 워프이다. 또한 워프가 모여서 스레드 블록을 이룬다. 블록 내에 스레드가 있다는 것은, 스레드 묶음인 워프로 이루어져 있다는 것이다. 그리드에서 GPU 1개의 GPU는 여러 개의 그리드를 처리할 수 있다. 1개의 그리드가 여러 개의 GPU를 동시에 사용하거나 옮겨가며 실행할 수 없다. 스레드 블록에서 SM 그리드의 스레드 블록은 그리드가 배정된 GPU 속 SM에서 처리한다. 스레드 블록을 처리하는 물리적 단위는 SM이다. 따라서 스레드 블록을 적절히 SM에 분배해야 한다. 스레드 블록들은 SM에 순차적으로 균등하게 분배되어 처리된다. (하나의 SM에 여러 개의 블록이 할당될 수 있다.) SM이 갖는 자원 양과 스레드 블록을 처리하기 위해, 필요한 자원의 양에 따라 한 SM이 동시에 처리할 수 있는 스레드 블록 수를 결정 워프 & 스레드 -> SM 속의 CUDA 코어 스레드 블록의 스레드는 워프로 묶을 수 있다. 워프는 32개 스레드로 구성되며, 스레드 각각 CUDA 코어 하나에서 처리한다. 워프는 하나의 명령어에 의해 움직인다. GPU가 SIMT 아키텍처라는 말이 나온 이유이다. CUDA 코어 그룹마다 워프 스케줄러와 명령어 전달 유닛이 1개씩 있다. 스레드의 실행 문맥 워프 내 스레드들은 하나의 명령어에 의해 움직이지만, 각 스레드는 독립적으로 처리될 수 있다. (스레드만의 실행 문맥) 실행 문맥은 작업 상태의 기록이다. GPU에서 각 스레드는 자신만의 작업 상황을 저장한다. 실제 스레드의 실행 문맥은 레지스터로 관리되며, 중요한 GPU 아키텍처 특징은 스레드 블록 내 모든 워프가 SM 내부 레지스터 파일을 나누어 사용한다. ...

[기술] GPU와 CUDA (5) - 스레드 레이아웃

스레드 레이아웃 스레드 레이아웃 결정 앞서 CUDA 커널의 레이아웃은 그리드와 블록의 형태로 결정한다고 하였다. 구체적으로 다음의 과정을 따른다. 블록 형태 결정 (즉, 스레드를 어떻게 배치할껀지 결정) 데이터의 크기 및 블록 형태에 따라 그리드 형태 결정 블록 형태는 커널의 알고리즘 특성과 GPU 환경을 고려하여야 한다. 이때 레지스터, 공유 메모리 크기 등도 고려해야 할 요소이다. 큰 벡터의 합을 연산하는 CUDA 커널 (2) 벡터 차원이 1,024보다 크면 블록을 여러 개 지정해야 한다. 하나의 블록이었다면, 각 스레드가 벡터의 첫 번째 원소부터 담당하여 연산한다. 블록이 여러 개라면 각 블록이 벡터의 서로 다른 영역을 연산해야 한다. 각 블록마다 같은 인덱스의 스레드가 존재하므로, 블록 인덱스 고려없이 스레드가 백터 원소를 담당하면, 모든 블록의 동일 인덱스 스레드는 벡터의 같은 원소에 접근한다. 따라서 모든 블록이 벡터의 같은 영역을 처리한다. 원하는 구현은 다음과 같다. ...

[기술] GPU와 CUDA (4) - CUDA 연산 구조

CUDA 스레드 계층 스레드 CUDA 스레드 계층에서 가장 작은 단위는 스레드이다. 따라서 CUDA 연산을 수행하거나, 코어를 사용하는 기본 단위이다. 커널 호출 시, CUDA 커널 코드는 모든 스레드에 공유된다. 각 스레드는 커널을 독립적으로 실행한다. 워프 CUDA 스레드 계층의 두 번째 계층은 워프이다. 워프는 32개 스레드를 하나로 묶은 단위이다. 중요한 점은 워프는 디바이스에서 하나의 제어 장치에 의해 제어된다. GPU의 SIMT 구조에서 멀티 스레드 단위가 바로 워프이다. 이 말은 1개의 명령어에 의해 32개 스레드가 동시에 움직이는 것을 의미한다. ...

[기술] GPU와 CUDA (3) - CPU와 GPU의 벡터 합 연산

벡터 합을 구하는 호스트 프로그램 #include <stdio.h> #inlcude <stdlib.h> #include <string.h> #define NUM_DATA 1024 int main(void) { int* a, * b, * c; int memSize = sizeof(int) * NUM_DATA a = new int[NUM_DATA]; memset(a, 0, memSize); b = new int[NUM_DATA]; memset(b, 0, memSize); c = new int[NUM_DATA]; memset(c, 0, memSize); for (int i = 0; i < NUM_DATA; i++) { a[i] = rand() % 10; b[i] = rand() % 10; } for (int i = 0; i < NUM_DATA; i++) { c[i] = a[i] + b[i]; } delete[] a; delete[] b; delete[] c; } 벡터 합을 구하는 디바이스 프로그램 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> #inlcude <stdlib.h> #include <string.h> #define NUM_DATA 1024 __global__ void vecAdd(int* _a, int* _b, int* _c) { int tID = threadIdx.x; _c[tID] = _a[tID] + _b[tID]; } int main(void) { int* a, * b, * c, * hc; int* da, * db, * dc; int memSize = sizeof(int) * NUM_DATA; printf("%d elements, memSize = %d byte\n", NUM_DATA, memSize); // 호스트 디바이스에 메모리 할당 a = new int[NUM_DATA]; memset(a, 0, memSize); b = new int[NUM_DATA]; memset(b, 0, memSize); c = new int[NUM_DATA]; memset(c, 0, memSize); hc = new int[NUM_DATA]; memset(hc, 0, memSize); // 데이터 값 할당 for (int i = 0; i < NUM_DATA; i++) { a[i] = rand() % 10; b[i] = rand() % 10; } // 호스트에서 벡터 합 계산 (성능 비교용) for (int i = 0; i < NUM_DATA; i++) { hc[i] = a[i] + b[i]; } // 디바이스 메모리를 할당 cudaMalloc(&da, memSize); cudaMemset(da, 0, memSize); cudaMalloc(&db, memSize); cudaMemset(db, 0, memSize); cudaMalloc(&dc, memSize); cudaMemset(dc, 0, memSIze); // 호스트에서 디바이스로 데이터 카피 cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice); cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice); // 커널 함수 호출 vecAdd <<<1, NUM_DATA>>>(da, db, dc); // 결과를 디바이스에서 호스트로 카피 cudaMemcpy(c, dc, memSize, cudaMemcpyDevcieToHost); // 디바이스 메모리 해제 cudaFree(da); cudaFree(db); cudaFree(dc); // 결과 체크 bool result = true; for (int i = 0; i < NUM_DATA; i++) { if (hc[i] != c[i]) { printf("[%d] The result is not matched! (%d, %d) \n", i, hc[i], c[i]) result = false } } if (result) printf("GPU works ~") // 호스트 메모리 해제 delete[] a;delete[] b; delete[] c; return 0; } 실행 시간 커널 실행 시간 커널 실행 시간을 측정하기 위해서는 커널 호출 전과 후에 시간을 측정해야 한다. 이를 통해 디바이스의 연산 실행 시간을 알 수 있다. 비동기성과 동기화 호스트는 커널을 호출하면 디바이스에게 커널 실행을 요청하고, 바로 다음 명령으로 넘어간다. 따라서 호스트는 디바이스와 별개로 동시에 작업이 가능하다. 따라서 커널 실행 시간을 정확히 측정하려면, 호스트가 커널 호출 후 바로 다음 작업으로 넘어가는 비동기성을 고려해야 한다. 호스트가 디바이스의 연산 수행이 끝나기까지 기다리게 하여 정확한 시간을 측정할 수 있다. 이를 동기화라고 한다. 주의사항 CUDA API 호출은 큐 구조를 통해 관리된다. 따라서 순차 실행을 한다. 호스트가 디바이스 제어만을 위한 경우, 동기화는 필요없다. 그러나 호스트, 디바이스 모두 연산에 관여하는 경우, 동기화가 필요하다. CUDA API 호출 자체에도 일정 시간의 소모 비용이 있다. 하지만 CUDA 연산의 정확성과 효율성을 모두 잡으려면, 스레드 계층과 스레드 인덱싱 방법이 필요하다. # cudaDeviceSynchronize를 통해 디바이스 연산이 끝날 때까지, 호스트를 기다리게 한다. veeAdd <<<1, NUM_DATA>>>(da, db, dc); cudaDeviceSynchronize(); 데이터 전송 CUDA 프로그램 흐름 호스트에서 디바이스로 데이터 복사 디바이스 연산 디바이스에서 호스트로 데이터 복사 데이터 복사 CUDA 알고리즘의 성능을 판단할 때는, 반드시 데이터 전송 시간을 고려해야 한다. 이 외에도 데이터 변환 등의 추가 작업이 있다면, 이의 소요 시간도 포함해야 한다.

[기술] GPU와 CUDA (2) - CPU와 GPU 통신

호스트와 디바이스 호스트 호스트는 일반적으로 CPU를 의미한다. 따라서 호스트 코드는 CPU에서 실행되는 코드를 의미한다. 또한 호스트 메모리는 CPU가 사용하는 시스템 메모리이다. (DRAM) 디바이스 디바이스는 일반적으로 GPU를 의미한다. 따라서 디바이스 코드는 GPU에서 실행되는 코드를 의미한다. 또한 디바이스 코드는 GPU가 사용하는 GPU 메모리이다. CUDA 프로그램 CUDA 프로그램은 호스트 코드와 디바이스 코드로 구성된다. 프로그램 실행 시 처음 호출되는 코드는 CPU에서 프로세스를 할당하기 때문에, 호스트 코드가 통상 같이 있다. CUDA 프로그램에서 호스트 코드는 gcc와 같은 컴파일러로, 디바이스 코드는 NVCC 컴파일러로 컴파일한다. C++에서 소스 파일은 “.cpp”, 헤더 파일은 “.h” 에 대응하면, CUDA는 소스 파일은 “.cu”, 헤더 파일은 “.cuh” 이다. ...

[기술] GPU와 CUDA (1) - GPU의 연산 개념

GPU에 관하여 GPU는 방대한 수학 연산을 가속하기 위해 설계된 전자 회로이다. GPU는 CPU에 비해 수천 개의 작은 코어(모델 및 사용 목적에 따라 다름)를 가지고 있기 때문에 GPU 아키텍처는 병렬 처리에 최적화되어 있다. GPU는 여러 작업을 동시에 처리할 수 있으며 그래픽 및 수학적 워크로드에서 더 빠르다. GPU vs CPU GPU vs CPU 기본적인 GPU 구조 Flynn’s Taxanomy 플린의 분류법은 스탠포드 대학교의 마이클 J. 플린이 컴퓨터 아키텍처를 분류한 것이다. 플린의 분류법의 기본 개념은 간단하다. 계산은 순차적으로(한 번에 하나의 스트림) 또는 병렬로(한 번에 여러 스트림) 처리할 수 있는 두 개의 스트림(데이터 및 명령어 스트림)으로 구성된다. 두 개의 데이터 스트림과 이를 처리할 수 있는 두 가지 방법은 플린의 분류법에서 4가지 범주로 이어진다. ...