3D Client Programmer

그림 - CUDA 메모리 계층 구조

그림 - CUDA 메모리 특성

액세스 속도 (느림 -> 빠름)

Global -> Local -> Shared -> Register

메모리 용량 (작음 -> 큼)

Register -> Shared -> Local -> Global

액세스 범위 (좁음 -> 넓음)

Register -> Local -> Shared -> Global

Register Memory는 온 칩 프로세서에 있는 메모리로 직접 연산을 수행하는 가장 빠른 메모리이다. CUDA C를 이용해 구현할 때도 몇 개의 레지스터를 사용하는 지와 레지스터가 어떻게 재사용 되는지 알 수 없다.

다만 커널 함수 내에서 로컬 변수를 많이 사용할수록 레지스터 사용량이 증가하고, GPU마다 레지스터의 수는 차이가 있다.

sin(), cos() 등 수학 함수를 사용할 때도 레지스터를 사용한다.

로컬 메모리(Local Memory)는 커널 함수 내에서 너무 많은 로컬 변수를 사용하거나, 배열형 변수로 큰 용량을 사용하면 프로세서 밖에 있는 DRAM에 메모리가 할당된다. 로컬 변수가 레지스터로 사용될지 로컬 메모리에 할당 될지 명확하지 않다.

Local Memory로 할당되는 경우는 다음과 같다.

1. 너무 많은 레지스터 변수를 사용했을 때

2. 너무 많은 로컬 변수를 사용했을 때

3. 로컬 변수로 배열을 사용했을 때

쉽게 말하면, 각 Thread는 자신만이 사용하기 위한 메모리이다.

공유 메모리(Shared Memory)는 CUDA의 큰 장점 중 하나이다. CUDA C 컴파일러는 공유 메모리의 변수들을 일반 변수들과는 다르게 취급한다. CUDA C 컴파일러는 GPU에서 실행되는 각 블록마다 변수의 복사본을 하나씩 생성한다. 블록 내의 각 스레드는 메모리를 서로 공유하지만, 다른 블록 내에서 보고 있는 변수의 복사본을 보거나 수정할 수는 없다. 하지만 공유 메모리는 블록 내의 스레드들이 계산하는 데 통신과 협력을 할 수 있게 함으로써 훌륭한 수단을 제공한다. 더욱이 공유메모리 버퍼는 오프칩(off-chip)인 DRAM에 상주하는 것과 달리 물리적으로 GPU 상에 상주한다. 이 때문에 공유 메모리의 접근 지연 시간은 블록당 공유 메모리를 효율적으로 생성하는 일반 버퍼, 소프트웨어로 관리되는 캐시 또는 스크래치패드(Scratchpad)보다 훨씬 짧은 경향이 있다.

GPU의 SM 안에 공유 메모리가 있고, SM은 블록 단위로 프로그램을 처리 하기 때문에 공유 메모리의 데이터는 동일한 블록 안에 있는 스레드 사이에서만 공유할 수 있다. 공유 메모리를 할당하는 방법은 정작할당, 동적할당이 있다.

쉽게 말하면, 각 Thread 간에 data 공유를 위한 메모리이다.

__shared__ int a; 

전역 메모리(Global Memory)는 비디오 카드에 장착된 DRAM 메모리를 의미한다. 전역 메모리는 GPU 칩 외부에 있기 때문에 액세스 속도는 레지스터나 공유 메모리에 비해 많이 느리지만, CPU의 메모리 보다는 빠르다.

쉽게 말하면, Block들의 집합인 Grid간에 data 공유를 위한 메모리이다.

상수 메모리(Constant Memory)는 DRAM에 있는 데이터를 읽기 전용으로 사용하며 캐시를 지원한다. NVIDIA의 그래픽 카드는 Constant Memory의 크기가 64KB이다. 최초 데이터는 DRAM에서 가져오기 때문에 속도가 좀 느리지만, 한번 캐시에 올라온 값을 반복하여 재 사용하기 때문에 속도가 빠르다. Host에서는 값을 쓰고, Device 영역에서는 값을 읽기만 할 수 있다.

