728x90

이전에 커널 함수에서 Kernel<<<Block, Thread>>> 이라고 잠깐 설명을 했었다.


여기서 Grid, Block, Thread에 대해 개념을 어느 정도 잡고 있어야 할 것 같아서 정리한다.

우선 한줄 요약을 하면 다음과 같다. "Thread가 모여서 Block이 되고, Block이 모여서 Grid가 된다."

즉, 'Thread -> Warp -> Block -> Grid' 라고 생각하면 된다.


그림 - 동작 단위



커널 함수에서 Kernel<<<N, 1>>> 은 N개의 블록이 각각 1개의 스레드를 가지고 있다는 뜻이다.

Kernel<<<1, N>>> 은 1개의 블록이 N개의 스레드를 가지고 있다는 뜻이다.


앞에서 blockIdx로 블록의 인덱스를 이용해서 데이터를 처리했었다. 

마찬가지로 threadIdx로 스레드의 인데스를 이용해서 데이터를 처리할 수 있다.

blockDim은 블록당 스레드의 수를 나타내고, gridDim은 그리드당 블록의 수를 나타낸다.


위에 있는 Warp에 대해 설명한다. CUDA 아키텍처에서의 워프는 "서로 엮여 있으면서" 정확히 동일한 방식으로 발맞추어 실행되는 32개의 스레드들의 집단과 관계가 있다. 프로그램 코드의 매 라인마다 하나의 워프 내의 각 스레드는 다른 데이터를 가지고 동일한 명령어를 수행한다.




728x90
728x90

10개의 원소를 갖는 두 벡터 a, b의 합을 그림으로 보면 다음과 같다.



그림 - 벡터의 합


이것을 우리가 흔히 사용했던 C 코드로 작성을 하면 다음과 같을 것이다.


그림 - 벡터의 합 C 코드



그림 - 벡터의 합 결과



이것을 이제 CUDA를 이용하여 병렬 프로그래밍을 해보도록 하자.


그림 - CUDA를 이용한 벡터의 합 Main


이전 C코드에서 Main부분은 바뀐 부분이 GPU에서 사용되는 변수가 추가된 것과 add() 함수가 커널 함수로 바뀐 것 밖에 없다.

커널 함수 add<<<N, 1>>>의 첫 번째 인자를 통해 N개의 병렬 블록(Block)들이 실행되도록 지시하고 뒤에 1은 스레드가 1개가 돈다는 뜻이다. 즉 add<<<Block 개수, Thread 개수>>> 이다. 

Block과 Thread는 다음 글에서 다시 설명하도록 하겠다.

이 블록(Block)이라는 것은 쉽게 말해 각각의 병렬 실행을 뜻한다. 즉 N개의 add()가 생성되어 병렬로 실행된다는 것이다.

블록의 개수 N은 65,535의 값을 초과할 수 없다.(하드웨어의 제약사항)



그림 - add 커널 함수


add() 함수는 tid 값에 처음 보는 blockIdx.x가 대입되었다. 앞에서 커널 함수 add()를 호출할 때 블록 인자를 넣어주었다.

blockIdx가 바로 디바이스 코드를 작동하는 블록의 인덱스 값이다. 그리고 'x, y'가 있지만, 여기서는 'x'만 이용하였을 뿐이다.(y는 어차피 1이니...)

그리고 밑에서 while 문이 if 문으로 바뀌었다. 이것은 블록별로 병렬 실행되므로 while로 계산할 필요없이, 해당하는 인덱스의 데이터만 처리하면 되기 때문이다.


각 스레드들의 blockIdx.x는 0에서 N - 1사이의 값을 가질 것이다. 따라서 변수 blockIdx.x가 서로 다른 값을 가지면서 동일한 디바이스 코드의 복사본들을 수행하는 네 개의 블록들을 가정해볼 수 있다.


그림 - GPU를 이용한 벡터의 합



그리고 이러한 병렬 블록들의 집합을 그리드(grid)라고 부른다. 



728x90
728x90

DirectX11을 사용할 때 반드시 이해해야 하는 두 가지 주된 인터페이스로 장치(Device)장치 문맥(Device Context)가 있다.


