Dispatch 모델

이 챕터에서 정복할 CG 개념

• GPU 병렬 처리의 작업 단위인 dispatch, workgroup, thread의 위계 구조.
• fragment shader의 암묵적 dispatch와 compute shader의 명시적 dispatch가 어떻게 다른가.
• "한 thread가 한 attribute element를 처리한다"는 약속이 왜 GPGPU의 기본 계약인가.
• workgroup size가 하드웨어(NVIDIA 32, AMD 64)에 묶이는 이유.
• TDIndex(), TDNumElements(), indirect dispatch의 정확한 의미.

왜 이게 중요한가

POP은 GPU 그래픽스 파이프라인의 데이터 모델을 노드 그래프로 노출시킨 학습 환경이다. 이 학습 환경의 핵심은 attribute라는 데이터 단위이며, GLSL POP은 그 attribute를 직접 손으로 쓰게 한다. 그러나 "쓴다"는 동작은 단순한 함수 호출이 아니다. GPU는 한 번에 수천 개의 thread를 띄워 같은 코드를 서로 다른 데이터에 적용하고, 이 동시 실행의 기본 단위가 dispatch이다. dispatch 모델을 이해하지 못한 채로 GLSL POP을 만지는 것은, 마치 for 루프 없이 배열을 인덱싱하려는 것과 같다. Ch.8부터 등장하는 모든 GLSL POP 코드는 이 챕터의 dispatch 모델을 전제로 동작한다. Unity Compute Shader의 [numthreads(8,8,1)], WebGPU의 @workgroup_size(64), CUDA의 <<<blocks, threads>>>도 같은 약속의 다른 표기일 뿐이다.

POP에서의 노출 지점

GLSL POP의 GLSL 페이지에는 dispatch와 관련된 파라미터가 모여 있다.

• Number of Threads (numthreadsmode) — 메뉴: auto, otherinputelements, numelems, numelemsattrib, manual.
• Threads Input (threadsinput) — otherinputelements 모드에서 어느 입력을 기준으로 할지.
• Number of Elements (numelems) — 수동 모드에서 직접 개수를 지정.
• Num Elements Attrib (numelemsclass, numelemsattr) — 다른 POP의 attribute 첫 값에서 개수를 가져옴 (indirect dispatch).
• Work Group Size (workgroupsize x,y,z) — manual 모드에서만 사용.
• Dispatch Size (dispatchsize x,y,z) — manual 모드에서 launch할 workgroup의 개수.

GLSL Advanced POP의 같은 페이지에서는 numthreadsmode 메뉴가 더 세분화된다: inputpoint, inputvert, inputprim, outputpoint, outputvert, outputprim, numelems, numelemsattrib, manual.

코드 측에서는 다음 두 built-in이 dispatch 모델을 노출한다:

uint TDNumElements();   // 요청된 thread 개수 (workgroup 단위로 올림 전 값)
uint TDIndex();         // 현재 thread의 1차원 인덱스

두 값 모두 manual 모드에서는 정의되지 않으며 컴파일 에러를 발생시킨다.

이론

왜 GPU는 한 명령으로 수천 개를 한 번에 처리하는가

CPU는 한 코어가 복잡한 한 작업을 빠르게 한다. branch prediction, out-of-order execution, 깊은 캐시 위계로 한 thread의 latency를 최소화한다. 일반 CPU 코어는 16개 안팎.

GPU는 정반대 설계다. 한 코어(엄밀히는 한 SM)에는 수십 개의 SIMD lane이 있다. 한 cycle에 한 명령이 32개(NVIDIA) 또는 64개(AMD)의 다른 데이터에 동시 적용된다. 코어 하나의 latency는 낮지 않다 — memory access가 일어나면 수백 cycle 멈춘다. 대신 GPU는 active warp을 수십 개 띄우고, 한 warp이 memory를 기다리는 동안 다른 warp을 돌린다. latency를 throughput으로 덮는 설계.

이 설계는 한 가지 조건에 의존한다 — 작업이 충분히 많고, 작업끼리 의존성이 적어야 한다. 작업이 16개뿐이라면 SIMD lane을 채우지 못해 GPU의 90%가 놀게 된다. 작업이 강한 의존성으로 묶여 있으면 동시 실행이 불가능하다.

