این صفحه به‌وسیله ‏Cloud Translation API‏ ترجمه شده است.

با GPU Compute در وب شروع کنید

این پست API آزمایشی WebGPU را از طریق مثال‌ها بررسی می‌کند و به شما کمک می‌کند تا با استفاده از GPU محاسبات موازی داده را شروع کنید.

François Beaufort

پس زمینه

همانطور که قبلاً می دانید، واحد پردازش گرافیکی (GPU) یک زیر سیستم الکترونیکی در یک رایانه است که در ابتدا برای پردازش گرافیک تخصصی بود. با این حال، در 10 سال گذشته، به سمت یک معماری انعطاف‌پذیرتر تکامل یافته است که به توسعه‌دهندگان اجازه می‌دهد تا انواع بسیاری از الگوریتم‌ها را پیاده‌سازی کنند، نه فقط گرافیک‌های سه‌بعدی را ارائه دهند، در حالی که از معماری منحصربه‌فرد GPU بهره می‌برند. این قابلیت ها به عنوان GPU Compute نامیده می شوند و استفاده از یک GPU به عنوان یک پردازنده مشترک برای محاسبات علمی همه منظوره برنامه نویسی GPU (GPGPU) همه منظوره نامیده می شود.

GPU Compute سهم قابل توجهی در رونق اخیر یادگیری ماشین داشته است، زیرا شبکه‌های عصبی کانولوشن و سایر مدل‌ها می‌توانند از مزیت معماری برای اجرای کارآمدتر بر روی GPU استفاده کنند. با توجه به اینکه پلتفرم وب فعلی فاقد قابلیت‌های GPU Compute است، گروه اجتماعی "GPU برای وب" W3C در حال طراحی یک API برای افشای APIهای GPU مدرن است که در اکثر دستگاه‌های فعلی موجود است. این API WebGPU نام دارد.

WebGPU یک API سطح پایین است، مانند WebGL. همانطور که خواهید دید بسیار قدرتمند و کاملاً پرحرف است. اما این اشکالی ندارد. آنچه ما به دنبال آن هستیم عملکرد است.

در این مقاله، من قصد دارم بر روی بخش GPU Compute WebGPU تمرکز کنم و صادقانه بگویم، من فقط سطح را خراش می‌دهم تا بتوانید به تنهایی شروع به بازی کنید. در مقاله‌های آینده عمیق‌تر غواصی خواهم کرد و رندر WebGPU (بوم، بافت و غیره) را پوشش خواهم داد.

به GPU دسترسی داشته باشید

دسترسی به GPU در WebGPU آسان است. فراخوانی 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 باشد. دستگاه 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 است و در خواندن/نوشتن از جاوا اسکریپت قابل دسترسی است. برای اینکه GPU بتواند به آن دسترسی داشته باشد، باید از نقشه برداری خارج شود که به سادگی فراخوانی gpuBuffer.unmap() است.

مفهوم نگاشت/نقشه برداری نشده برای جلوگیری از شرایط مسابقه که در آن GPU و CPU همزمان به حافظه دسترسی دارند، مورد نیاز است.

خواندن حافظه بافر

حالا بیایید ببینیم چگونه یک بافر GPU را در بافر GPU دیگری کپی کنیم و آن را دوباره بخوانیم.

از آنجایی که ما در اولین بافر GPU می نویسیم و می خواهیم آن را در بافر GPU دوم کپی کنیم، یک پرچم استفاده جدید GPUBufferUsage.COPY_SRC مورد نیاز است. بافر دوم GPU در حالت بدون نقشه این بار با device.createBuffer() ایجاد می شود. پرچم استفاده آن GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ به عنوان مقصد اولین بافر GPU استفاده می شود و پس از اجرای دستورات کپی GPU در جاوا اسکریپت خوانده می شود.

// 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() برگردانده می شود، شی جاوا اسکریپت است که دسته ای از دستورات "بافر" را می سازد که در نقطه ای به GPU ارسال می شود. از سوی دیگر، متدهای موجود در GPUBuffer "unbuffered" هستند، به این معنی که در زمانی که فراخوانی می شوند به صورت اتمی اجرا می شوند.