쉽게 말하면 파이프라인이 사용할 자원들을 생성할 때는 장치(Device) 인터페이스가 사용되고, 생성된 자원들을 실제로 사용하거나, 파이프라인 자체를 조작할 때에는 장치 문맥(Device Context) 인터페이스가 사용된다.


DirectX11의 장치 인터페이스(ID3D11Device)는 셰이더 프로그램 객체, 자원, 상태 객체, 질의 객체 등의 생성을 위한 여러 메서드들을 제공한다. 또한 여러 하드웨어 기능들의 사용 가능 여부를 점검하는 메서드들과 진단 및 디버깅 관련 메서드들도 제공한다.


장치 문맥(ID3D11DeviceContext)은 장치로 생성한 자원이나 상태 객체를 파이프라인에 묶는 데 쓰인다. 또한 장치 문맥은 렌더링 파이프라인과 계산 파이프라인의 실행을 제어하는 수단을 제공하며, 장치로 생성한 자원을 조작하는 수단도 제공한다. 응용 프로그램은 인터페이스(ID3DDeviceContext)를 통해서 장치 문맥에 접근한다. DirectX11은 다중 스레드 렌더링을 좀 더 잘 지원하기 위해서 두 종류의 장치 문맥을 제공하는데, 하나는 즉시 문맥(Immediate Context)과 또 하나는 지연 문맥(Deferred Context)이다.


즉시 문맥(Immediate Context)은 파이프라인에 직접 연결되는 통로라고 생각할 수 있다. 이 문맥에서 어떤 메서드를 호출하면 그 호출은 DirectX 11 실행시점 모듈에 '즉시' 제출되어서 해당 명령이 구동기에서 실행된다. 하나의 응용 프로그램은 단 하나의 즉시 문맥만 사용할 수 있으며, 이 즉시 문맥은 장치가 생성될 때 함께 만들어진다. 이 문맥을 파이프라인의 모든 구성요소와 직접 상호작용하기 위한 인터페이스라고 생각하면 된다. 이 문맥은 반드시 주 렌더링 스레드에서, GPU에 대한 1차적인 인터페이스로서 사용해야 한다.


지연 문맥(Deferred Context)는 2차적인 장치 문맥으로 주 렌더링 스레드 이외에 2차적인 스레드들이 보낸 일련의 명령들을 스레드에 안전하게 기록하는 메커니즘을 제공한다. 이 문맥으로 소위 '명령 목록(Command List)' 객체를 생성하고, 나중에 주 스레드에서 그 명령 목록을 '재생'하는 것이 가능하다. 또한 DirectX 11은 여러 스레드들에서 자원들을 비동기적으로 생성할 수 있다.


그림 - Immediate Context와 Deferred Context의 관계


728x90
728x90

컴퓨터 프로그램을 사용하여 Model 또는 Scene으로 부터 영상을 만들어 내는 과정을 Rendering이라고 하고, 한 데이터 처리 단계의 출력이 다음 단계의 입력으로 이어지는 형태로 연결된 구조를 Pipeline이라고 한다.


그림 - Pipeline


즉, Rendering Pipeline 이라는 것은 메모리 자원들을 GPU로 처리해서 하나의 렌더링된 이미지를 만들어 내는 데 쓰이는 메커니즘을 뜻한다. 


DirectX11의 Rendering Pipeline을 대략적으로 보면 다음과 같다.

그림 - 대략적인 Rendering Pipeline


대략적인 Rendering Pipeline의 그림을 보면 우선 Rendering 과정이 초록색과 빨간색으로 나뉘어져 있다.

초록색은 Fixed Function(고정 기능 단계)로 미리 정해진 특정한 연산들만 수행이 가능하고, 실행 시점에서 기능은 바꿀 수 없지만, 상태 객체라는 개념을 이용하여 설정들의 변경이 가능하며, 이 Fixed Function의 단계는 임의로 실행을 하지 않는 방법은 없다.