attribute element 처리는 정확히 이 조건을 만족한다. 백만 개의 point가 있고 각 point가 독립이면, GPU는 어떤 순서로든 처리할 수 있다. dispatch 모델은 이 가정을 "한 thread, 한 element"로 직접 표현한 것이다.

Compute shader가 다른 셰이더 stage와 어떻게 다른가

전통적 graphics pipeline은 vertex → tessellation → geometry → rasterizer → fragment stage로 흐른다. 각 stage는 정해진 입력과 출력이 있고, GPU 드라이버가 stage 사이의 데이터 흐름을 강제한다.

compute shader는 이 흐름 밖에 있다. 한 dispatch 호출이 들어오면 GPU는 정해진 수의 thread를 띄우고, 셰이더는 SSBO / image / sampler 같은 binding을 자유롭게 read/write한다. 출력 위치도, 입력 위치도, 호출 횟수도 작성자가 정한다.

이 자유가 compute shader를 다양한 용도에 쓰이게 한다:

• 입자 시뮬 (Ch.10)
• 물리 시뮬 (cloth, fluid)
• image processing (blur, tone mapping)
• particle / mesh culling
• spatial hashing (Ch.13)
• prefix sum, sorting
• ML inference / kernel evaluation

POP의 GLSL POP은 compute shader를 attribute 갱신에 특화한 셸이다. 작성자는 dispatch와 binding을 직접 다루지 않고, attribute class와 출력 이름만 정한다. 그러나 안에서 도는 코드는 일반 compute shader와 같다.

Fragment shader의 dispatch는 보이지 않는다

평소 작성하는 fragment shader는 dispatch를 의식하지 않는다. Render TOP이 viewport를 1920×1080으로 띄우면 GPU 드라이버는 약 200만 개의 픽셀 각각에 fragment shader 한 호출을 자동으로 배정한다. 어떤 픽셀이 어떤 SM(Streaming Multiprocessor)에서 돌고, 어떤 warp(NVIDIA) / wavefront(AMD)로 묶이는지는 드라이버 몫이다. 작성자는 그저 gl_FragCoord를 읽어 한 픽셀의 색을 반환한다.

Compute shader의 dispatch는 직접 명시한다

compute shader는 그릴 픽셀 같은 자연스러운 작업 단위가 없다. 그래서 작성자가 직접 dispatch 차원을 명시한다. dispatch는 세 단계 위계로 구성된다.

dispatch (호스트가 한 번 launch)
└─ workgroup × (dispatchSize.x × dispatchSize.y × dispatchSize.z)
   └─ thread × (workGroupSize.x × workGroupSize.y × workGroupSize.z)

총 thread 개수는 두 차원의 곱이다.

totalThreads = dispatchSize.x * dispatchSize.y * dispatchSize.z
             * workGroupSize.x * workGroupSize.y * workGroupSize.z

같은 workgroup 안의 thread들은 같은 SM에서 돈다. 그래서 같은 workgroup 안에서만 shared 메모리와 barrier()로 동기화가 가능하다. workgroup 사이에는 동기화가 없으며, dispatch 호출이 끝나기 전까지 순서를 가정할 수 없다.

다이어그램 — dispatch / workgroup / thread

Dispatch (한 번의 vkCmdDispatch 호출)
│
├── Workgroup (0,0,0)              ┐
│   ├── thread (0,0,0) → P[0]      │   같은 SM
│   ├── thread (1,0,0) → P[1]      │   shared mem 공유
│   ├── thread (2,0,0) → P[2]      │   barrier() 가능
│   └── ... (32 threads on NVIDIA) ┘
│
├── Workgroup (1,0,0)              ┐
│   ├── thread (0,0,0) → P[32]     │   별도 SM 가능
│   ├── thread (1,0,0) → P[33]     │   다른 workgroup과
│   └── ...                        ┘   동기화 불가
│
└── Workgroup (2,0,0)
    └── ...

workgroup size는 하드웨어 SIMD 폭에 맞춰져 있다. NVIDIA GPU의 warp는 32 thread, AMD GPU의 wavefront는 64 thread 단위로 실행되므로, workgroup size를 그 배수로 잡지 않으면 SM의 일부 lane이 놀게 된다. TouchDesigner wiki는 다음을 명시한다: "The actual number of threads is most likely higher, since it is rounded up to the work group size, which is set to 32 for NVIDIA hardware and 64 for AMD hardware."

