정의를 모두에게

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..

가속화된 코드를 최적화하려는 시도가 실제로 성공적이었음을 확인할 수 있는 유일한 방법은 어플리케이션의 성능에 대한 정량적인 지표 및 정보에 대해 응용 프로그램을 프로파일링하는 것이다. nsys(Nsight Systems command line tool)은 가속화된 어플리케이션을 프로파일링하기 위한 강력한 도구로 CUDA toolkit과 함께 제공된다. nsys의 가장 기본적인 사용법은 단순히 nvcc로 컴파일된 실행 파일의 경로를 전달하는 것이다. nsys는 어플리케이션 실행을 진행한 후 어플리케이션의 GPU 활동, CUDA API 호출에 대한 요약 출력 및 Unified Memory 활동에 대한 정보를 보여준다. 어플리케이션을 가속화하거나 이미 가속화 된 어플리케이션을 최적화 할 때, 과학적으로 그리고..

애플리케이션 개발자가 GPU 가속화의 이점을 가장 쉽게 얻을 수있는 코드 부분을 신속하게 식별하고, 그 이점을 빠르게 인식하고, 활용을 시작할 수 있도록 지원하기 위해 애플리케이션에 대한 APOD (Assess, Parallelize, Optimize, Deploy) 설계주기를 소개한다. 가능한한 빠른 생산(개발) 속도를 성취해낼 수 있도록 APOD를 활용해야 한다! APOD는 Cyclical(주기적인)한 프로세스이다. 최소한의 초기 시간 투자만으로 초기 속도 향상을 달성, 테스트 및 배포 할 수 있다. 또한 어느 시점에서 최적화를 수행할 지 식별하고, 추가 속도 향상을 확인한 다음, 더 빠른 버전을 배포하여 APOD 주기를 다시 시작함으로써 생산(개발) 내 빠른 버전의 응용 프로그램을 배포할 수 있다...

Grid와 Block은 3차원까지 정의 가능하다. 다차원을 사용하는 데 있어 성능 상의 이점은 없지만, 행렬/영상 등과 같은 다차원 데이터를 처리하는 데에는 유용하다. grid나 block을 2, 3차원으로 정의하기 위해서는 dim3 type을 아래와 같이 사용한다. dim3 threads_per_block(16, 16, 1); dim3 number_of_blocks(16, 16, 1); someKernel(); 위와 같은 예시에서, kernel 내의 gridDim.x, gridDim.y, blockDim.x, blockDim.y 는 모두 16으로 같다. 2차원 행렬 곱셈을 가속화하기 위한 두 개의 예제코드를 NVIDIA DLI 예제로부터 그대로 가져왔다. #include #define N 64 __gl..

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(); // ..