웹에서 GPU 컴퓨팅 시작하기

이 게시물에서는 예시를 통해 실험용 WebGPU API를 살펴보고 GPU를 사용하여 데이터 병렬 계산을 수행하는 것으로 시작합니다.

François Beaufort
François Beaufort

배경

아시다시피 그래픽 처리 장치 (GPU)는 원래 처리 작업에 특화된 컴퓨터 내의 하위 시스템으로서 있습니다. 그러나 지난 10년 동안 생성형 AI는 더 유연한 방식으로 이 아키텍처를 통해 개발자는 단순한 프로그래밍 언어뿐만 아니라 API의 고유한 아키텍처를 활용하여 3D 그래픽을 사용할 수 있습니다. 이러한 기능을 GPU 컴퓨팅이라고 하며, GPU는 범용 과학 컴퓨팅을 위한 보조 프로세서를 범용 과학 컴퓨팅이라고 하며 GPU (GPGPU) 프로그래밍

GPU 컴퓨팅은 최근 머신러닝 호황에 크게 기여했습니다. 왜냐하면 컨볼루션 신경망과 기타 모델에서 더 효율적으로 실행되도록 설계되었습니다 현재의 웹 플랫폼에서 W3C의 '웹용 GPU'는 GPU 컴퓨팅 기능이 없는 커뮤니티 그룹 대부분의 컴퓨팅 환경에서 사용할 수 있는 최신 GPU API를 노출하는 있습니다. 이 API를 WebGPU라고 합니다.

WebGPU는 WebGL과 같은 하위 수준 API입니다. 이것은 매우 강력하고 장황합니다. 확인할 수 있습니다. 하지만 괜찮습니다. 우리가 바라는 것은 바로 성능입니다.

이 문서에서는 WebGPU의 GPU 컴퓨팅 부분에 중점을 두고 솔직히 말씀드리자면 그냥 겉핥기식으로 있습니다. 앞으로 WebGPU 렌더링 (캔버스, 텍스처, 등)를 포함해야 합니다.

GPU 액세스

WebGPU에서는 GPU에 쉽게 액세스할 수 있습니다. navigator.gpu.requestAdapter()님에게 전화 거는 중 GPU와 비동기식으로 확인되는 JavaScript 프로미스를 반환합니다. 어댑터에 연결합니다. 이 어댑터를 그래픽 카드라고 생각하세요. GKE는 클러스터 자체를 (CPU와 동일한 칩에 있음) 또는 별개의 장치인 성능이 더 우수하지만 전력 소비량이 더 많음).

GPU 어댑터가 있으면 adapter.requestDevice()를 호출하여 프로미스를 가져옵니다. GPU 계산을 수행하는 데 사용할 GPU 장치로 확인될 것입니다.

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();

두 함수 모두 원하는 기기(확장 프로그램, 한도)를 설정할 수 있습니다. 대상 이 도움말에서는 편의상 기본 옵션을 사용합니다.

버퍼 메모리 쓰기

JavaScript를 사용하여 GPU의 메모리에 데이터를 쓰는 방법을 살펴보겠습니다. 이 최신 웹에서 사용되는 샌드박스 모델 때문에 프로세스가 간단하지 않습니다. 있습니다.

아래 예는 액세스 가능한 메모리에 4바이트를 쓰는 방법을 보여줍니다. GPU에서 추출해야 합니다. device.createBuffer()를 호출하여 버퍼 및 그 사용을 나타냅니다. 사용 플래그 GPUBufferUsage.MAP_WRITE가 이 특정 호출에는 필요하지 않으므로, 이 버퍼에 전달합니다. 그 결과, GPU 버퍼 객체가 생성될 때 mappedAtCreation를 true로 설정합니다. 그런 다음 연결된 원시 바이너리 데이터 버퍼는 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가 동시에 사용할 수 있습니다

버퍼 메모리 읽기

이제 GPU 버퍼를 다른 GPU 버퍼에 복사하고 다시 읽는 방법을 살펴보겠습니다.

첫 번째 GPU 버퍼에 쓰기 때문에 두 번째 GPU 버퍼에 GPU 버퍼를 사용하려면 새 사용 플래그 GPUBufferUsage.COPY_SRC가 필요합니다. 두 번째 이번에는 GPU 버퍼가 매핑되지 않은 상태로 생성됩니다. device.createBuffer() 사용 플래그는 첫 번째 GPU의 대상으로 사용되므로 GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ입니다. GPU 복사 명령이 실행되면 JavaScript에서 버퍼링 및 읽기를 수행합니다.

// 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에서 GPU 명령어 인코더는 device.createCommandEncoder()는 배치를 빌드하는 JavaScript 객체입니다. '버퍼링됨' 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 버퍼를 읽으려면 다음을 사용하여 gpuReadBuffer.mapAsync()를 호출합니다. GPUMapMode.READ GPU 버퍼가 해제될 때 확인되는 프로미스를 반환합니다. 있습니다. 그런 다음 gpuReadBuffer.getMappedRange()를 사용하여 매핑된 범위를 가져옵니다. 큐에 추가된 모든 GPU 명령어가 있으면 첫 번째 GPU 버퍼와 동일한 값을 포함합니다. 확인할 수 있습니다

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));

이 샘플을 사용해 볼 수 있습니다.

간단히 말해, 버퍼 메모리 작업과 관련하여 기억해야 할 사항은 다음과 같습니다.

  • GPU 버퍼는 기기 대기열 제출에 사용하려면 매핑을 해제해야 합니다.
  • 매핑되면 GPU 버퍼를 JavaScript로 읽고 쓸 수 있습니다.
  • GPU 버퍼는 mapAsync()createBuffer()이 true로 설정된 mappedAtCreation가 호출됩니다.