"한 thread, 한 attribute element"라는 약속

GLSL POP의 핵심 모델은 단순하다. attribute class가 Point이고 입력에 1024개의 point가 있다면, GPU는 1024개의 thread를 띄운다. 각 thread는 자기 인덱스 TDIndex()에 해당하는 한 element만 책임진다.

const uint id = TDIndex();
if (id >= TDNumElements()) return;
// 이 thread는 오직 [id]번째 element만 다룬다

이 단순성이 GPU 병렬화의 본질이다. element 간 의존성이 없으면 GPU는 1024개를 순서 없이, 동시에, 자유롭게 처리한다. element 간 의존성이 필요해지는 순간(이웃 평균, prefix sum 등) 별도의 메커니즘이 필요해진다 — Ch.11의 ping-pong, Ch.13의 neighborhood가 그 답이다.

SIMT 실행과 워프 발산

GPU thread는 독립 실행 단위처럼 보이지만 하드웨어는 32개(NVIDIA warp) 또는 64개(AMD wavefront)를 묶어 한 instruction을 한 cycle에 같이 실행한다. SIMT(Single Instruction, Multiple Threads) 모델이다. 같은 warp 안의 thread들이 모두 같은 분기를 타면 한 cycle에 32개가 끝난다. 분기가 갈리면 두 분기 모두를 순차 실행하면서 해당 lane을 mask 한다 — warp divergence.

if (id < 100) {
    p *= 2.0;       // 일부 lane만 활성
} else {
    p *= sin(p);    // 나머지 lane이 활성, 이전 lane은 idle
}

위 분기에서 같은 warp 안에 id < 100인 thread와 그렇지 않은 thread가 섞이면 두 branch가 차례로 실행된다. 같은 warp의 모든 thread가 한쪽 branch만 타면 다른 branch는 아예 fetch되지 않는다. 그래서 분기는 가능한 warp-granular하게 묶이는 것이 좋다 — bounds-check if (id >= TDNumElements()) return;이 warp의 끝부분 lane에서만 발생하기 때문에 본 셰이더의 분기 비용은 거의 0에 가깝다.

Occupancy — 왜 너무 큰 workgroup이 나쁜가

한 SM은 register, shared memory, warp slot이 모두 한정되어 있다. workgroup이 크면 SM 위에 동시에 올릴 수 있는 workgroup 수가 줄어든다. 동시 active warp 수가 줄면 memory latency를 다른 warp로 가릴 여지가 사라진다.

위는 register 등 다른 제약을 무시한 단순 예시다. 실전에서는 셰이더의 register 사용량이 occupancy의 진짜 병목인 경우가 많다. GLSL POP의 auto 모드는 작성자가 이 계산에 신경 쓰지 않게 32 또는 64 단위로 자동 설정한다. manual 모드를 쓰는 거의 유일한 이유는 2D/3D 격자 dispatch (예: 텍스처-비슷한 attribute layout)다.

Latency hiding의 직관

한 thread가 SSBO에서 vec3를 읽으면 GPU 메모리 컨트롤러는 그 요청을 보내고 응답을 기다린다. 응답이 오기까지 약 200–400 cycle. 이 동안 그 thread는 멈춘다.

GPU의 답은 단순하다 — 그 동안 다른 warp을 돌린다. SM 위에 8개의 active warp이 있으면, 한 warp이 memory를 기다리는 동안 다른 7개 warp이 ALU를 사용. SM의 ALU lane은 거의 100% 가동 상태 유지.

이 메커니즘이 작동하려면 active warp 수가 충분해야 한다. occupancy(SM에 동시 active한 warp의 비율)가 25%면 latency hiding이 부분적. 50% 이상이면 안정적. 80% 이상이면 거의 완벽.

작성자가 occupancy에 영향을 미치는 두 변수:

• workgroup size: 큰 workgroup은 SM 위에 적은 수가 올라감.
• register 사용량: 셰이더가 큰 로컬 변수 / 큰 array를 쓰면 thread당 register 증가, 동시 thread 수 감소.