빨간색은 Programmable(프로그래밍 가능 단계)로 HLSL(High Level Shading Language, 고수준 셰이딩 언어)로 셰이더 프로그래밍이 가능한 단계이다. 물론 특정 Programmable 단계를 비활성화 하는 것도 가능하다.


Rendering Pipeline의 진입점은 입력 조립기(Input Assembler)이다. 이 단계는 자원들로 부터 입력 자료를 읽어 들여서 파이프라인의 이후 단계들이 사용할 정점들을 긁어 모으는 작업을 한다. 또한 입력 자원들에 기초해서 정점들 사이의 연결성을 파악하고 바람직한 렌더링 구성을 결정한다. 이 단계는 취합한 정점들과 기본도형(Primitive) 연결성 정보를 다음 단계인 정점 셰이더로 넘겨준다.


다음 단계인 정점 셰이더(Vertex Shader) 단계는 입력 조립기가 넘겨준 정점 자료의 정점들을 한 번에 하나 씩 처리한다. 각 입력 정점마다 현재 지정된 정점 셰이더 프로그램이 적용된다. 이 정점 셰이더에서 하는 일은 저의 경험으로 비추어 볼때 대체로 변환 행렬을 적용하거나, 정점 별 조명계산 수행 등 Pixel Shader에서의 불필요한 연산들을 미리 할 때 사용한다.


다음 세 단계는 하드웨어의 테셀레이션(Tessellation) 기능을 활용하기 위해서 최근에 파이프라인에 추가된 것으로 반드시 세 단계 (Hull Shader, Tessellatior, Domain Shader)을 함께 사용하거나, 사용하지 않아야 한다.


테셀레이션(Tessellation)의 첫 번째 단계인 덮개 셰이더(Hull Shader)는 정점 셰이더가 넘겨준 기본 도형들을 받아서 두 가지 작업을 수행한다. 첫 작업은 각 기본 도형마다 실행되는 것으로, 여기서 덮개 셰이더는 기본도형에 대한 일단의 테셀레이션 계수들을 결정한다. 이 계수들은 이후 과정에서 해당 기본 도형을 얼마나 세밀하게 분할해야 하는지를 파악하는데 쓰인다. 두 번째 작업은 바람직한 출력 제어 패치 구성의 각 제어점마다 실행되는 것으로, 여기서 덮개 셰이더는 이후 영역 셰이더 단계(Domain Shader)에서 기본도형을 실제로 분할하는 데 사용할 제어 점들을 만들어낸다.


테셀레이션의 두 번째 단계인 테셀레이터(Tessellator)는 현재 기본 도형 종류에 적합한 표본 추출 패턴(sampling pattern)을 결정한다. 테셀레이터 단계는 테셀레이션 계수들과 자신만의 구성을 이용해서, 현재 기본 도형의 정점들 중 기본 도형을 더 작은 조각으로 분할하기 위한 표본으로 사용할 정점들을 결정한다. 그 정점들로부터 산출한 일단의 무게 중심 좌표들을 다음 단계인 영역 셰이더(Domain Shader)에게 넘긴다.


테셀레이션의 세 번째 단계인 영역 셰이더(Domain Shader)는 무게 중심 좌표(barycentric coordinates)들과 덮개 셰이더가 생성한 제어점들을 입력으로 받아서 새 정점들을 생성한다. 이 단계에서 현재 기본 도형에 대해 생성된 제어점들 전체와 텍스처, 절차적 알고리즘 등을 이용해서, 테셀레이션된 각 점마다 무게중심 '위치'들을 출력 기하구조(geometry)로 변환해서 다음 단계로 넘겨준다. 테셀레이션 단계에서 증폭된 정점 자료들로부터 출력 기하구조를 생성하는 부분의 유연성 덕분에 파이프라인 안에서 좀 더 다양한 테셀레이션 알고리즘을 구현할 수 있는 여지가 생긴다.


다음 단계인 기하 셰이더(Geometry Shader)는 완성된 형태의 기본 도형(다각형)들을 처리하거나 생성한다. 이 단계에서는 파이프라인에 새로운 자료 요소를 추가하거나 제거할 수 있으며, 덕분에 전통적인 렌더링 파이프라인에서는 불가능했던 흥미로운 응용이 가능해진다.


