Race Condition 경쟁 상태 병렬 프로그램에서 Race condition (경쟁 상태)이 흔히 발생하는데 이는 여러 개의 스레드가 하나의 공유 메모리를 동시에 읽고 쓸 때 발생한다. 덧셈과 뺄셈같이 한 번에 이루어지는 계산은 아토믹 함수를 이용하고 더 넓은 범위의 코드는 크리티컬 섹션을 사용하여 경쟁 상태를 해결한다. 이 두 방법은 손쉬운 방법이지만 불필요하게 많이 사용하게 되면 오히려 프로그램의 효율이 떨어진다. //아래 코드 실행 시 10000*512개의 스레드가 실행되지만 Race condition(경쟁 상태) 발생으로 //결과로 더 적은 수가 출력됨. __global__ void ThreadCountDataRace(int* nThreadCount) { (*nThreadCount)++; }..

CUDA 프로그램의 구조 1. 호스트에서 디바이스로 입력 데이터 전송 2. 커널 함수에서 데이터 처리 3. 결과 데이터를 디바이스에서 호스트로 전송 위 과정을 차례대로 처리하면 입력 데이터가 전송되는 시간동안 GPU 프로세서는 대기하게 된다. 만일 입력 데이터를 좀 더 작게 나누고 입력 데이터 전송이 완료된 것부터 GPU에서 계산을 하는 동시에 데이터를 전송하면 프로그램의 전체 효율이 좋아진다. 데이터 입력을 받으면서 GPU 프로세서가 함께 계산하며 동시성을 높여 주는 것이 CUDA의 스트림이다. 데이터 처리 과정을 여러 개로 나누어 데이터 입력 과정과 동시에 처리할 수 있도록 하려면 데이터의 의존성(Dependency)이 없어야 한다. 앞서 입력된 데이터의 결과가 뒤에 입력될 데이터에 사용되지 않아야 ..

프로그램의 성능을 측정하기 위해 소요 시간을 계산하는 방법 1. 소요시간 측정 CPU 타이머 활용 (NVCC(CUDA compiler)가 Windows API를 간단하게 사용할 수 있는 환경이 아니므로 비추천) GPU 타이머 활용 //Measure the execution time of a program, unit : msec #include //CUDA utility tools unsigned int timer; cutCreateTimer(&timer); //Create a timer cutStartTimer(timer); //Start ... cutStopTimer(timer); //stop double elapsed_time = cutGetTimerValue(timer); CUDA 타이머 사용 시 ..

CUDA API 아키텍처 CUDA API는 3가지 그룹으로 이루어진다. CUDA Driver API CUDA Runtime API CUDA Libraries CUDA API의 세 그룹은 서로 통신을 하면서 GPU를 제어하여 CUDA 프로그램을 동작시킨다. CUDA API가 다루는 분야는 다음과 같다. 디바이스 관리 문맥 관리 메모리 관리 코드 모듈 관리 실행 제어 텍스처 관리 OpenGL과 Direct3D 상호 운용 CUDA 드라이버 API CUDA 드라이버 API는 가장 낮은 수준에 있는 API로, GPU를 직접 제어한다. 드라이버 API로 프로그램을 구현하면 더 많은 함수 호출과 관리가 필요하다. 드라이버 API는 디바이스를 제어하는 가장 낮은 수준의 명령어이기 때문에 가장 빠르게 OpenGL이나 D..

CUDA의 메모리는 GPU의 구조에 따라 계층 구조로 되어 있음. 그래픽 카드에 있는 메모리는 레지스터, 로컬 메모리, 공유 메모리, 글로벌 메모리, 상수 메모리, 텍스처 메모리로 구성되어 있음. 개별 스레드 영역 - 레지스터, 로컬 메모리 동일 블록 내에 있는 스레드 간에 공유 - 공유 메모리 그리드를 이루는 모든 스레드 간에 공유 - 글로벌 메모리, 상수 메모리, 텍스처 메모리 액세스 속도 : 글로벌 메모리 로컬 메모리 > 공유 메모리 > 레지스터 액세스 범위 : 글로벌 메모리 > 로컬 메모리 > 공유 메모리 > 레지스터 Register 레지스터 On chip processor memory로 직접 연산을 수행하는 가장 빠른..

Thread Context Switching GPU는 대량의 레지스터를 지원하므로 스레드 교환을 위해 할당할 공간을 더 많이 가지고 있어 스레드를 많이 생성해도 부담이 적음. 동시에 처리할 수 있는 스레드의 수보다 많은 스레드를 이용하여 프로그램을 구성해도 CPU와 같이 동작하지 않음. Thread Context Switching을 위해 레지스터와 메모리 간의 스위칭 교환으롭 라생할 부하를 막도록 최대 생성 가능 스레드만 생성하고 나머지 스레드는 대기시킴. 최대 스레드 개수 만큼 처리하고 난 뒤 나머지 스레드를 다시 생성하여 스레드 컨텍스트 스위칭을 위한 부하를 방지. 이를 블록과 워프 단위로 처리. SP의 스레드 스케줄링 CUDA의 최소 단위 코어인 SP(Streaming Processor)는 한 번에..