본 챕터의 단순 셰이더는 register 사용량이 매우 적고 workgroup size도 작으므로 occupancy가 거의 100%. 복잡한 셰이더에서 register pressure를 진단하려면 RenderDoc / Nsight 같은 GPU 프로파일러가 필요하다.

Indirect dispatch — GPU에 있는 숫자로 GPU를 launch

본 책의 한 단어 정의: "다음 dispatch의 thread 개수가 이전 dispatch의 결과에 의존하는데, CPU가 그 결과를 알 필요가 없는 경우." 예시:

• particle compaction: 1,000,000 입자 중 alive인 것만 다음 step에 갱신하고 싶다. alive 카운트는 GPU의 atomic counter SSBO에 있고, CPU는 그 값을 안 받아도 된다.
• adaptive subdivision: 전체 mesh를 voronoi cell 단위로 dispatch하는데, 각 cell의 vertex 수는 GPU의 prefix-sum 결과로 결정된다.

numelemsattrib 모드는 다른 POP의 한 attribute 첫 element 값을 dispatch thread 수로 쓴다. wiki 발췌: "If the number of elements for the selected class and input is known on the GPU and not on the CPU, an indirect command buffer is filled, and the compute shader is launched with an indirect dispatch."

CPU 왕복이 없어진다는 것은 한 frame 안에서 여러 단계의 GPU pipeline을 끊김 없이 chaining 할 수 있다는 뜻이다. real-time procedural geometry, GPU culling, neural radiance field의 GPU 측 ray batching이 모두 이 패턴 위에서 돈다.

Workgroup size 선택 — 32, 64, 128, 256, 512, 1024 중 어느 것

GLSL POP의 auto 모드는 작성자에게 이 결정을 숨긴다. manual 모드를 쓸 때만 직접 고른다. 결정을 위한 기준:

• NVIDIA에서 32의 배수, AMD에서 64의 배수. 그 이하면 SIMD lane 일부가 idle.
• 너무 작으면: workgroup 개수가 많아져 SM 스케줄링 overhead 증가, occupancy는 좋아질 수 있음.
• 너무 크면: SM 위에 동시에 올라가는 workgroup 수가 줄어 occupancy 감소. memory latency를 가릴 active warp이 부족해짐.
• 합리적 범위: 32 ~ 256. 256을 넘으면 register pressure가 occupancy를 더 줄임.

본 핸드북의 권장: auto로 시작. manual이 필요하면 (64, 1, 1) 또는 (128, 1, 1)에서 시작.

shared memory를 활용하는 패턴(neighborhood, prefix sum)에서는 workgroup size가 shared memory tile 크기와 직접 연결되어 결정이 더 미묘해진다 — 그 패턴은 Ch.13에서 다룬다.

bounds-check이 왜 의무인가

TDNumElements()는 작성자가 요청한 thread 개수다. 그러나 실제로 launch되는 thread 개수는 workgroup size의 배수로 올림된다. element가 100개일 때 NVIDIA에서 workgroup size 32라면, 실제로는 4 × 32 = 128 thread가 뜨고 그중 28개는 "유효 범위 밖" thread이다. 이 28개가 SSBO에 쓰면 메모리 손상이 일어난다.

if (id >= TDNumElements()) return;

TouchDesigner wiki는 단언한다: "The if condition prevents the shader to write outside of the bounds of the allocated SSBO, which can lead to unpredictible results and crashes."

Number of Threads 모드별 의미

indirect dispatch가 중요하다. GPU에서 만들어진 결과(예: 살아있는 입자의 개수)를 다시 CPU로 가져오지 않고 다음 dispatch의 thread 수로 쓸 수 있다. CPU 왕복이 사라진다.

Dispatch의 크기 결정 — 어떤 N이 좋은가

TDNumElements()로 띄울 thread 수를 결정할 때, 실용적 가이드라인:

• point 개수와 같게: 가장 일반적. 한 point에 한 thread. auto 모드.
• point 개수의 약수: 한 thread가 여러 point를 묶어 처리 (loop). throughput은 같지만 launch overhead 절감.
• point 개수의 배수: 한 point에 여러 thread (예: 한 point에 두 noise sample). 드물지만 가능.
• GPU에 있는 동적 카운트: indirect dispatch. CPU 왕복 없음.