쉽게 말하면, Half - Warp내의 모든 스레드가 동일한 주소의 데이터를 상수 메모리로 부터 요청하면 GPU는 단 한번의 읽기만 요청한 후 모든 스레드에게 그 데이터를 알려주는 메모리이다.

__constant__ int a;

텍스처 메모리(Texture Memory)는 상수 메모리처럼 칩에 캐싱된다. 따라서 어떤 경우엔 텍스처 메모리는 오프칩인 DRAM에 대한 메모리 요청을 줄임으로써 매우 효과적인 대역폭을 제공할 수 있다. 텍스처 캐시는 공간 구역성(spatial locality)를 자주 드러내는 메모리 접근 패턴을 가진 그래픽스 어플리케이션들을 위해 설계되었다. 

texture<자료형> a;

출처 : http://blog.naver.com/riverrun17/220420579990

CUDA 메모리에 대해 공부하다가 어떤 분이 잘 정리해 놓으셔서 긁어왔습니다....

이번에는 앞전에 했던 블록을 이용한 벡터의 합을 스레드를 이용한 벡터의 합으로 구현해보려고 한다. 

사실 코드는 몇줄 바뀌지 않는다. 단순하게 블록에서 하던 일을 스레드에서 하게 바꾸는 작업만 하면 되는 일이니까...

앞에서 작성했던 코드와 비교해서 바뀐 부분은 빨간색 네모를 확인하면 된다.

그림 - add()에서 바뀐 부분

그림 - main()에서 바뀐 부분

앞에서 하드웨어의 제약으로 인해 한 번에 실행 가능한 블록의 수는 65,535개를 초과할 수 없다고 했었다. 동일하게, 하나의 커널을 실행하는 데 있어서 블록당 스레드 수 역시 하드웨어의 제약을 받는다. 구체적으로, 이 수는 우리가 3장에서 보았던 디바이스 속성 구조체의 'maxThreadsPerBlock' 멤버의 값을 초과해서는 안된다. 현재 이용가능한 많은 그래픽스 프로세서들에서 블록당 스레드의 최대 개수는 512개이다.

그렇다면 배열의 개수가 512개보다 더 클 때 어떻게 스레드 기반으로 두 벡터의 합을 수행할 수 있는지 생각해보자. 뭐 방법은 간단하다. 스레드와 블록을 같이 사용하는 방법이다.

그림 - 스레드와 블록을 같이 사용

여기서 'blockDim'이라는 내장 변수를 사용하는데 앞에서 설명했다시피, 각 블록당 스레드의 개수 이다. 여기서는 1차원의 블록을 사용하고 있기 때문에 그냥 'blockDim.x'만 사용했고, 2차원이라면 'y'도 사용 해야 할 것이다.

또 다른 변경은 커널 실행 자체에 있다. 실행할 N개의병렬 스레드들이 여전히 필요하지만, 우리에게 부과된 512개의 스레드 제한 수를 초과하지 않도록 다수의 블록들을 통해 스레드들을 발동시키길 원한다. 하나의 해결책은 임의로 블록 크기를 고정된 스레드의 개수로 지정하는 것이다. 예를 들어, 블록당 128개의 스레드들을 사용한다고 고정시켜보자.

그리고 블록을 생성하는 개수를 정할 차례인데, 여기서 묘안은 N/128이 정수의 나눗셈이라는 것이다. 만약 N이 127이라면 N/128은 정수로 0이 된다. 즉 0개의 블록을 생성하게 되므로 스레드도 실행되지 않게 된다. 하지만, 이렇게 구현해버리면, N이 정확히 128의 배수여야만 올바른 블록의 개수가 실행될 것이다. 이것은 좋지 못하다. 이를 해결하기 위한 일반적인 요령은 다음과 같다.