هنگامی که رمزگذار دستور GPU را دارید، مانند تصویر زیر copyEncoder.copyBufferToBuffer() را فراخوانی کنید تا این دستور را برای اجرای بعدی به صف فرمان اضافه کنید. در نهایت، دستورات رمزگذاری را با فراخوانی copyEncoder.finish() به پایان برسانید و آنها را به صف فرمان دستگاه GPU ارسال کنید. صف مسئول رسیدگی به ارسال های انجام شده از طریق device.queue.submit() با دستورات GPU به عنوان آرگومان است. این کار تمام دستورات ذخیره شده در آرایه را به ترتیب به صورت اتمی اجرا می کند.

// 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 را می توان در جاوا اسکریپت خواند و نوشت.
زمانی که mapAsync() و createBuffer() با mappedAtCreation روی true فراخوانی می شوند، بافرهای GPU نگاشت می شوند.

برنامه نویسی شیدر

برنامه‌هایی که روی GPU اجرا می‌شوند و فقط محاسبات را انجام می‌دهند (و مثلث نمی‌کشند) شیدر محاسباتی نامیده می‌شوند. آنها به طور موازی توسط صدها هسته GPU (که کوچکتر از هسته های CPU هستند) اجرا می شوند که با هم کار می کنند تا داده ها را خرد کنند. ورودی و خروجی آنها بافرهایی در WebGPU هستند.

برای نشان دادن استفاده از شیدرهای محاسباتی در WebGPU، ما با ضرب ماتریس، یک الگوریتم رایج در یادگیری ماشین که در زیر نشان داده شده است، بازی خواهیم کرد.

به طور خلاصه، این چیزی است که ما می خواهیم انجام دهیم:

سه بافر GPU ایجاد کنید (دو تا برای ضرب ماتریس ها و یکی برای ماتریس نتیجه)
ورودی و خروجی را برای شیدر محاسباتی توضیح دهید
کد شیدر محاسباتی را کامپایل کنید
یک خط لوله محاسباتی راه اندازی کنید
دستورات کدگذاری شده را به صورت دسته ای به 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
});

طرح بندی گروه و گروه صحافی

مفاهیم چیدمان گروه bind و گروه bind مختص WebGPU هستند. طرح بندی گروه bind رابط ورودی/خروجی مورد انتظار یک سایه زن را تعریف می کند، در حالی که یک گروه bind داده های ورودی/خروجی واقعی یک سایه زن را نشان می دهد.

در مثال زیر، طرح‌بندی گروه bind دو بافر ذخیره‌سازی فقط خواندنی در اتصال‌های ورودی شماره‌دار 0 ، 1 و یک بافر ذخیره‌سازی در 2 برای سایه‌زن محاسباتی انتظار دارد. از طرف دیگر، گروه bind که برای این طرح‌بندی گروه bind تعریف شده است، بافرهای GPU را به ورودی‌ها مرتبط می‌کند: gpuBufferFirstMatrix به binding 0 ، gpuBufferSecondMatrix به binding 1 و resultMatrixBuffer به binding 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 نوشته شده است، WebGPU Shader Language، که به طور ساده به SPIR-V قابل ترجمه است. بدون پرداختن به جزئیات، باید در زیر سه بافر ذخیره‌سازی شناسایی شده با var<storage> را بیابید. این برنامه از firstMatrix و secondMatrix به عنوان ورودی و resultMatrix به عنوان خروجی استفاده خواهد کرد.

توجه داشته باشید که هر بافر ذخیره‌سازی دارای تزئینات binding است که با همان شاخصی که در طرح‌بندی‌های گروه bind و گروه‌های bind اعلام شده در بالا تعریف شده است، مطابقت دارد.

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() آن را ایجاد کنید. دو آرگومان نیاز دارد: طرح‌بندی گروه bind که قبلاً ایجاد کردیم، و یک مرحله محاسباتی که نقطه ورودی شیدر محاسباتی ما (عملکرد main WGSL) را تعریف می‌کند و ماژول سایه‌زن محاسباتی واقعی که با device.createShaderModule() ایجاد شده است.

const computePipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module: shaderModule,
    entryPoint: "main"
  }
});

ارسال دستورات

پس از نمونه سازی یک گروه bind با سه بافر GPU و یک خط لوله محاسباتی با طرح بندی گروه bind، زمان استفاده از آنها فرا رسیده است.

بیایید یک رمزگذار محاسباتی قابل برنامه ریزی را با commandEncoder.beginComputePass() شروع کنیم. ما از این برای رمزگذاری دستورات GPU استفاده می کنیم که ضرب ماتریس را انجام می دهند. خط لوله آن را با passEncoder.setPipeline(computePipeline) و گروه bind آن را با passEncoder.setBindGroup(0, bindGroup) در اندیس 0 تنظیم کنید. شاخص 0 مربوط به دکوراسیون group(0) در کد WGSL است.

حالا بیایید در مورد نحوه اجرای این شیدر محاسباتی بر روی GPU صحبت کنیم. هدف ما این است که این برنامه را به صورت موازی برای هر سلول از ماتریس نتیجه، مرحله به مرحله اجرا کنیم. به عنوان مثال، برای یک ماتریس نتیجه با اندازه 16 در 32، برای رمزگذاری دستور اجرا، در @workgroup_size(8, 8) ، passEncoder.dispatchWorkgroups(2, 4) یا passEncoder.dispatchWorkgroups(16 / 8, 32 / 8) را فراخوانی می کنیم. passEncoder.dispatchWorkgroups(16 / 8, 32 / 8) . آرگومان اول "x" بعد اول است، آرگومان دوم "y" بعد دوم است و آخرین آرگومان "z" بعد سوم است که به طور پیش فرض 1 است زیرا در اینجا به آن نیاز نداریم. در دنیای محاسبات GPU، رمزگذاری دستوری برای اجرای یک تابع هسته بر روی مجموعه ای از داده ها، Dispatching نامیده می شود.

اندازه شبکه گروه کاری برای شیدر محاسباتی ما (8, 8) در کد WGSL ما است. به همین دلیل، "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() را فراخوانی کنید. سپس، یک بافر GPU ایجاد کنید تا از آن به عنوان مقصد برای کپی کردن بافر ماتریس نتیجه با copyBufferToBuffer استفاده کنید. در نهایت، کدگذاری دستورات را با copyEncoder.finish() تمام کنید و با فراخوانی device.queue.submit() با دستورات GPU، آن ها را به صف دستگاه 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]);

ماتریس نتیجه را بخوانید

خواندن ماتریس نتیجه به آسانی فراخوانی gpuReadBuffer.mapAsync() با GPUMapMode.READ و انتظار برای حل شدن وعده بازگشتی است که نشان می دهد بافر GPU اکنون نقشه برداری شده است. در این مرحله، می توان محدوده نگاشت شده را با gpuReadBuffer.getMappedRange() دریافت کرد.

در کد ما، نتیجه وارد شده در کنسول جاوا اسکریپت DevTools "2، 2، 50، 60، 114، 140" است.

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

تبریک می گویم! تو درستش کردی می توانید با نمونه بازی کنید .

آخرین ترفند

یکی از راه‌های ساده‌تر خواندن کد، استفاده از روش کاربردی getBindGroupLayout خط لوله محاسباتی برای استنتاج طرح‌بندی گروه bind از ماژول سایه‌زن است. این ترفند نیاز به ایجاد یک طرح بندی گروهی باند سفارشی و تعیین طرح خط لوله در خط لوله محاسباتی شما را برطرف می کند، همانطور که در زیر می بینید.

تصویری از 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 زمانی که اندازه ماتریس ها بزرگتر از 256 در 256 باشد یک انتخاب واضح به نظر می رسد.

این مقاله تنها آغاز سفر من برای کاوش در WebGPU بود. به زودی منتظر مقالات بیشتری باشید که شامل غواصی عمیق تر در GPU Compute و نحوه عملکرد رندر (بوم، بافت، نمونه) در WebGPU می شود.