대체로 N은 매우 큰 값(만 단위 이상)이 좋다. SM 하나당 active warp이 충분히 깔려야 latency hiding이 일어난다. N=100은 비효율 (NVIDIA 기준 한 SM에 4 warp = 한 SM만 사용). N=10,000은 충분히 많은 warp으로 모든 SM을 채움.

Dispatch와 attribute class의 짝

GLSL POP의 Attribute Class와 Number of Threads는 짝으로 움직인다. attrclass가 point이고 numthreadsmode가 auto이면, 입력의 point 개수만큼 thread가 뜬다. 같은 셰이더에서 attrclass를 vertex로 바꾸면 vertex 개수만큼 thread가 뜬다. 두 경우의 동일한 코드가 다른 element 집합을 처리한다는 뜻이다.

이 짝은 한 GLSL POP 안에서 깨지지 않는다. 점 attribute에 쓰면서 prim attribute에 쓸 수는 없다. 두 class에 동시에 쓰려면 GLSL Advanced POP의 single dispatch mode가 필요하다. 그 안에서는 한 dispatch가 point/vertex/primitive 모두에 쓸 수 있고, dispatch thread 개수는 inputpoint/inputvert/inputprim/outputpoint/outputvert/outputprim 중 하나로 명시한다.

Per Primitive Batch 모드는 한 차원 더 추가한다. wiki 발췌: "In per primitive batch mode, runs the shader once per primitive batch." 한 mesh가 여러 종류의 primitive batch(triangle / quad / line strip 등)로 구성될 때, 같은 셰이더가 각 batch에 대해 별도 dispatch로 호출된다. Ch.12에서 다룬다.

1D vs 2D vs 3D dispatch의 실용적 의미

Dispatch Size는 (x, y, z) 세 차원을 받는다. 본 챕터의 모든 예시는 1D dispatch ((N, 1, 1))를 사용한다. 2D / 3D dispatch는 어디서 쓰이나?

• 2D dispatch: 텍스처 비슷한 attribute layout (예: NxM 격자 sampling, image processing 스타일 작업). compute shader가 fragment shader처럼 픽셀 격자를 다룰 때 자연스럽다.
• 3D dispatch: 3D 격자 (voxel field, 3D 텍스처 채우기). volumetric 작업의 자연스러운 표현.

POP의 attribute는 본질적으로 1D 배열이므로 2D/3D dispatch는 manual 모드에서 작성자가 직접 인덱스 계산을 해야 한다. wiki는 manual 모드의 TDIndex()가 undefined임을 명시한다. 작성자는 gl_GlobalInvocationID.x + gl_GlobalInvocationID.y * dispatchSize.x 같은 수식을 직접 만들어야 하지만, TD가 gl_GlobalInvocationID를 노출하는지는 wiki 미확정 — 실험 필요.

본 챕터의 실용 권장: GLSL POP에서 1D auto로 시작하고, 2D/3D 격자가 필요해질 때만 manual로 넘어간다.

한 dispatch 안에서 일어나는 일의 시간 순서

작성자가 직관적으로 가정하기 쉬운 잘못된 모델이 있다 — "thread 0이 먼저 끝나고 thread 1이 그 다음에…". GPU의 실제 모델은 그렇지 않다.

dispatch가 launch되면:

1. GPU 드라이버가 workgroup들을 SM에 배정. 어느 workgroup이 어느 SM으로 가는지는 동적 결정.
2. 각 SM은 자기 workgroup의 warp을 active 슬롯에 올린다. SM이 가진 슬롯 수만큼 동시 active.
3. SM의 instruction scheduler가 매 cycle 한 active warp을 골라 명령 한 줄을 실행.
4. warp이 memory 또는 long-latency 명령에 막히면 다른 active warp으로 전환.
5. 모든 workgroup이 끝나면 dispatch 완료.

이 흐름에서 thread 0과 thread 1024가 같은 SM의 다른 warp에 들어가 거의 동시에 끝날 수도 있고, 매우 다른 시간에 끝날 수도 있다. 순서를 가정하지 말 것.

가정 가능한 것:

• 같은 warp 안의 32 thread는 lockstep으로 같이 진행.
• 같은 workgroup의 thread들은 barrier()로 한 지점에서 만남.
• dispatch 호출이 끝났을 때 모든 thread가 완료된 상태.

