이 게시물에서는 예시를 통해 실험용 WebGPU API를 살펴보고 GPU를 사용하여 데이터 병렬 계산을 시작하는 데 도움이 됩니다.
배경
이미 알고 계시겠지만 그래픽 처리 장치 (GPU)는 원래 그래픽 처리에 특화된 컴퓨터 내의 전자 하위 시스템입니다. 그러나 지난 10년 동안 보다 유연한 아키텍처로 발전하여 개발자가 GPU의 고유한 아키텍처를 활용하면서 3D 그래픽을 렌더링할 뿐만 아니라 다양한 유형의 알고리즘을 구현할 수 있게 되었습니다. 이러한 기능을 GPU 컴퓨팅이라고 하며, GPU를 범용 과학 컴퓨팅의 보조 프로세서로 사용하는 것을 범용 GPU (GPGPU) 프로그래밍이라고 합니다.
컨볼루션 신경망 및 기타 모델이 이 아키텍처를 활용하여 GPU에서 보다 효율적으로 실행할 수 있기 때문에 GPU 컴퓨팅은 최근 머신러닝 호황에 크게 기여했습니다. 현재 웹 플랫폼에는 GPU 컴퓨팅 기능이 부족하기 때문에 W3C의 '웹용 GPU' 커뮤니티 그룹은 대부분의 최신 기기에서 사용할 수 있는 최신 GPU API를 노출하는 API를 설계하고 있습니다. 이 API를 WebGPU라고 합니다.
WebGPU는 WebGL과 같은 하위 수준 API입니다. 보시다시피 매우 강력하고 장황합니다. 하지만 괜찮습니다. 우리가 찾고 있는 것은 성능입니다.
이 문서에서는 WebGPU의 GPU 컴퓨팅 부분을 중점적으로 설명하겠습니다. 솔직히 말씀드리자면 사용자가 직접 플레이를 시작할 수 있도록 일부 내용을 간략하게만 살펴봤습니다. 향후 도움말에서 WebGPU 렌더링 (캔버스, 텍스처 등)에 대해 자세히 알아보고 설명하겠습니다.
GPU 액세스
WebGPU에서는 GPU에 쉽게 액세스할 수 있습니다. navigator.gpu.requestAdapter()를 호출하면 GPU 어댑터를 사용하여 비동기식으로 확인되는 자바스크립트 프로미스가 반환됩니다. 이 어댑터를 그래픽 카드라고 생각하세요. CPU와 동일한 칩에 통합되거나 분리(일반적으로 더 성능이 우수하지만 더 많은 전력을 사용하는 PCIe 카드)일 수 있습니다.
GPU 어댑터가 있으면 adapter.requestDevice()를 호출하여 일부 GPU 계산에 사용할 GPU 기기로 확인되는 프로미스를 가져옵니다.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
두 기능 모두 원하는 어댑터 종류 (전원 환경설정)와 기기 (확장 프로그램, 한도)를 구체적으로 지정할 수 있는 옵션을 사용합니다. 편의상 이 문서에서는 기본 옵션을 사용합니다.
버퍼 메모리 쓰기
자바스크립트를 사용하여 GPU의 메모리에 데이터를 쓰는 방법을 알아보겠습니다. 최신 웹브라우저에서 사용되는 샌드박스 모델 때문에 이 프로세스는 간단하지 않습니다.
아래 예는 GPU에서 액세스할 수 있는 버퍼 메모리에 4바이트를 쓰는 방법을 보여줍니다. 버퍼 크기와 사용량을 가져오는 device.createBuffer()를 호출합니다. 이 특정 호출에는 GPUBufferUsage.MAP_WRITE 사용 플래그가 필요하지 않지만 이 버퍼에 쓰려고 함을 명시적으로 가정해 보겠습니다. 이렇게 하면 mappedAtCreation가 true로 설정되어 생성 시 GPU 버퍼 객체가 매핑됩니다. 그런 다음 GPU 버퍼 메서드 getMappedRange()를 호출하여 연결된 원시 바이너리 데이터 버퍼를 검색할 수 있습니다.
바이트 쓰기는 이미 ArrayBuffer를 사용해 본 적이 있다면 익숙할 것입니다. TypedArray를 사용하여 값을 복사하세요.
// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuBuffer = device.createBuffer({
mappedAtCreation: true,
size: 4,
usage: GPUBufferUsage.MAP_WRITE
});
const arrayBuffer = gpuBuffer.getMappedRange();
// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);
이 시점에서 GPU 버퍼가 매핑됩니다. 즉, CPU가 소유하며 JavaScript에서 읽기/쓰기로 액세스할 수 있습니다. GPU에서 액세스할 수 있도록 하려면 gpuBuffer.unmap()를 호출하는 것만큼 간단하게 매핑을 해제해야 합니다.
매핑됨/매핑되지 않음의 개념은 GPU와 CPU가 동시에 메모리에 액세스하는 경합 상태를 방지하는 데 필요합니다.
버퍼 메모리 읽기
이제 GPU 버퍼를 다른 GPU 버퍼에 복사하여 다시 읽는 방법을 살펴보겠습니다.
첫 번째 GPU 버퍼를 작성하고 두 번째 GPU 버퍼에 복사하려고 하므로 새 사용 플래그 GPUBufferUsage.COPY_SRC가 필요합니다. 두 번째 GPU 버퍼는 이번에는 device.createBuffer()를 사용하여 매핑되지 않은 상태로 생성됩니다. 사용 플래그는 첫 번째 GPU 버퍼의 대상으로 사용되고 GPU 복사 명령어가 실행되면 JavaScript에서 읽히므로 GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ입니다.
// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuWriteBuffer = device.createBuffer({
mappedAtCreation: true,
size: 4,
usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC
});
const arrayBuffer = gpuWriteBuffer.getMappedRange();
// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);
// Unmap buffer so that it can be used later for copy.
gpuWriteBuffer.unmap();
// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
size: 4,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});
GPU는 독립적인 보조 프로세서이므로 모든 GPU 명령어는 비동기식으로 실행됩니다. 이러한 이유로 GPU 명령어 목록이 빌드되어 필요할 때 일괄적으로 전송됩니다. WebGPU에서 device.createCommandEncoder()에 의해 반환되는 GPU 명령어 인코더는 특정 시점에 GPU로 전송될 '버퍼링된' 명령어 배치를 빌드하는 자바스크립트 객체입니다. 반면 GPUBuffer의 메서드는 '버퍼링되지 않음', 즉 호출 시 원자적으로 실행됩니다.
GPU 명령어 인코더가 있으면 아래와 같이 copyEncoder.copyBufferToBuffer()를 호출하여 나중에 실행할 수 있도록 이 명령어를 명령어 큐에 추가합니다.
마지막으로 copyEncoder.finish()를 호출하여 명령어 인코딩을 완료하고 GPU 기기 명령어 큐에 제출합니다. 큐는 GPU 명령어를 인수로 사용하여 device.queue.submit()를 통해 제출된 작업을 처리하는 역할을 합니다.
이렇게 하면 배열에 저장된 모든 명령어가 순서대로 원자적으로 실행됩니다.
// Encode commands for copying buffer to buffer.
const copyEncoder = device.createCommandEncoder();
copyEncoder.copyBufferToBuffer(
gpuWriteBuffer /* source buffer */,
0 /* source offset */,
gpuReadBuffer /* destination buffer */,
0 /* destination offset */,
4 /* size */
);
// Submit copy commands.
const copyCommands = copyEncoder.finish();
device.queue.submit([copyCommands]);
이 시점에서 GPU 큐 명령어가 전송되었지만 반드시 실행되지는 않습니다.
두 번째 GPU 버퍼를 읽으려면 GPUMapMode.READ로 gpuReadBuffer.mapAsync()를 호출합니다. GPU 버퍼가 매핑되면 해결될 프로미스를 반환합니다. 그런 다음 큐에 추가된 모든 GPU 명령어가 실행되면 첫 번째 GPU 버퍼와 동일한 값이 포함된 gpuReadBuffer.getMappedRange()를 사용하여 매핑된 범위를 가져옵니다.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
이 샘플을 사용해 볼 수 있습니다.
요약하면, 버퍼 메모리 연산과 관련하여 기억해야 할 사항은 다음과 같습니다.
- 기기 대기열 제출에 사용하려면 GPU 버퍼를 매핑 해제해야 합니다.
- 매핑되면 GPU 버퍼를 JavaScript로 읽고 쓸 수 있습니다.
- GPU 버퍼는
mappedAtCreation가 true로 설정된mapAsync()및createBuffer()가 호출될 때 매핑됩니다.
셰이더 프로그래밍
GPU에서 실행되는 프로그램을 계산만 하고 삼각형을 그리지 않는 프로그램을 컴퓨팅 셰이더라고 합니다. 함께 작동하여 데이터를 처리하는 수백 개의 GPU 코어 (CPU 코어보다 작음)에 의해 병렬로 실행됩니다. 입력 및 출력은 WebGPU의 버퍼입니다.
WebGPU에서 컴퓨팅 셰이더를 사용하는 방법을 설명하기 위해 아래 그림에서 머신러닝의 공통 알고리즘인 행렬 곱셈을 사용해 보겠습니다.
간단히 말씀드리면 다음과 같습니다.
- GPU 버퍼 3개 (곱할 행렬용 2개, 결과 행렬용 1개)를 만듭니다.
- 컴퓨팅 셰이더의 입력 및 출력 설명
- 컴퓨팅 셰이더 코드 컴파일
- 컴퓨팅 파이프라인 설정
- 인코딩된 명령어를 GPU에 일괄 제출
- 결과 행렬 GPU 버퍼 읽기
GPU 버퍼 생성
편의상 행렬은 부동 소수점 수의 목록으로 표현됩니다. 첫 번째 요소는 행 수, 두 번째 요소는 열 수, 나머지는 행렬의 실제 숫자입니다.
컴퓨팅 셰이더에서 데이터를 저장하고 검색해야 하므로 세 개의 GPU 버퍼는 스토리지 버퍼입니다. 이는 GPU 버퍼 사용 플래그에 모든 버퍼에 대해 GPUBufferUsage.STORAGE가 포함된 이유를 설명합니다. 결과 행렬 사용 플래그에는 GPUBufferUsage.COPY_SRC도 있습니다. 모든 GPU 큐 명령어가 모두 실행되면 읽기를 위해 다른 버퍼에 복사되기 때문입니다.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
// First Matrix
const firstMatrix = new Float32Array([
2 /* rows */, 4 /* columns */,
1, 2, 3, 4,
5, 6, 7, 8
]);
const gpuBufferFirstMatrix = device.createBuffer({
mappedAtCreation: true,
size: firstMatrix.byteLength,
usage: GPUBufferUsage.STORAGE,
});
const arrayBufferFirstMatrix = gpuBufferFirstMatrix.getMappedRange();
new Float32Array(arrayBufferFirstMatrix).set(firstMatrix);
gpuBufferFirstMatrix.unmap();
// Second Matrix
const secondMatrix = new Float32Array([
4 /* rows */, 2 /* columns */,
1, 2,
3, 4,
5, 6,
7, 8
]);
const gpuBufferSecondMatrix = device.createBuffer({
mappedAtCreation: true,
size: secondMatrix.byteLength,
usage: GPUBufferUsage.STORAGE,
});
const arrayBufferSecondMatrix = gpuBufferSecondMatrix.getMappedRange();
new Float32Array(arrayBufferSecondMatrix).set(secondMatrix);
gpuBufferSecondMatrix.unmap();
// Result Matrix
const resultMatrixBufferSize = Float32Array.BYTES_PER_ELEMENT * (2 + firstMatrix[0] * secondMatrix[1]);
const resultMatrixBuffer = device.createBuffer({
size: resultMatrixBufferSize,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
});
그룹 레이아웃 및 그룹 바인드
바인드 그룹 레이아웃과 바인드 그룹의 개념은 WebGPU에만 적용됩니다. 바인드 그룹 레이아웃은 셰이더에서 예상하는 입력/출력 인터페이스를 정의하는 반면, 바인드 그룹은 셰이더의 실제 입력/출력 데이터를 나타냅니다.
아래 예에서 바인드 그룹 레이아웃은 번호가 지정된 항목 바인딩 0, 1에 읽기 전용 스토리지 버퍼 두 개와 컴퓨팅 셰이더의 2에 스토리지 버퍼를 예상합니다.
반면 이 바인드 그룹 레이아웃에 대해 정의된 바인드 그룹은 GPU 버퍼를 항목에 연결합니다. gpuBufferFirstMatrix는 바인딩 0에, gpuBufferSecondMatrix는 바인딩 1에, resultMatrixBuffer는 바인딩 2에 연결합니다.
const bindGroupLayout = device.createBindGroupLayout({
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: "read-only-storage"
}
},
{
binding: 1,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: "read-only-storage"
}
},
{
binding: 2,
visibility: GPUShaderStage.COMPUTE,
buffer: {
type: "storage"
}
}
]
});
const bindGroup = device.createBindGroup({
layout: bindGroupLayout,
entries: [
{
binding: 0,
resource: {
buffer: gpuBufferFirstMatrix
}
},
{
binding: 1,
resource: {
buffer: gpuBufferSecondMatrix
}
},
{
binding: 2,
resource: {
buffer: resultMatrixBuffer
}
}
]
});
셰이더 코드 계산
행렬 곱셈을 위한 컴퓨팅 셰이더 코드는 SPIR-V로 쉽게 변환할 수 있는 WebGPU 셰이더 언어인 WGSL로 작성됩니다. 자세히 설명하지 않아도 var<storage>로 식별된 세 개의 스토리지 버퍼를 아래에서 확인할 수 있습니다. 프로그램은 firstMatrix 및 secondMatrix를 입력으로 사용하고 resultMatrix를 출력으로 사용합니다.
각 스토리지 버퍼에는 위에서 선언된 바인드 그룹 레이아웃과 바인드 그룹에 정의된 동일한 색인에 상응하는 binding 장식이 사용됩니다.
const shaderModule = device.createShaderModule({
code: `
struct Matrix {
size : vec2f,
numbers: array<f32>,
}
@group(0) @binding(0) var<storage, read> firstMatrix : Matrix;
@group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
@group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;
@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) global_id : vec3u) {
// Guard against out-of-bounds work group sizes
if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
return;
}
resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);
let resultCell = vec2(global_id.x, global_id.y);
var result = 0.0;
for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
let a = i + resultCell.x * u32(firstMatrix.size.y);
let b = resultCell.y + i * u32(secondMatrix.size.y);
result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
}
let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
resultMatrix.numbers[index] = result;
}
`
});
파이프라인 설정
컴퓨팅 파이프라인은 수행할 컴퓨팅 작업을
실제로 설명하는 객체입니다 device.createComputePipeline()를 호출하여 만듭니다.
앞서 만든 바인드 그룹 레이아웃과 컴퓨팅 셰이더의 진입점 (main WGSL 함수)을 정의하는 컴퓨팅 단계, device.createShaderModule()로 만든 실제 컴퓨팅 셰이더 모듈의 두 가지 인수를 사용합니다.
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
명령어 제출
세 개의 GPU 버퍼와 바인드 그룹 레이아웃이 있는 컴퓨팅 파이프라인으로 바인드 그룹을 인스턴스화한 후 이를 사용할 차례입니다.
commandEncoder.beginComputePass()로 프로그래밍 가능한 컴퓨팅 패스 인코더를 시작해 보겠습니다. 이를 사용하여 행렬 곱셈을 수행할
GPU 명령어를 인코딩합니다 passEncoder.setPipeline(computePipeline)로 파이프라인을 설정하고 색인 0에 바인드 그룹을 passEncoder.setBindGroup(0, bindGroup)로 설정합니다. 색인 0은 WGSL 코드의 group(0) 장식에 상응합니다.
이제 이 컴퓨팅 셰이더가 GPU에서 어떻게 실행되는지 알아보겠습니다. Google의 목표는 결과 행렬의 각 셀에 대해 이 프로그램을 단계별로 병렬로 실행하는 것입니다. 예를 들어 크기가 16x32인 결과 행렬의 경우 실행 명령어를 인코딩하려면 @workgroup_size(8, 8)에서 passEncoder.dispatchWorkgroups(2, 4) 또는 passEncoder.dispatchWorkgroups(16 / 8, 32 / 8)를 호출합니다.
첫 번째 인수 'x'는 첫 번째 차원이고, 두 번째 인수 'y'는 두 번째 차원이며, 마지막 인수 'z'는 세 번째 차원이며 여기에서는 기본값인 1로 설정됩니다.
GPU 컴퓨팅 환경에서는 데이터 세트에서 커널 함수를 실행하기 위해 명령어를 인코딩하는 것을 디스패치라고 합니다.
컴퓨팅 셰이더의 작업 그룹 그리드 크기는 WGSL 코드에서 (8, 8)입니다. 따라서 첫 번째 행렬의 행 수와 두 번째 행렬의 열 수에 해당하는 'x'와 'y'가 각각 8로 나뉩니다. 따라서 이제 passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8)를 사용하여 컴퓨팅 호출을 전달할 수 있습니다. 실행할 작업 그룹 그리드의 수는 dispatchWorkgroups() 인수입니다.
위 그림에서 볼 수 있듯이 각 셰이더는 계산할 결과 행렬 셀을 파악하는 데 사용되는 고유한 builtin(global_invocation_id) 객체에 액세스할 수 있습니다.
const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const workgroupCountX = Math.ceil(firstMatrix[0] / 8);
const workgroupCountY = Math.ceil(secondMatrix[1] / 8);
passEncoder.dispatchWorkgroups(workgroupCountX, workgroupCountY);
passEncoder.end();
컴퓨팅 패스 인코더를 종료하려면 passEncoder.end()를 호출합니다. 그런 다음 copyBufferToBuffer를 사용하여 결과 행렬 버퍼를 복사하는 대상으로 사용할 GPU 버퍼를 만듭니다. 마지막으로 copyEncoder.finish()로 명령어 인코딩을 완료하고 GPU 명령어로 device.queue.submit()를 호출하여 해당 명령어를 GPU 기기 큐에 제출합니다.
// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
size: resultMatrixBufferSize,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});
// Encode commands for copying buffer to buffer.
commandEncoder.copyBufferToBuffer(
resultMatrixBuffer /* source buffer */,
0 /* source offset */,
gpuReadBuffer /* destination buffer */,
0 /* destination offset */,
resultMatrixBufferSize /* size */
);
// Submit GPU commands.
const gpuCommands = commandEncoder.finish();
device.queue.submit([gpuCommands]);
결과 행렬 읽기
결과 행렬을 읽는 것은 GPUMapMode.READ로 gpuReadBuffer.mapAsync()를 호출하고 반환 프로미스가 확인될 때까지 기다리면 GPU 버퍼가 이제 매핑되었음을 나타내기만 하면 됩니다. 이 시점에서 gpuReadBuffer.getMappedRange()를 사용하여 매핑된 범위를 가져올 수 있습니다.
코드에서 DevTools JavaScript 콘솔에 기록된 결과는 '2, 2, 50, 60, 114, 140'입니다.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));
수고하셨습니다 기본 검색 엔진이 Google로 설정된 것입니다. 샘플을 사용해 볼 수 있습니다.
마지막 유용한 정보
코드를 더 쉽게 읽을 수 있도록 하는 한 가지 방법은 컴퓨팅 파이프라인의 편리한 getBindGroupLayout 메서드를 사용하여 셰이더 모듈에서 바인드 그룹 레이아웃을 추론하는 것입니다. 이 방법을 사용하면 아래와 같이 커스텀 바인드 그룹 레이아웃을 만들고 컴퓨팅 파이프라인에서 파이프라인 레이아웃을 지정할 필요가 없습니다.
이전 샘플의 getBindGroupLayout 그림을 사용할 수 있습니다.
const computePipeline = device.createComputePipeline({
- layout: device.createPipelineLayout({
- bindGroupLayouts: [bindGroupLayout]
- }),
compute: {
-// Bind group layout and bind group
- const bindGroupLayout = device.createBindGroupLayout({
- entries: [
- {
- binding: 0,
- visibility: GPUShaderStage.COMPUTE,
- buffer: {
- type: "read-only-storage"
- }
- },
- {
- binding: 1,
- visibility: GPUShaderStage.COMPUTE,
- buffer: {
- type: "read-only-storage"
- }
- },
- {
- binding: 2,
- visibility: GPUShaderStage.COMPUTE,
- buffer: {
- type: "storage"
- }
- }
- ]
- });
+// Bind group
const bindGroup = device.createBindGroup({
- layout: bindGroupLayout,
+ layout: computePipeline.getBindGroupLayout(0 /* index */),
entries: [
실적 발견 항목
그렇다면 GPU에서 행렬 곱셈을 실행하는 것은 CPU에서 실행하는 것과 어떻게 다를까요? 알아보기 위해 방금 CPU에 대해 설명한 프로그램을 작성했습니다. 또한 아래 그래프에서 볼 수 있듯이 행렬 크기가 256x256보다 크면 GPU의 성능을 최대한 활용하는 것이 명백합니다.
이 문서는 WebGPU 탐색 여정의 시작에 불과합니다. GPU 컴퓨팅 및 WebGPU에서의 렌더링(캔버스, 텍스처, 샘플러) 작동 방식에 대한 자세한 내용이 담긴 추가 도움말이 곧 제공될 예정입니다.