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 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 다음