셰이더 프로그래밍

계산만 수행하고 GPU 그리기는 수행하지 않는 GPU에서 실행 중인 프로그램 삼각형)을 컴퓨팅 셰이더라고 합니다. 이러한 함수는 수백 개의 크런치를 위해 함께 작동하는 GPU 코어 (CPU 코어보다 작음)의 수 데이터를 수집하는 데 사용됩니다 입력 및 출력은 WebGPU의 버퍼입니다.

WebGPU에서 컴퓨팅 셰이더를 사용하는 방법을 설명하기 위해 행렬을 사용해 보겠습니다. 곱셈의 공식이 사용됩니다.

<ph type="x-smartling-placeholder">
</ph> 행렬 곱셈 다이어그램
행렬 곱셈 다이어그램

간단히 말하면 다음과 같습니다.

  1. 3개의 GPU 버퍼 (행렬 곱셈용 2개와 결과 행렬)
  2. 컴퓨팅 셰이더의 입력 및 출력 설명
  3. 컴퓨팅 셰이더 코드 컴파일
  4. 컴퓨팅 파이프라인 설정
  5. 인코딩된 명령어를 GPU에 일괄적으로 제출
  6. 결과 행렬 GPU 버퍼 읽기

GPU 버퍼 생성

편의상 행렬은 부동 소수점 수 목록으로 표현됩니다. 사용할 수 없습니다. 첫 번째 요소는 행 개수이고 두 번째 요소는 나머지는 행렬의 실제 숫자입니다.

<ph type="x-smartling-placeholder">
</ph> JavaScript로 행렬을 간단하게 표현하고 이에 상응하는 수학 표기법으로 표현
JavaScript로 행렬을 간단하게 표현하고 이에 상응하는 수학 표기법으로 표현

3개의 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에만 적용됩니다. 바인드 그룹 레이아웃은 셰이더에서 예상하는 입력/출력 인터페이스를 정의하고 바인드 그룹은 셰이더의 실제 입력/출력 데이터를 나타냅니다.

아래 예에서 바인드 그룹 레이아웃은 컴퓨팅 셰이더의 2에 번호가 매겨진 항목 바인딩 0, 1, 스토리지 버퍼. 반면에 이 바인드 그룹 레이아웃에 대해 정의된 바인드 그룹은 GPU는 결합 0의 항목(gpuBufferFirstMatrix)에 버퍼링됩니다. 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
      }
    }
  ]
});

컴퓨팅 셰이더 코드

행렬을 곱하기 위한 컴퓨팅 셰이더 코드는 WGSL로 작성됩니다. SPIR-V로 쉽게 번역할 수 있는 WebGPU 셰이더 언어입니다. 제외 자세히 살펴보면, 아래에 확인된 3개의 스토리지 버퍼는 (var<storage> 포함) 프로그램에서는 다음과 같이 firstMatrixsecondMatrix를 사용합니다. 입력 및 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 버퍼 3개와 컴퓨팅 1개로 바인드 그룹을 인스턴스화한 후 바인드 그룹 레이아웃이 있는 경우 이를 사용해야 합니다.

다음 명령어로 프로그래밍 가능한 컴퓨팅 패스 인코더를 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 컴퓨팅 환경에서는 데이터 세트에서 커널 함수를 실행하기 위해 명령어를 인코딩하는 것을 디스패칭이라고 합니다.

<ph type="x-smartling-placeholder">
</ph> 각 결과 매트릭스 셀에 대해 병렬 실행
각 결과 행렬 셀에 대해 병렬 실행

컴퓨팅 셰이더의 작업 그룹 그리드 크기는 WGSL에서 (8, 8)입니다. 있습니다. 따라서 'x'는 'y' 각 열의 행 수와 첫 번째 행렬과 두 번째 행렬의 열 개수를 있습니다. 이를 통해 이제 다음을 사용하여 컴퓨팅 호출을 발송할 수 있습니다. 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()를 호출합니다. 그런 다음 결과 행렬 버퍼를 복사하는 대상으로 사용할 GPU 버퍼입니다. copyBufferToBuffer 마지막으로 copyEncoder.finish()를 호출하고 다음 호출을 통해 GPU 기기 큐에 제출합니다. GPU 명령어가 포함된 device.queue.submit()

// 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]);

결과 행렬 읽기

결과 행렬을 읽는 것은 다음과 같이 gpuReadBuffer.mapAsync()를 호출하는 것만큼 간단합니다. GPUMapMode.READ이고 반환되는 프라미스가 해결을 기다리는 중입니다. 이제 GPU 버퍼가 매핑됩니다. 이 시점에서 이 레이어의 모든 매핑된 gpuReadBuffer.getMappedRange()로 범위를 지정합니다.

<ph type="x-smartling-placeholder">
</ph> 행렬 곱셈 결과
행렬 곱셈 결과

코드에서 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에 대해 설명한 프로그램을 작성했습니다. 그리고 최대한 GPU를 최대한 활용하는 것은 당연한 선택처럼 보입니다 행렬의 크기가 256x256보다 크면 이 값을 계산합니다.

<ph type="x-smartling-placeholder">
</ph> GPU vs CPU 벤치마크
GPU 대 CPU 벤치마크

이 문서는 WebGPU 탐색 여정의 시작에 불과합니다. 기대되는 성과 GPU 컴퓨팅에 대한 심층 분석과 렌더링 방법에 대해 (캔버스, 텍스처, 샘플러)는 WebGPU에서 작동합니다.