โพสต์นี้จะสำรวจ WebGPU API รุ่นทดลองผ่านตัวอย่างต่างๆ และช่วยให้คุณเริ่มต้นคำนวณคู่ขนานข้อมูลโดยใช้ GPU
ที่มา
อย่างที่คุณอาจทราบอยู่แล้วว่าหน่วยประมวลผลกราฟิก (GPU) เป็นระบบย่อยอิเล็กทรอนิกส์ภายในคอมพิวเตอร์ที่แต่เดิมมีความเชี่ยวชาญด้านการประมวลผลกราฟิก อย่างไรก็ตาม ในช่วง 10 ปีที่ผ่านมา การทดสอบได้พัฒนาไปเป็นสถาปัตยกรรมที่ยืดหยุ่นมากขึ้นทำให้นักพัฒนาซอฟต์แวร์สามารถนำอัลกอริทึมหลายประเภทไปใช้ ซึ่งไม่ใช่แค่แสดงผลกราฟิก 3 มิติเท่านั้น แต่ยังใช้ประโยชน์จากสถาปัตยกรรมที่มีเอกลักษณ์ของ 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() จะส่งกลับคำสัญญา JavaScript ที่จะแก้ไขแบบอะซิงโครนัสด้วยอะแดปเตอร์ GPU อะแดปเตอร์นี้เป็นการ์ดแสดงผล โดยสามารถผสานรวม (บนชิปเดียวกันกับ CPU) หรือแบบแยกต่างหาก (โดยปกติจะเป็นการ์ด PCIe ที่มีประสิทธิภาพสูงกว่าแต่ใช้พลังงานมากกว่า)
เมื่อคุณมีอะแดปเตอร์ 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() ซึ่งจะใช้ขนาดของบัฟเฟอร์และการใช้งาน แม้ว่าไม่จำเป็นต้องระบุ Flag การใช้งาน GPUBufferUsage.MAP_WRITE สำหรับการเรียกนี้ โปรดระบุอย่างชัดเจนว่าเราต้องการเขียนลงในบัฟเฟอร์นี้ ซึ่งส่งผลให้มีการแมปออบเจ็กต์บัฟเฟอร์ GPU ขณะสร้างด้วยการตั้งค่า mappedAtCreation เป็น "จริง" จากนั้น คุณจะเรียกบัฟเฟอร์ข้อมูลไบนารีดิบที่เกี่ยวข้องได้โดยเรียกใช้เมธอดบัฟเฟอร์ 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 เข้าถึง GPU ได้ คุณต้องยกเลิกการแมป GPU ซึ่งง่ายมาก เพียงแค่เรียกใช้ gpuBuffer.unmap()
แนวคิดของการแมป/ไม่ได้แมปนั้นจำเป็นต่อการป้องกันเงื่อนไขการแข่งขันที่ GPU และ CPU เข้าถึงหน่วยความจำพร้อมกัน
อ่านหน่วยความจำบัฟเฟอร์
มาดูวิธีคัดลอกบัฟเฟอร์ GPU ไปยังบัฟเฟอร์ GPU อื่นและอ่านอีกครั้ง
เนื่องจากเรากำลังเขียนในบัฟเฟอร์ GPU แรกและต้องการคัดลอกไปยังบัฟเฟอร์ GPU ที่สอง จึงจำเป็นต้องมี GPUBufferUsage.COPY_SRC แฟล็กการใช้งานใหม่ บัฟเฟอร์ GPU ที่สองจะสร้างขึ้นในสถานะที่ไม่ได้แมปด้วย device.createBuffer() ในครั้งนี้ แฟล็กการใช้งานคือ GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ เนื่องจากระบบจะใช้เป็นปลายทางของบัฟเฟอร์ GPU แรกและอ่านใน JavaScript เมื่อเรียกใช้คำสั่งการคัดลอก 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() คือออบเจ็กต์ JavaScript ที่สร้างกลุ่มคำสั่ง "บัฟเฟอร์" ซึ่งจะส่งไปยัง GPU ในบางจุด ในทางตรงกันข้าม เมธอดใน GPUBuffer จะ "ไม่บัฟเฟอร์" ซึ่งหมายความว่าเมธอดจะทำงานแบบอะตอมในเวลาที่มีการเรียกใช้
เมื่อคุณมีตัวเข้ารหัสคำสั่ง 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 รายการที่ 2 ให้เรียกใช้ 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()ที่ตั้งค่าmappedAtCreationเป็น "จริง"
โปรแกรมให้เฉดสี
โปรแกรมที่ทำงานบน GPU ซึ่งทำงานด้านการคำนวณเท่านั้น (และไม่ได้วาดรูปสามเหลี่ยม) เรียกว่าตัวปรับเฉดสีสำหรับการประมวลผล ตัวประมวลผลเหล่านี้ทำงานแบบคู่ขนานด้วยแกน GPU (ซึ่งมีขนาดเล็กกว่าแกนของ CPU) หลายร้อยแกนที่ทำงานร่วมกันเพื่อประมวลผลข้อมูล อินพุตและเอาต์พุตเป็นบัฟเฟอร์ใน WebGPU
มาดูตัวอย่างการใช้ตัวปรับแสงเงาจากการประมวลผลใน WebGPU กัน โดยจะมาเล่นการคูณเมทริกซ์ซึ่งเป็นอัลกอริทึมทั่วไปในแมชชีนเลิร์นนิงตามที่แสดงไว้ด้านล่าง
กล่าวโดยสรุปคือสิ่งที่เราจะทำมีดังนี้
- สร้างบัฟเฟอร์ GPU 3 รายการ (2 รายการสำหรับเมทริกซ์เพื่อคูณ และอีก 1 รายการสำหรับเมทริกซ์ผลลัพธ์)
- อธิบายอินพุตและเอาต์พุตสำหรับตัวปรับแสงเงาในการประมวลผล
- คอมไพล์โค้ดตัวปรับแสงเงาการประมวลผล
- ตั้งค่าไปป์ไลน์การประมวลผล
- ส่งคำสั่งที่เข้ารหัสเป็นชุดไปยัง GPU
- อ่านบัฟเฟอร์ GPU ของเมทริกซ์ผลลัพธ์
การสร้างบัฟเฟอร์ GPU
เพื่อความสะดวก เมทริกซ์จะแสดงเป็นรายการจำนวนจุดลอยตัว องค์ประกอบแรกคือจำนวนแถว องค์ประกอบที่ 2 คือจำนวนคอลัมน์ และองค์ประกอบที่เหลือคือจำนวนจริงของเมทริกซ์
บัฟเฟอร์ GPU ทั้ง 3 ตัวเป็นบัฟเฟอร์พื้นที่เก็บข้อมูลที่เราต้องการจัดเก็บและเรียกข้อมูลในตัวสร้างเงาการประมวลผล นี่เป็นคำอธิบายเหตุผลที่แฟล็กการใช้งานบัฟเฟอร์ 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 โดยเฉพาะ เลย์เอาต์กลุ่มการเชื่อมโยงจะกำหนดอินเทอร์เฟซอินพุต/เอาต์พุตที่คาดไว้โดยตัวปรับแสงเงา ขณะที่กลุ่มการเชื่อมโยงจะแสดงข้อมูลอินพุต/เอาต์พุตจริงสำหรับตัวปรับแสงเงา
ในตัวอย่างด้านล่าง เลย์เอาต์ของ Bind Group คาดหวังบัฟเฟอร์พื้นที่เก็บข้อมูลแบบอ่านอย่างเดียว 2 รายการในการเชื่อมโยงรายการตัวเลข 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
}
}
]
});
โค้ดตัวปรับแสงเงาในการประมวลผล
โค้ดตัวปรับแสงเงาการประมวลผลสำหรับเมทริกซ์การคูณเขียนด้วย WGSL ซึ่งเป็นภาษา Shader ของ WebGPU โดยแปลเป็น SPIR-V ได้ 3 รายการ หากไม่ลงรายละเอียด คุณจะเห็นบัฟเฟอร์พื้นที่เก็บข้อมูล 3 รายการที่ระบุด้วย 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()
ใช้อาร์กิวเมนต์ 2 ตัว ได้แก่ เลย์เอาต์กลุ่มการเชื่อมโยงที่เราสร้างไว้ก่อนหน้านี้ และขั้นตอนการประมวลผลที่กำหนดจุดแรกเข้าของตัวสร้างเงาการประมวลผล (main ฟังก์ชัน WGSL) และโมดูลตัวปรับเงาการประมวลผลจริงที่สร้างด้วย device.createShaderModule()
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
การส่งคำสั่ง
หลังจากสร้างกลุ่มการเชื่อมโยงด้วยบัฟเฟอร์ GPU 3 รายการและไปป์ไลน์การประมวลผลที่มีเลย์เอาต์กลุ่มการเชื่อมโยงแล้ว ก็ถึงเวลาใช้งานกลุ่มเหล่านั้น
มาเริ่มโปรแกรมเปลี่ยนไฟล์ที่ตั้งโปรแกรมได้สำหรับการประมวลผลที่มีโปรแกรมด้วย commandEncoder.beginComputePass() กัน เราจะใช้โค้ดนี้เข้ารหัสคำสั่ง GPU
ที่จะทำการคำนวณการคูณเมทริกซ์ ตั้งค่าไปป์ไลน์ด้วย passEncoder.setPipeline(computePipeline) และกลุ่มการเชื่อมโยงของดัชนีที่ดัชนี 0 ด้วย passEncoder.setBindGroup(0, bindGroup) ดัชนี 0 สอดคล้องกับการตกแต่ง group(0) ในโค้ด WGSL
คราวนี้เราจะมาพูดคุยว่าตัวปรับแสงเงาการประมวลผลนี้จะทำงานบน GPU อย่างไร เป้าหมายของเราคือการใช้โปรแกรมนี้พร้อมกันกับเมทริกซ์ผลลัพธ์แต่ละเซลล์ตามลำดับ ตัวอย่างเช่น สำหรับเมทริกซ์ผลลัพธ์ขนาด 16 x 32 หากต้องการเข้ารหัสคำสั่งการดำเนินการใน @workgroup_size(8, 8) เราจะเรียกใช้ passEncoder.dispatchWorkgroups(2, 4) หรือ passEncoder.dispatchWorkgroups(16 / 8, 32 / 8)
อาร์กิวเมนต์แรก "x" คือมิติข้อมูลแรก มิติข้อมูลที่ 2 "y" คือมิติข้อมูลที่ 2 และมิติข้อมูลล่าสุด "z" คือมิติข้อมูลที่ 3 ซึ่งมีค่าเริ่มต้นเป็น 1 เนื่องจากเราไม่จำเป็นต้องใช้
ในโลกการประมวลผลของ GPU การเข้ารหัสคำสั่งเพื่อเรียกใช้ฟังก์ชันเคอร์เนลในชุดข้อมูลเรียกว่าการจ่ายงาน
ขนาดของตารางกริดกลุ่มงานสำหรับตัวปรับแสงเงาการประมวลผลคือ (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();
หากต้องการสิ้นสุดการใช้โปรแกรมเปลี่ยนไฟล์ Compute Pass ให้โทรหา passEncoder.end() จากนั้นสร้างบัฟเฟอร์ GPU เพื่อใช้เป็นปลายทางในการคัดลอกบัฟเฟอร์เมทริกซ์ผลลัพธ์ด้วย copyBufferToBuffer สุดท้าย ใช้คำสั่งการเข้ารหัสให้เสร็จสิ้นด้วย 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]);
อ่านเมทริกซ์ผลลัพธ์
การอ่านเมทริกซ์ผลลัพธ์นั้นง่ายพอๆ กับการเรียกใช้ gpuReadBuffer.mapAsync() ด้วย GPUMapMode.READ และรอให้สัญญาที่ส่งคืนมาแก้ไขปัญหา ซึ่งระบุว่าตอนนี้บัฟเฟอร์ GPU แล้ว ในจุดนี้ คุณจะได้ช่วงที่แมปด้วย gpuReadBuffer.getMappedRange()
ในโค้ด ผลลัพธ์ที่เข้าสู่ระบบคอนโซล JavaScript ของ DevTools คือ "2, 2, 50, 60, 114, 140"
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));
ยินดีด้วย คุณทำสำเร็จแล้ว ก็ลองเล่นตัวอย่างได้
เคล็ดลับสุดท้าย
วิธีหนึ่งที่จะทำให้โค้ดอ่านง่ายขึ้นคือการใช้เมธอด 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 ได้อย่างเต็มที่จะดูเหมือนว่าเป็นตัวเลือกที่ชัดเจนเมื่อขนาดของเมทริกซ์มากกว่า 256 x 256
บทความนี้เป็นเพียงจุดเริ่มต้นในเส้นทางการสำรวจ WebGPU ของผมเท่านั้น พบกับบทความอื่นๆ ในเร็วๆ นี้ จะมีการเจาะลึกรายละเอียดมากขึ้นใน GPU Compute และวิธีการทำงานของการแสดงผล (Canvas, พื้นผิว, Sampler) ใน WebGPU