CUDA Event 2019. 5. 2 CUDA Event → Stream 내부 명령에 흐름 중, 특정 지점에 표시를 남기는 것(일종의 표시 → marker) Creation and Destruction cudaError_t cudaEventCreate (cudaEvent_t* event); cudaError_t cudaEventDestory (cudaEvent_t event); Recording Events cudaError_t cudaEventRecord (cudaEvent_t event, cudaStream_t stream = 0); stream에 event 발생 위치를 표시함 stream에서 꺼내질 때 정보가 event에 기록됨 Synchronization / Querying cudaError_t cudaEventSynchronize .. CUDA Stream & Concurrent Execution 2019. 5. 2 CUDA Stream → Host에서 Device로 명령을 보내는 통로(Host에서 호출하는 명령들이 차례대로 들어감) Types of CUDA Stream NULL stream → 암묵적으로 선언된 stream 사용할 stream을 명시하지 않은 경우(default stream) Non-NULL stream → 명시적으로 선언된 stream 명시적으로 생성 및 사용 하나의 stream에 들어온 명령은 순서대로 Device에 의해 처리된다. → Synchronous execution C() → B() → A() 하지만 서로 다른 stream에 있는 명령들은 순서가 정해져 있지 않다.(비동기적) 즉, stream들이 명시되어 있으면 비동기적인 stream이 될 수 있다. Concurrent Executio.. Synchronization in CUDA 2019. 5. 2 CUDA에서 Synchronization 하는 방법은 총 3가지가 있다. Synchronization functions Atomic functions Manual control Synchronization Function __synchthreads() Intra-block synchronization(Block 내 모든 thread가 도달) __syncwarp Inter-warp synchronization(Warp 내 모든 thread가 도달) Device code에서 사용 Kernel or device functions Atomic Functions → 여러 스레드가 들어와도 한 스레드가 한 번의 연산을 하는 것(한 번에 처리)을 보장하는 function On 32-bit or 64-bit words.. Maximizing Memory Throughput 2019. 5. 1 Memory Throughput을 높이는 방법은 Memory Access Pattern과 관련이 있다. 하나의 warp는 32개의 thread로 구성되어 있다. 따라서 동시에 32개의 요청이 발생할 수 있다. 이 요청의 방식에 따라 성능이 크게 달라질 수 있다. 예를들어 4개의 스레드가 요청이 왔을 때 위 모양보다 아래 모양에서 더 효율적인 성능이 나올 수 있을 것이다. Global Memory를 살펴보자면 Global Memory의 transcation은 L2 Cache를 통해 이루어지는데 L2 Cache는 최대 32-byte memory를 활용할 수 있다. (최근에는 L1 Cache를 활용하는 것도 있다. - 128-byte) 그렇다면 Global Memory의 접근 방법을 살펴보자. Aligned .. CUDA Memory Model 2019. 5. 7 CUDA Memory Hierarchy Per-Thread (local Memory) Registers 가장 빠르고, 가장 작은 메모리 보통 블록 1개가 8k-64k 32bit register를 사용한다. 각 SM 내부에 존재 thread들이 register를 나누어서 가져간다.(register의 영역을 나누어서 사용) 따라서 context switching이 필요 없다. Local Memory Off-chip memory SM 내부에 없고, DRAM 영역 내부에 존재 register에 다 올라가지 못하는 것들은 Local Memory에 올려서 동작 Register에 비해 느리지만 더 많은 공간을 활용할 수 있다.(크게 제약이 없다.) Per-Block (shared Memory) Shared Memor.. CUDA Execution Model 2019. 5. 7 ceil() → 올림 해주는 함수 이 함수를 사용해 필요한 Grid의 개수를 설정해줄 수 있다. 2D Grid with 2D Block의 index를 찾아갈 때 1차원적으로 배열을 생각하고 찾아갈 수도 있지만 2차원적으로 생각해서 Row와 Col을 각각 구하고 이를 통해 index를 찾아가는 2차원적 방법도 생각해볼 수 있다. 만약 위 그림과 같은 구조에서 index를 찾아간다고 하면 Row(iy)와 Col(ix)를 각각 구한 다음 iy만큼 nCol을 곱하고 ix만큼 더해주면 원하는 index가 된다. 이를 코드로 옮긴다면 iy, ix를 구하는 방법은 다음과 같고 index를 구하는 방법은 결국 다음과 같다. GPU Architecture (Fermi) CUDA core Basic processing u.. CUDA Thread 2019. 4. 2 기존 C에서 "Hello World"를 찍기 위한 코드는 다음과 같다. 1 2 3 4 int main(void) { printf("Hello World!\n"); return 0; } cs CUDA에서 "Hello World"를 찍기 위한 코드는 다음과 같다. CUDA를 컴파일하기 위해선 NVIDIA compiler(nvcc)가 필요하다. 1 2 3 4 5 6 7 8 __global__ void mykernel(void){ // empty kernel } int main(void) { mykernel (); printf("Hello CUDA!\n"); return 0; } cs CUDA의 기본적인 코드들을 살펴보자. __host__ → Host에서 호출 가능한 code(기본값) __device__ → D.. 이전 1 다음