François Beaufort

پس زمینه

به GPU دسترسی داشته باشید

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

حافظه بافر بنویسید

// 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 را در بافر GPU دیگری کپی کنیم و آن را دوباره بخوانیم.

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

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

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

می توانید این نمونه را امتحان کنید .

به طور خلاصه، در اینجا چیزی است که باید در مورد عملیات حافظه بافر به خاطر بسپارید:

بافرهای GPU باید نقشه برداری نشده باشند تا در ارسال صف دستگاه استفاده شوند.
هنگام نقشه برداری، بافرهای GPU را می توان در جاوا اسکریپت خواند و نوشت.
زمانی که mapAsync() و createBuffer() با mappedAtCreation روی true فراخوانی می شوند، بافرهای GPU نگاشت می شوند.

برنامه نویسی شیدر

به طور خلاصه، این چیزی است که ما می خواهیم انجام دهیم:

سه بافر GPU ایجاد کنید (دو تا برای ضرب ماتریس ها و یکی برای ماتریس نتیجه)
ورودی و خروجی را برای شیدر محاسباتی توضیح دهید
کد شیدر محاسباتی را کامپایل کنید
یک خط لوله محاسباتی راه اندازی کنید
دستورات کدگذاری شده را به صورت دسته ای به GPU ارسال کنید
بافر GPU ماتریس نتیجه را بخوانید

ایجاد بافرهای 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
});

طرح بندی گروه و گروه صحافی

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

محاسبه کد سایه زن

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;
    }
  `
});

راه اندازی خط لوله

const computePipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module: shaderModule,
    entryPoint: "main"
  }
});

ارسال دستورات

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();

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

ماتریس نتیجه را بخوانید

در کد ما، نتیجه وارد شده در کنسول جاوا اسکریپت DevTools "2، 2، 50، 60، 114، 140" است.

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

تبریک می گویم! تو درستش کردی می توانید با نمونه بازی کنید .

آخرین ترفند

تصویری از 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: [

یافته های عملکرد

François Beaufort

پس زمینه

به GPU دسترسی داشته باشید

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

حافظه بافر بنویسید

// 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 را در بافر GPU دیگری کپی کنیم و آن را دوباره بخوانیم.

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

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

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

می توانید این نمونه را امتحان کنید .

به طور خلاصه، در اینجا چیزی است که باید در مورد عملیات حافظه بافر به خاطر بسپارید:

بافرهای GPU باید نقشه برداری نشده باشند تا در ارسال صف دستگاه استفاده شوند.
هنگام نقشه برداری، بافرهای GPU را می توان در جاوا اسکریپت خواند و نوشت.
زمانی که mapAsync() و createBuffer() با mappedAtCreation روی true فراخوانی می شوند، بافرهای GPU نگاشت می شوند.

برنامه نویسی شیدر

به طور خلاصه، این چیزی است که ما می خواهیم انجام دهیم:

سه بافر GPU ایجاد کنید (دو تا برای ضرب ماتریس ها و یکی برای ماتریس نتیجه)
ورودی و خروجی را برای شیدر محاسباتی توضیح دهید
کد شیدر محاسباتی را کامپایل کنید
یک خط لوله محاسباتی راه اندازی کنید
دستورات کدگذاری شده را به صورت دسته ای به GPU ارسال کنید
بافر GPU ماتریس نتیجه را بخوانید

ایجاد بافرهای 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
});

طرح بندی گروه و گروه صحافی

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

محاسبه کد سایه زن

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;
    }
  `
});

راه اندازی خط لوله

const computePipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module: shaderModule,
    entryPoint: "main"
  }
});

ارسال دستورات

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();

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

ماتریس نتیجه را بخوانید

در کد ما، نتیجه وارد شده در کنسول جاوا اسکریپت DevTools "2، 2، 50، 60، 114، 140" است.

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

تبریک می گویم! تو درستش کردی می توانید با نمونه بازی کنید .

آخرین ترفند

تصویری از 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: [