그림 - 커널의 다른 실행 방법

여기서 N이 10개 이므로, 블록은 10개 스레드는 128개가 된다. 즉 1280개의 스레드가 생성되어 돌아가기 때문에 잘못된 계산이 나올수 있으므로 예외 처리를 해줘야 한다. 하지만, add() 커널 함수에서 이미 예외 처리를 해놨었다.

그림 - 예외 처리가 된 부분

tid가 위에서 계산되었을 때 N보다 클 때에는 if문에서 걸러지기 때문에 올바르게 작동하게 된다.

하지만 이렇게 구현하게 되면, 하드웨어 제약사항으로 인한 블록의 개수(65,535) * 스레드의 개수(512) 개를 초과하게 되면, 병렬 처리가 정상적으로 처리가 되지 않게 된다. 그러므로 더 수정을 해야 한다.

그림 - 더 업그레이드 된 add() 커널 함수

위와 같이 각 스레드가 현재 인덱스에서의 작업을 끝내면 그리드 내에서 작동 중인 스레드의 총 수만큼 각 인덱스를 증가시킬 필요가 있다. 즉, while() 문에서 tid 값을 위와 같이 증가 시키면 각 스레드는 유니크한 인덱스를 구할 수 있을 것이다.

그림 - 고정된 개수의 블록과 스레드를 실행하는 add() 커널 함수

위와 같이 블록의 개수(128) * 스레드의 개수(128) = 16,384개 이지만, 3만개든 10만개든 덧셈이 동시에 가능해진다.

람다(lambda)는 '람다 함수' 또는 '이름 없는 함수'라고 부르며, 그 성질은 함수 객체와 같다.

기존 C++에서 STL의 find_if, sort 등의 알고리즘을 사용할 때 특정 조건자를 사용하려면 함수 객체를 정의해야 했다.

그런데 STL 알고리즘 함수에서만 사용하려고 따로 함수 객체를 만들려니 귀찮을 수 밖에 없었다.

그래서 보통은 함수를 따로 만들거나, 함수를 정의하는 것도 귀찮아서 그냥 STL을 사용안하고 직접 컨테이너를 다루었다.

람다 덕택에 이제는 이런 수고를 할 필요가 없어졌다.

람다의 기본 문법은 다음과 같다.

다음은 람다의 간단한 사용 예다.

auto를 사용하면 람다를 변수에 대입 할 수도 있다.

람다는 일반 함수처럼 파라미터를 정의 할 수도 있다.

반환 값도 사용할 수 있다.

find_if를 쓸 때에는 더욱 유용한 것을 느낄 수 있다.

기존에는 다음과 같이 find_if 알고리즘을 이용해서 '죽은 유저'를 찾으려면 다음과 같이 함수 객체를 정의해야 했다. 

하지만 람다를 사용하게 되면, 한 줄로 간단하게 끝낼 수 있다.

* 람다는 'decltype' 과 'sizeof' 에서는 사용할 수 없다.

아마도 C++ 11에서 가장 인기 있는 기능이 auto일 것이다. auto를 가장 잘 말해주는 예는

위 처럼 써야 할 것을 C++ 11에서는 

이렇게 쓸 수 있다. 

역시 auto는 그냥 일반 자료형 보다는 STL을 사용할 때 훨씬 편하게 코딩할 수 있게 만들어 준다.

* auto는 지역 변수에서만 사용할 수 있으므로, 클래스의 멤버 변수나 전역 변수, 함수의 인자로는 사용할 수 없다.

* auto는 '변수를 정의할 때 명시적으로 형을 지정하지 않고 컴파일을 할 때 자동으로 결정해주는 키워드' 이다.

이전에 커널 함수에서 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개의 스레드들의 집단과 관계가 있다. 프로그램 코드의 매 라인마다 하나의 워프 내의 각 스레드는 다른 데이터를 가지고 동일한 명령어를 수행한다.

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)라고 부른다. 

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의 관계

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

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