본 챕터의 단순 셰이더는 thread 간 의존성이 없으므로 이 모든 미묘함이 보이지 않는다. 그러나 neighborhood나 누적이 들어가면 즉시 중요해진다.

손작업 (Hands-on)

시작 파일

- Sphere POP1 (Primitive Type: Geodesic, Frequency: 4)
- GLSL POP1
- Text DAT1 (Compute Shader 코드)
- Geometry COMP1 (Sphere의 출력을 받음)
- Render TOP1

연결: Sphere POP1 → GLSL POP1 → Geometry COMP1 (POP input) → Render TOP1

GLSL POP1 파라미터:

• Compute Shader (computedat): Text DAT1
• Attribute Class (attrclass): point
• Number of Threads (numthreadsmode): auto
• Output Attributes (outputattrs): P
• Initialize Output Attributes (initoutputattrs): On

Step 1 — 통로만 만든다

Text DAT1에 다음을 붙여 넣는다.

// dispatch 셸: 입력 P를 그대로 출력에 쓴다. 변형은 없다.
void main() {
    const uint id = TDIndex();
    if (id >= TDNumElements()) return;
    P[id] = TDIn_P();
}

이 코드가 GPU에서 실제로 하는 일: GPU는 auto 모드에 따라 입력 point 개수만큼 thread를 띄운다. NVIDIA 기준 workgroup size 32로 올림되어 실제 thread는 그보다 많을 수 있다. 각 thread는 자기 TDIndex() 위치의 입력 P 값을 SSBO에서 읽어, 같은 위치의 출력 P SSBO에 쓴다. bounds-check이 28개의 "넘친" thread가 메모리 밖에 쓰는 것을 막는다.

Render TOP에 sphere가 그대로 보이면 dispatch가 정상 동작한 것이다.

Step 2 — numelems 모드 비교

같은 셰이더를 그대로 두고, Number of Threads를 numelems로 바꾸고 Number of Elements를 입력 sphere의 정확한 point 개수로 직접 입력한다. Sphere POP의 info CHOP에서 point 개수를 읽거나 viewer에서 확인.

결과: 동일.

관찰: dispatch 모드를 바꿔도 결과가 같다. 약속이 같기 때문이다 — "한 thread, 한 element".

Step 3 — manual 모드는 어떻게 다른가

Number of Threads를 manual로 바꾼다. Work Group Size를 (32, 1, 1)로, Dispatch Size를 ceil(numPoints / 32)로 설정.

이대로 컴파일하면 에러가 난다. wiki 발췌: "With the 'Manual' mode, the work group size and number of dispatches can be specified in 3 dimensions, in that case TDNumElements() and TDIndex() are undefined and will causes a compile error."

manual 모드에서는 작성자가 직접 인덱스를 계산해야 한다. wiki는 gl_GlobalInvocationID가 노출되는지 명시하지 않는다 — wiki 미확정 — 실험 필요.

Step 4 — workgroup size 32×1×1 vs 8×8×1

manual 모드를 유지하고 두 가지 설정을 비교:

• (A) Work Group Size = (32, 1, 1), Dispatch Size = (ceil(N/32), 1, 1)
• (B) Work Group Size = (8, 8, 1), Dispatch Size = (ceil(N/64), 1, 1)

두 경우 모두 총 thread ≥ N이어야 한다. attribute가 1D 배열이므로 두 경우 모두 같은 element를 같은 방식으로 다루지만 thread를 GPU 위에 배치하는 방식만 다르다. 결과는 같다. 차이는 occupancy와 캐시 동작에 있다 — element 간 의존성이 없는 본 셸에서는 차이가 거의 없다.

TDIndex의 실체 — 어떤 built-in을 wrapping하는가

TDIndex()는 TouchDesigner가 Vulkan compute shader의 raw built-in 위에 만든 wrapper다. 추론적으로 그 안에는 다음과 유사한 식이 들어 있을 것이다:

uint TDIndex() {
    return gl_GlobalInvocationID.x;  // auto / non-manual 모드 가정
}

