정의를 모두에게

from Steve Rennich, NVIDIA

arxiv.org/pdf/1910.07787v2.pdf github.com/ProfHubert/NAMF ProfHubert/NAMF This project is for our paper "NAMF: A NON-LOCAL ADAPTIVE MEAN FILTER FOR SALT-AND-PEPPER NOISE REMOVAL" - ProfHubert/NAMF github.com

CUDA 프로그래밍에서 Stream은 순서대로 실행되는 일련의 명령체계이다. CUDA 어플리케이션에서 Kernel 실행과 일부 메모리 전송은 CUDA Stream 내에서 발생한다. 명시적으로 CUDA Stream에 대해 작성하지 않았을 때에는 실제로 CUDA Code는 default stream이라고 불리는 stream 내에서 kernel이 실행된다. CUDA 프로그래밍에서는 default stream 외에 non-default CUDA stream을 생성하고 활용할 수 있는데, 이를 활용하면 여러 kernel을 동시에 여러 개의 다른 stream에서 실행하는 등의 여러 작업을 수행할 수 있다. 다중 stream을 사용하면 가속화된 어플리케이션에 병렬화 계층을 추가할 수 있고 어플리케이션 최적화를 위한 ..

Host to Device 혹은 Device to Host 메모리 전송 모두의 경우에서 page fault (페이지 폴트)와 on-demand memory migration (메모리 마이그레이션)의 오버헤드를 줄이는기술을 Asynchronous Memory Prefetching (비동기 메모리 프리페칭) 이라고 부른다. 이 기술을 사용하면 프로그래머가 어플리케이션 코드에서 사용하기 전 백그라운드에서 Unified Memory (UM)를 시스템 내의 CPU 혹은 GPU 장치로 비동기적으로 마이그레이션하는 것이 가능하다. 이렇게 하면, page fault 및 on-demand 데이터 마이그레이션 오버헤드가 줄어들어 GPU Kernel과 CPU function 성능이 향상될 수 있다. Prefetching은 ..

지금까지 cudaMallocManaged를 사용해 Host 혹은 Device 에서 사용할 메모리를 할당해왔다. cudaMallocManaged가 실제로 어떻게 동작하는지, Unified Memory(UM) 할당 방법에 대한 세부 정보를 모르는 상태에서도 자동 메모리 마이그레이션, 프로그래밍 용이성 등의 이점을 얻었다. nsys 프로필은 가속화된 어플리케이션의 UM 관리에 대한 세부 정보 또한 제공하고, 이 정보를 UM 작동 방식에 대한 보다 자세한 이해와 함께 사용한다면 가속화된 어플리케이션을 최적화할 수 있는 추가적인 기회를 얻을 수 있다. (이렇게 쓰고 보니 게임 같네..) 아래 슬라이드는 향후 설명할 내용을 시각적으로 보여주니 참고하자. UM이 할당되면, 메모리는 host나 device에 아직 상주..

Streaming Multiprocessors and Warps GPU 하드웨어의 특징을 이해하여 최적화를 촉진하는 방법을 살펴보자. Streaming Multiprocessors(SMs)를 알고 나면, 지금까진 작업해온 응용 프로그램을 더욱 최적화할 수 있다. CUDA 어플리케이션이 실행되는 GPU들에는 Streaming Multiprocessor라는 처리 장치가 있는데, Kernel을 실행하는 동안 thread들을 구성하는 block들이 SM에 전달되어 실행하도록 구성된다. 여기에서 GPU가 가능한한 많은 병렬 처리를 수행하도록 하기 위해, 주어진 GPU 장치에서 SM 수의 배수가 되는 블록의 수를 grid 크기로서 선택함으로써 성능을 향상시킬 수 있다. 또한, SM은 warps라고 불리는 block..

cudaError_t err; err = cudaMallocManaged(&a, N) // Assume the existence of `a` and `N`. if (err != cudaSuccess) // `cudaSuccess` is provided by CUDA. { printf("Error: %s\n", cudaGetErrorString(err)); // `cudaGetErrorString` is provided by CUDA. } 위 코드처럼 cudaError_t type을 통해 오류를 확인하는 것이 가능한데, 우리가 종종 만들게될 kernel의 경우 return type이 void 이므로 위와 같이 확인이 불가능하다. 그 때에는 아래 코드와 같이 오류를 확인한다. someKernel(); // ..

CUDA 함수의 실행 환경(Execution Configuration)이, 루프를 병렬화하기 위해 필요한 thread의 수와 정확히 일치하지 않는 경우가 있다. 최적의 block size들을 선택하는 데 있어서는 GPU 하드웨어 특성 상 thread의 32의 배수로 설정하는 것이 성능 향상에 있어 일반적이다. 일례로, 1000개의 병렬 작업을 수행한다고 가정할 경우 1000은 32의 배수로 나누어 떨어지지 않기 때문에 block size를 선택하는 데 어려움이 있다. 이와 같은 경우는 다음과 같은 방법으로 쉽게 해결할 수 있다. 할당된 병렬 작업을 수행하는 데 필요한 수(N)보다 많은 thread를 생성하는 실행 환경을 작성한다. Kernel에 파라미터를 추가하여 N을 인수로서 전달한다. Grid 내에서..

CUDA에서 CPU와 GPU 장치에서 메모리를 보다 쉽게 이용하기 위해 CUDA는 다양한 메모리 할당 기법을 지원한다. 가장 기본부터 살펴보자면, Host(CPU)와 Device(GPU)에서 참조될 포인터에 대하여 할당(allocate) 및 해제(free)하기 위해서, malloc과 free 를 cudaMallocManaged와 cudaFree로 대체할 수 있다. int N = 2

기존에 CPU 전용 응용 프로그램으로 작성된 For loop 코드는 CUDA를 이용해 가속화하기에 적합한 예이다. Loop의 반복적인 직렬 Process를 자체 thread에서 병렬로 Processing하는 것이 가능하다. Loop가 실행되는 횟수를 제어하고, 각 반복에 대해 발생될 일을 고려하고 수정해야 한다. void loop(int N) { for (int idx = 0; idx < N; ++idx) printf("%d\n", idx); } void main() { int TOTAL_ITER = 10; loop(TOTAL_ITER); } __global__ void loop() { printf("%d\n", threadIdx.x); } void main() { int TOTAL_ITER = 10; ..