이제 래스터화기 단계(Rasterizer Stage)에서는 주어진 기하 구조가 렌더 대상의 어떤 픽셀들을 덮는지 파악해서 그 픽셀들에 대한 단편 자료를 산출한다. 각 단편은 모든 정점별(per-vertex) 특성을 해당 픽셀 위치에 맞게 보간한 값들을 가진다. 이렇게 생성된 단편은 픽셀 셰이더로 넘어간다.


픽셀 셰이더(Pixel Shader)는 연결된 각 렌더 대상을 위한 색상 값을 출력한다.


마지막 단계인 출력 병합기(Output Merger)는 픽셀 셰이더 출력을 파이프라인에 연결된 깊이/스텐실 자원 및 렌더 대상 자원에 제대로 '합치는' 작업을 담당한다. 이 과정에서 출력 병합기는 깊이 판정과 스텐실 판정, 혼합 함수 적용 등을 수행하며, 최종적으로는 출력을 해당 자원에 실제로 기록한다.



또 다른 셰이더가 한개 더 존재한다. 그것은 계산 셰이더(Compute Shader)이다. 이 단계는 통상적인 렌더링 패러다임에서 벗어나는 계산을 수행하기 위해 만들어진 것으로, 그런 만큼 통상적인 렌더링 파이프라인과 개별적으로 수행된다고 생각하면 된다. 계산 셰이더 단계는 전적으로 범용 계산에만 쓰이는 하나의 독립적인 파이프라인이라고 할 수 있다. 이전에는 셰이더 프로그램의 호출(실행)이 특정 단계에서 입력이 처리되는 방식에 제약을 받았다. 그리고 스레드가 쓰이는 방식을 개발자가 직접 제어하는 것이 불가능했다. 하지만 계산 셰이더 단계에서는 더 이상 그렇지 않다. 또 다른 기능은 '그룹 공유 메모리' 블록이다. 한 스레드 그룹에 속하는 모든 스레드는 그 그룹의 공유 메모리에 접근할 수 있다. 실행 도중 스레드들은 이 메모리 블록을 통해서 서로 자료를 주고받을 수 있는데, 전통적인 렌더링 파이프라인에서는 이러한 소통이 불가능했다. 그런 소통이 가능해지면, 적재한 자료나 중간 계산 결과를 공유함으로써 효율성을 더욱 개선할 여지가 생긴다.


마지막으로 계산 셰이더 단계에서는 자원들에 대한 임의 읽기 접근과 쓰기 접근이 동시에 가능하다.



728x90
728x90

CUDA에서는 CPU를 호스트, GPU는 디바이스라고 부른다.

그리고 호스트 코드는 CPU에서 수행될 코드를 의미하고, 디바이스 코드는 CUDA가 가능한 그래픽 칩셋에서 수행될 코드를 의미한다.


__global__ : 커널(Kernel) 함수임을 컴파일러에게 명시

__device__ :  디바이스(GPU)에서만 사용하는 함수라고 명시 (생략이 가능함)

__host__호스트(CPU)에서 실행되는 함수라고 명시

커널(Kernel) 함수 : 디바이스에서 실행되는 함수이며, 호스트에서만 호출 됨

<<<A,B>>> : CUDA 런타임 시스템에 넘겨질 매개변수


아래 예제는 단순하게 A + B = C 예제이다.



cudaMalloc() : 디바이스 메모리(전역 메모리)를 할당

cudaMemcpy() : 호스트와 디바이스간의 메모리 복사

 - cudaMemcpyDeviceToHost

 - cudaMemcpyHostToDevice

 - cudaMemcpyDeviceToDevice

cudaFree() : 디바이스 메모리 해제


* cudaMalloc()으로 할당한 메모리 포인터를 커널로 전달할 수 있다.

* 커널에서 cudaMalloc()으로 할당한 메모리 포인터를 이용하여 메모리를 읽거나 쓸 수 있다.

* cudaMalloc()으로 할당한 메모리 포인터를 호스트 함수로 전달할 수 있다.