auto 모드에서는 dispatch가 1D이고 workgroup도 (N, 1, 1)이므로 gl_GlobalInvocationID.x가 곧 1차원 thread index가 된다. manual 모드에서 dispatch가 2D / 3D가 되면 단순한 wrapping이 깨지므로 TD는 TDIndex()를 undefined로 둔다 — 작성자가 의도한 인덱스가 2D 격자 위치인지 1D linear 인덱스인지 모르기 때문이다.

TDNumElements()도 비슷하다. 입력 attribute 개수에서 결정되는 "논리적 thread 요청 수"이며, 실제 launch된 thread 수(workgroup 단위로 올림된 값)보다 작다. bounds-check이 이 차이를 처리한다.

추가로 wiki에서 보장하는 입력 카운팅 helper:

uint TDInputNumPoints(uint inputIndex);
uint TDInputNumPrims(uint inputIndex);
uint TDInputNumVerts(uint inputIndex);
TDInputNumPoints() = TDInputNumPoints(0);  // 기본 input 0
uint TDInputNumElements();   // GLSL POP에서 attr class에 해당하는 카운트 반환

이들은 입력별 / class별 카운트를 직접 얻는다. TDIndex()와 다른 입력의 element 개수를 비교해 분기를 짤 때 유용하다.

메모리 모델 — 어디에 어떤 메모리가 있는가

compute shader 작성자가 다루는 메모리는 세 종류다.

SSBO read/write는 GB/s 단위로 측정되는 가장 느린 메모리 access다. shared memory는 100배쯤 빠르다. register는 거의 무료에 가깝다. 셰이더 최적화의 황금률은 SSBO read 횟수를 줄이고 shared memory와 register를 최대한 활용하는 것이다.

본 챕터의 단순 셰이더(P[id] = TDIn_P())는 SSBO read 한 번, SSBO write 한 번. element 간 의존성이 없으므로 shared memory의 필요가 없다. neighborhood operation(Ch.13)이 들어오면 같은 P를 여러 thread가 반복 read하게 되고, 그때 shared memory tiling이 중요해진다.

SSBO read 패턴과 coalescing

GPU memory access는 인접 thread가 인접 메모리를 읽을 때 가장 효율적이다. 한 warp의 32 thread가 32개 연속 vec3를 읽으면 GPU는 그것을 한두 번의 memory transaction으로 묶는다 — coalesced access. 같은 32 thread가 흩어진 메모리 위치를 읽으면 32번의 transaction이 필요하다.

POP의 attribute는 element 인덱스 순서로 SSBO에 저장된다. TDIn_P(0, TDIndex())는 thread id 순서로 P를 읽으므로 coalesced된다. 반면 TDIn_P(0, hash(TDIndex())) 같은 random access는 coalesce되지 않는다.

본 챕터의 단순 셰이더는 coalesced access이므로 memory bandwidth 효율이 매우 좋다. neighborhood 패턴(Ch.13)이 들어가면 같은 P를 여러 thread가 다른 인덱스로 읽게 되고, 그때 coalescing이 깨진다. 그래서 spatial hash 같은 자료구조로 access 패턴을 재배치하는 것이 의미를 가진다.

동기화 도구들

같은 workgroup 안에서만 사용 가능:

• barrier() — 같은 workgroup의 모든 thread가 같은 지점에 도달할 때까지 대기.
• memoryBarrierShared() — shared memory write의 가시성 보장.
• memoryBarrierBuffer() — SSBO write의 가시성 보장.

workgroup 간에는 위 함수들이 도움이 안 된다. dispatch가 끝나야 다른 workgroup의 write가 보장된다. 즉 multi-step 계산은 여러 dispatch로 쪼개야 한다. ping-pong 패턴(Ch.11)의 본질이 이것이다.

TD의 GLSL POP Passes 파라미터는 같은 셰이더를 N번 dispatch하는 기능이고, 이는 workgroup 간 동기화의 한 표현이다.

노드 ↔︎ GLSL 매핑

이 챕터의 주제는 dispatch 자체이므로 노드 vs GLSL 비교 대신, 같은 GLSL 셰이더를 서로 다른 Number of Threads 모드로 dispatch했을 때 결과가 동일한가를 검증한다.

dispatch 모드는 "몇 개의 thread를, 어떤 모양으로 띄울 것인가"의 약속만 바꾼다. 약속 안에 들어가는 일은 같다.