* 호스트 함수에서 cudaMalloc()으로 할당한 메모리 포인터를 읽거나 쓸 수 없다.



728x90
728x90

object라는 클래스가 있다는 가정하에 헤더파일에 선언할 수 있는 4가지 타입의 전역변수는 아래와 같습니다.

1. object var;
2. static object var;
3. exturn object var;
4. typedef Singleton <object> var;

1번과 같은 경우 2개 이상의 cpp파일에서 해당 헤더파일을 포함하는 경우, 변수가 미리 정의되어 있다는 에러를 발생시킵니다. 즉, 1개의 cpp파일을 위해서만 정의할 수 있는, 단순히 cpp파일에 선언한 것을 헤더파일로 옮긴거 외에는 의미를 가지지 않습니다.

뭐랄까? 1개의 cpp파일을 위한 전역변수가 되겠죠.

2번과 같은 경우 2개 이상의 cpp파일에서 해당 변수를 사용할 수 있습니다. 왜냐하면 헤더파일을 포함한 cpp파일의 개수만큼 객체가 생성되기 때문입니다. 전역변수기는 한데 각각의 cpp파일마다 각각의 전역변수가 생성됩니다. 뭐, 좀 더 정확하게 표현하자면 헤더파일을 포함하는 각각의 cpp파일에서만 사용하는 로컬변수라고 부를 수 있겠죠. 용어에서 헷갈리는 이유는 변수의 생존 범위인데, 쥔장의 개념에서 지역변수의 개념은 함수안에서만 존재하는 변수를 지역변수라고 생각했기 때문입니다. 그리고는 단순하게 로컬변수가 아닌 변수는 전역변수라고 생각했기 때문에... (흠... 지금도 헷갈림...켁~)

3번과 같은 경우는 cout, cin과 같은 우리가 흔희 보는 STL의 전역변수들이 사용하는 기법입니다. 진정으로 여러 cpp파일에서 공유하는 한개의 전역변수가 생성됩니다. 즉, C/C++에서 전역변수라고 불리는 것은 3번과 같은 경우가 될 것입니다.

4번과 같은 경우는 Design Pattern의 하나로서 1개의 객체만을 생성하게 됩니다. 전역변수의 개념이나 3번과 같이 compiler에 의해서가 아닌 테크닉에 의하여 생성되게 됩니다. 실제적으로는 클래스의 typedef이며 var::Instance() 와 같이 클래스의 static함수를 사용하여 1개의 전역변수를 사용하게 됩니다.

패턴주의자라면 4번을 선호할 것이며 C에서 부터 C++로 넘어오신 분이면 3번을 선호할 것 입니다.

초급을 넘어서면 심심할 때 고민할 이 둘의 차이에 대해서 말하자면 구글형이 가르켜준 Singleton VS global object 라는 토론에 잘 나와 있습니다. 인상깊었던 내용의 일부를 이야기하자면 표면적으로 비슷하지만 Singleton 쪽이 좀 깔끔하게 정의할 수 있고 (extern은 특정 cpp파일에서 실제 객체의 정의해야 하며, 전역객체가 많아지면 특정 객체의 정의 위치를 찾는 것도 한 일 되겠죠?) 그 외에 config파일 사용, 늦은 초기화 등의 테크닉을 사용할 수 있다는 장점이 있습니다. 개인적으로 그 외의 차이점은 전역변수 같은 경우는 전역변수를 다루는 메모리에 들어가는데 비하여 Singleton은 생성된 객체가 힙메모리에 위치하는 차이점이 있는듯 보입니다.


- 출처 : http://linkmemo.tistory.com/entry/C-%ED%97%A4%EB%8D%94%ED%8C%8C%EC%9D%BC%EC%97%90-%EC%84%A0%EC%96%B8%ED%95%A0-%EC%88%98-%EC%9E%88%EB%8A%94-4%EA%B0%80%EC%A7%80-%ED%83%80%EC%9E%85%EC%9D%98-%EC%A0%84%EC%97%AD%EB%B3%80%EC%88%98

728x90

+ Recent posts