Dispatch가 다른 환경에서 어떻게 보이는가

같은 약속을 다른 표기로 보여주면 본질이 더 또렷이 드러난다.

Unity Compute Shader (HLSL)

#pragma kernel CSMain

RWStructuredBuffer<float3> positions;

[numthreads(64, 1, 1)]
void CSMain (uint3 id : SV_DispatchThreadID) {
    uint i = id.x;
    if (i >= numElements) return;
    positions[i] *= 1.01;
}

C# 측에서 Dispatch(threadGroupsX, 1, 1) 호출. [numthreads(64,1,1)]은 TouchDesigner의 Work Group Size = (64,1,1). SV_DispatchThreadID는 TDIndex()에 해당하지만, Unity는 manual mode가 기본이므로 작성자가 직접 인덱스를 계산한다.

WebGPU (WGSL)

@group(0) @binding(0) var<storage, read_write> positions: array<vec3f>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3u) {
    let i = gid.x;
    if (i >= arrayLength(&positions)) { return; }
    positions[i] = positions[i] * 1.01;
}

JS 측에서 passEncoder.dispatchWorkgroups(workgroupsX, 1, 1) 호출. @workgroup_size(64)는 TD의 Work Group Size, global_invocation_id는 TDIndex() 위치에 해당하는 built-in.

CUDA

__global__ void update(float3* positions, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= N) return;
    positions[i].x *= 1.01f;
}
// 호스트:
update<<<(N + 63) / 64, 64>>>(d_positions, N);

<<<blocks, threads>>>의 첫 인자가 Dispatch Size, 둘째 인자가 Work Group Size. blockIdx, blockDim, threadIdx의 곱이 TDIndex()에 해당.

동일한 약속을 표 한 장으로

용어가 다르고 표기가 다르지만 모델은 같다. 한 환경에서 익히면 다른 환경에 그대로 옮겨진다.

확인 질문 (Self-check)

1. auto 모드에서 1000개의 point가 있고 GPU가 NVIDIA(workgroup size 32)일 때, 실제로 launch되는 thread 개수는?
2. TDIndex()가 manual 모드에서 정의되지 않는 이유는 무엇인가? compute shader 자체의 어떤 built-in을 GLSL POP이 wrap하고 있다고 추측할 수 있는가?
3. workgroup size 32×1×1과 8×8×1로 같은 1D attribute 작업을 dispatch했을 때 결과가 같다면, 두 설정의 실제 차이는 어디서 측정되어야 하는가? (힌트: SIMD lane utilization, cache locality)
4. indirect dispatch (numelemsattrib)가 CPU 왕복을 없애는 시나리오를 하나 구체적으로 들어보라.
5. fragment shader에서 dispatch가 "보이지 않는" 것은 왜인가? compute shader에서 dispatch를 명시해야 하는 것과 어떤 점에서 본질적으로 다른가?

연결 고리

• LearnOpenGL — Compute Shaders / Introduction (Guest Article, 2022): https://learnopengl.com/Guest-Articles/2022/Compute-Shaders/Introduction. "Compute Shader Stage", "Altering local size" 절이 workgroup의 의미를 GL 맥락에서 그대로 풀어낸다.
• WebGPU Fundamentals — Compute Shader Basics: https://webgpufundamentals.org/webgpu/lessons/webgpu-compute-shaders.html. "Workgroup Fundamentals", "Builtin Variables" 절에 local_invocation_id / workgroup_id / global_invocation_id가 정리되어 있다. WebGPU의 @workgroup_size는 TouchDesigner의 Work Group Size 파라미터와 같은 것을 다른 표기로 한 것이다.
• Khronos OpenGL Wiki — Compute Shader: https://wikis.khronos.org/opengl/Compute_Shader. dispatch / shared memory / barrier의 normative reference.
• Real-Time Rendering 4ed §3 "The Graphics Processing Unit" — SIMT 실행, warp/wavefront, occupancy의 책 측 framing.

이 챕터의 한 줄 명제

한 thread가 한 attribute element를 맡는다는 약속이 GLSL POP의 모든 코드를 떠받친다. workgroup은 그 약속을 GPU 하드웨어에 묶는 단위다.