WebGPU’da hisoblash asoslari
Muhammadxon NajimovUshbu maqola tajribali dasturchilar uchun GPU’da hisoblashni amalga bo’yicha boshlang’ich ma’lumotlarni taqdim etadi. Agar siz dasturlashni endi o’rganayotgan bo’lsangiz hechqisi yo’q. Faqat Arrays, Loops, Conditions, Functions mavzularini bilishingizga ishonch hosil qiling.
Bor diqqatingizni hisoblashga qaratish maqsadida, men WebGPU nima ekanligi va yoki uning tarixini muhokama qilmayman. Ammo men to’liqroq bilishingiz kerak deb hisoblagan har qanday narsa uchun albatta havola qoldirishga harakat qilaman. Siz ular bilan maqolani to’liq o’qib bo’lganingizdan keyin ham tanishib chiqishingiz mumkin. Shuningdek men taqdim qilmoqchi bo’lgan kod sodda bo’lishi uchun brauzeringiz WebGPU-ni qo‘llab-quvvatlashini tekshirish yoki shu kabi sodda qismlarni o‘tkazib yuboraman.
Men qanday aniq vazifani o’z zimmamga olaman?
Men sizga qandaydir murakkab muammoga yechim berish haqida o’ylamadim. Chunki siz avval nimani hisoblashni emas, balki qanday hisoblashni tushunishingiz kerak. Shuning uchun men, juda sodda vazifaga e’tiboringizni qaratmoqchiman: To’plam (array) ichidagi sonlar yig’indisini hisoblash. O’ylaymanki, bu algoritm har qanday dasturchiga tanish va GPU’da umumiy maqsadli hisoblash (General-Purpose Computing on Graphics Processing Units) misolini ko’rsatish uchun yetarlidir. Biroq darhol kod yozishdan avval, ba’zi tushunchalarni bilishingiz zarur. Shuning uchun men ularni quyidagicha tartibda taqdim qilmoqchiman:
- Hisoblash modulini yaratish
- Resurslar guruhini yaratish
- Hisoblash quvurini sozlash
Hisoblash modulini yaratish
Hisoblash moduli WGSL deb nomlanuvchi dasturlash tilida yoziladi. Bu til sintaktik jihatdan Rust dasturlash tiliga juda yaqin va u ham strict. Shuningdek u ba’zi matematik obyektlarni o’z ichiga oladi. Masalan vec3 (3 o’lchamli vektor), mat3x3 (3 ga 3 o’lchamli Matritsa). GPU uchun yozilgan kodlar odatda sheyderlar deyiladi. Demak bizning birinchi sheyderimiz quyidagicha:
fn sum() {}
Ajoyib! Sheyder rasman to’g’ri yozilgan ammo, u hisoblash sheyderi bo’lish uchun yetarli emas. WebGPU’da 3 xil, Compute, Vertex va Fragment sheyderlari mavjud. Vertex va Fragment sheyderlari — Renderlash mavzusiga doir bo’lgani uchun, men ularni bu yerda muhokama qilmayman. Compute sheyderi esa — umumiy maqsadlar uchun hisoblashni amalga oshirishda ishlatiladi va bu hozir ayni maqolaga aloqador mavzu.
WGSL haqida ko'proq bilishni istasangiz, bu foydali bo'lishi mumkin.
Shunday qilib siz qaysi turdagi sheyder e’lon qilayotganingizni WGSL ga tegishli atributlar bilan belgilab ketishingiz zarur. Bular @vertex, @fragment va@compute. Mana bizning yaxshilangan kodimiz:
@compute
fn sum() {}
Har qanday e’lon qilinadigan funksiya uchun - ularning turini belgilovchi atributlar e’lon qilinishi majburiy emas. Ular shunchaki yordamchi funksiyalar sifatida ham e’lon qilinishi mumkin. Masalan boshqa dasturlash tillaridagi private funksiyalar kabi.
Resurslar guruhini yaratish
WGSL turli resurslar (masalan xotira bufferlari, texture, sampler) bilan ishlashda — ularni turli guruhlarga ajratish imkonini beradi. Resurs guruhlari alohida shaklda yaratilmaydi. Bu xuddi aviakassa hodimiga “Men uchun Biznes sinfidan, oyna oldidagi joyni olib bering” degandek oson.
WebGPU’da resurs guruhlari va ularga tegishli joylar soni cheklangan (Samolyotlarda ham 🤪) va bu sizning grafik qurilmangizga bog’liq. Siz har qanday cheklovlarni bu yerdan bilishingiz mumkin. Buni kod yordamida aniqlash, hozircha ortiqcha.
Xo’sh, nazariy jihatdan tushunarli ammo WGSL da qanday yoziladi?
@group(0) @binding(0) var input: array<f32>;
@group(0) @binding(1) var result: f32;
@compute
fn sum() {}
Yaxshi. Kod tushunarli ammo, bu ishlamaydi.
Agar e’lon qilinayotgan binding - texture va sampler’ga doir ma’lumot turi bo’lmasa, WebGPU ulardan qanday foydalanish kerakligini bildirishingizni talab qiladi. Ya’ni, deklaratsiya paytida — GPU tomonidan faqat o’qish yoki o’qish va o’zgartira olish (qayta ishlash) imkoniyatiga ega ekanligini e’lon qilish kerak. Bu xuddi Rust’da “mutable” va “immutable” binding’lar e’lon qilinishiga o’xshaydi.
To’g’rilangan kod:
@group(0) @binding(0) var<storage, read> input: array<f32>;
@group(0) @binding(1) var<storage, read_write> f32;
@compute
fn sum() {}
Kodimiz bir qancha muhokamalarga muhtoj. Masalan binding’lardan funksiya ichida foydalansak, nega ularni o’sha funksiyaning parametri sifatida e’lon qilmadik? Nima uchun array’ning hajmi berilmagan? Buning javoblarini keyingi maqolalarda yoritib berishga harakat qilaman. Hozircha binding’larni sheyderning boshida, o’rtasida yoki oxirida ham yozish mumkinligini va WGSL da 2 xil fixed-size va runtime-sized array turlari mavjudligini bilishingiz kifoya. Array ichidagi elementlar sonining miqdorini aniqlash, mana biz yetib kelgan hisoblash quvurini (ComputePipeline) yaratish bosqichida aniqlanadi.
Hisoblash quvurini sozlash
WebGPU — Quvur (pipeline) deb nomlanuvchi konsepsiyani taqdim qiladi. Hozircha 2 xil RenderPipeline va ComputePipeline quvurlari mavjud. Quvurlar bir nechta bosqich (stage) lardan iborat bo’lishi mumkin. Har bir bosqichda GPU’da ma’lum bir vazifa bajariladi. Odatda ComputePipeline’da bitta bosqich ishlaydi — Compute stage. RenderPipeline esa bir nechta bosqichlardan iborat - Vertex Stage, Primitive Assembly, Rasterization, Fragment Stage, Output Merger. Hisoblash quvuridagi compute state - renderlash quvuridagi bosqichlardan farqli o’laroq, grafik primitivlar bilan ishlashga emas, umumiy maqsadli hisoblashlarni amalga oshirishga mo’ljallangan. Quyida compute stage’ning asosiy elementlari ya’ni operatsiyalari bilan qisqacha tanishamiz:
- Dispatching: Hisoblash (@compute) shader’ni bajarish uchun dispatch buyrug’i ishlatiladi. Shuningdek hisoblash sheyderini qancha miqdorda "ishchi guruhi" (workgroup) orqali ishlatish kerakligini belgilaydi.
- Workgroups: Workgroup — bu hisoblash sheyderini bajarish uchun ishlatiladigan asosiy ishchi guruh. Har bir workgroup ko’plab parallel thread’larni o’z ichiga oladi.
- Work Items: Har bir work item (yoki thread) — bu workgroup ichidagi bitta hisoblash birligi. Har bir thread o’z identifikatoriga ega va o’z vazifasini bajaradi.
Qo’shimcha: Hisoblash sheyderi GPU xotirasiga kirish huquqiga ega. Quyida turli xil xotira turlari haqida ma’lumot beraman:
- Global Memory: Barcha workgroup’lar uchun umumiy bo’lgan xotira.
- Shared Memory: Faqat bitta workgroup ichidagi work item’lar tomonidan umumiy foydalaniladigan xotira.
- Local Memory: Har bir thread uchun xos bo’lgan xotira.
Workgroup ichidagi work item’lar o’zaro xotira almashish jarayonida sinxronlashtirilishi mumkin. Bu sinxronizatsiya barrier buyruqlari yordamida amalga oshiriladi, bu esa barcha thread’lar ma'lum bir nuqtaga yetishini ta'minlaydi. barrier buyruqlarining ham bir nechta turlari (masalan MemoryBarrier, WorkgroupBarrier) mavjud. Batafsilroq yozishni istardim ammo, bu - ushbu maqoladan fokusimizni uzoqlarga olib ketishi aniq. Mening hozirgi vazifam esa — sizga GPU’da oddiy hisoblashni amalga oshirishni ko’rsatishdan iborat.
ComputePipeline WGSL tilida alohida shaklda yaratilmaydi. Ammo uning compute bosqichida bajariladigan operatsiyalar uchun muhim ma’lumotlarni sozlashimiz kerak:
@group(0) @binding(0) var<storage, read> input: array<f32>;
@group(0) @binding(1) var<storage, read_write> f32;
@compute
@workgroup_size(1)
fn sum() {}
Unutmang, workgroup_size WGSL'da compute sheyderlarini sozlashda juda muhim atribut hisoblanadi. U — parallel hisoblashlarni tashkil etish va boshqarish uchun ishlatiladi. To'g'ri sozlangan o’lcham, hisoblash samaradorligini oshirishga va GPU resurslaridan optimal foydalanishga yordam beradi. Bu o’lcham uch o’lchovli bo’lishi mumkin: X, Y, Z. Yuqorida yozgan kodimizda, biz faqat X uchun qiymat berdik ammo, odatiy holatda Y va Z qiymatlari 1 ga teng. Ya’ni @workgroup_size(1) bayonoti @workgroup_size(1, 1, 1) bilan bir xil.
WorkGroup — alohida katta mavzu. Men uni hozir muhokama qilmoqchi emasman. Ammo bu yerda batafsilroq ma'lumotlar bor. Hozircha@workgroup_size attributini bilishimiz yetarli. Bu hisoblash sheyderimiz nechta thread’dan foydalanish kerakligini belgilash uchun zarur. Umumiy thread’lar soni X, Y va Z qiymatiga bergan sonlarimizning ko’paytmasiga teng:
@workgroup_size(2) // 2 * 1 * 1 = 2 ta thread @workgroup_size(1,2) // 1 * 2 * 1 = 2 ta thread @workgroup_size(1,2,3) // 1 * 2 * 3 = 6 ta thread
Workgroup ga beriladigan eng katta son ham grafik qurilmangizga bog’liq. Odatda 1024 ga teng bo’ladi. Shunday qilib biz WGSL modulimizni yakunlab qo’yamiz:
@group(0) @binding(0) var<storage, read> input: array<f32>;
@group(0) @binding(1) var<storage, read_write> result: f32;
@compute
@workgroup_size(1)
fn sum() {
for (var i = 0u; i < arrayLength(&input); i++) {
result += input[i];
}
}
// 0u = unsigned 0
Yuqoridagi kod parallel dasturlashga misol bo’la olmaydi.
Agar raqamli ma’lumotlarni e’lon qilishda, ularning turlari aniq ko’rsatilmasa ular — Abstract Numeric Type ko’rinishida bo’ladi. Masalan 123 soni AbstractInt va 3.4 soni AbstractFloat hisoblanadi, Integer yoki Float emas.
Hisoblash sheyderidan foydalanish
Hisoblash sheyderidan foydalanish uchun ham bizga yana reja kerak:
- Hisoblash modulini yaratish (GPUShaderModule)
- Resurslar guruhini yaratish (BindGroup)
- Hisoblash quvurini yaratish (ComputePipeline)
- Buyruqlar kodlovchisini yaratish (CommandEncoder)
- Buyruqlarni GPU’ga yuborish (Submit)
- Natijani olish
1–2–3 bandlar yana takorlanayaptimi? Yo’q, bu safar boshqacharoq bo’ladi. Endi JavaScript’da yozamiz. Faqat men kodlarimni Mr.Doob ning kod usulida yozishga odatlanganman. Deyarli…
Hisoblash modulini yaratish (GPUShaderModule)
WGSL da biz hisoblash modulining manbaa kodini tamomladik. Uni ishlaydigan module sifatida e’lon qilish quyidagicha bo’ladi:
const wgsl = `
// bu yerga WGSL kodimizni yozamiz
`
const adapter = await navigator.gpu.requestAdapter()
const device = await adapter.requestDevice()
const module = device.createShaderModule( { code: wgsl } )
Resurslar guruhini yaratish (BindGroup)
Dastavval hisoblamoqchi bo’lgan to’plamimizni TypedArray yordamida e’lon qilamiz:
const inputData = new Float32Array( [ 1, 2, 3, 4, 5 ] )
WGSL da alohida guruh uchun 2 ta binding (input, result) e’lon qilganimiz yodingizdami? Endi ular bilan ishlash uchun bufferlar yaratib olishimiz kerak:
const inputBuffer = device.createBuffer( {
size: inputData.byteLength, // 20 = 4 byte * 5 (input'da 5 ta son bor)
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
mappedAtCreation: true,
} )
const resultBuffer = device.createBuffer( {
size: 4,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
} )
const stagingBuffer = device.createBuffer( {
size: 4,
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
} )
GPUDevice’ning createBuffer metodi yordamida 2 ta emas 3 ta buffer yaratib oldik. Nima uchun 3 ta? Birinchi buffer WGSL da e’lon qilingan input bilan ishlash uchun, ikkinchisi result bilan ishlash uchun. Tushunarli. Lekin stagingBuffer nega kerak?
Staging buffer - bu CPU va GPU o'rtasida ma'lumotlarni ko'chirish uchun ishlatiladigan vaqtinchalik xotira hududi. Bu buffer CPU tomonidan yoziladigan va GPU tomonidan o'qiladigan yoki aksincha bo'lishi mumkin.
Bizning holatimizda hisoblash shaderi o’z ishini tugatganidan keyin, natijani result qiymatidan to’g’ridan-to’g’ri olishning imkoni yo’q. Shuning uchun undagi buffer ma’lumotlarini stagingBuffer ga ko’chiramiz. Bu jarayon Download deyiladi. Ya’ni GPU’dagi ma’lumotni CPU’ga yuklab olamiz. Aksincha bo’lsa Upload.
Mana nima uchun har bir bufferda usage xususiyati bor. Ya’ni bu bilan bufferdan qanday maqsadda foydalanayotganimizni bildiramiz. Agar qiymat GPUBufferUsage.STORAGE bo’lsa — xotira ombori sifatida va agar GPUBufferUsage.MAP_READ bo’lsa bu o’qish uchun map qilish maqsadida foydalanishimizni bildiradi. Nima? Map qilish?
WebGPU’da map va unmap qilish buferning ma'lumotlariga kirish imkoniyatini olish va bu kirishni tugatishni anglatadi. Buferni map qilish buferning ma'lumotlariga CPU tomonidan kirishni amalga oshiradi va unmap qilish bu kirishni yopadi.
Buferni map qilish ikki xil usulda amalga oshiriladi:
- GPUBufferUsage.MAP_READ: Buferni o'qish uchun map qilish. Bu holatda GPU'da saqlangan ma'lumotlarni CPU o'qishi mumkin.
- GPUBufferUsage.MAP_WRITE: Buferni yozish uchun map qilish. Bu holatda CPU ma'lumotlarni buferga yozishi mumkin.
Ma’lumotlar map qilinganidan keyin unmap qilinishi kerak, aks holda buferni boshqa operatsiyalar uchun ishlatib bo'lmaydi.
Input buffer yaratilayotgan payt mappedAtCreation xususiyatiga true qiymatini berdik. Chunki u yaratilayotgan paytidayoq CPU ga kerak. Ya’ni keyingi qatorlarda inputBuffer ga inputData dagi sonlarni ko’chirishimiz kerak. Buning uchun inputBufferning ArrayBufferini olamiz. Uni to’g’ridan to’g’ri manipulyatsiya qilishning imkoni yo’qligi bois, yangi Float32Array yaratish usuli bilan ko’chirib olamiz. ArrayBuffer’ning WebGPU ga aloqasi yo’q. Siz u haqida batafsil ma’lumotni bu yerdan topasiz. So’ng inputBuffer ni unmap qilamiz. Ya’ni CPU dan kirish to’xtatilib GPU da foydalanish yoqilishi uchun:
const arrayBuffer = inputBuffer.getMappedRange() new Float32Array( arrayBuffer ).set( inputData ) inputBuffer.unmap()
GPUBufferUsage.COPY_DST va GPUBufferUsage.COPY_SRC WebGPU'da buferlarni qanday ishlatishni aniqlaydigan flag'lar hisoblanadi. Bu flag'lar ma'lumotlarni GPU buferlari o'rtasida nusxalash (copy) operatsiyalarini bajarish uchun ishlatiladi.
COPY_DST flag'iga ega bufer, ma'lumotlarni nusxalashda manzil sifatida ishlatiladi. Boshqa buferlardan yoki manbalardan ma'lumotlarni bu shu ko'chirish uchun ishlatiladi. COPY_SRC flag'iga ega bufer esa, ma'lumotlarni nusxalashda manba sifatida ishlatiladi. Ya’ni shu bufer ma'lumotlarni boshqa buferlarga ko'chirish uchun manba sifatida ishlatiladi.
Hozirchaga biz faqat bufferlarni yaratdik. Biroq hali ham BindGroup (resurslar guruhini) yaratmadik. Endi ayni vaqti:
const bindGroupLayout = device.createBindGroupLayout( {
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: "read-only-storage" }
},
{
binding: 1,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: "storage" }
},
],
} )
const bindGroup = device.createBindGroup( {
layout: bindGroupLayout,
entries: [
{
binding: 0,
resource: { buffer: inputBuffer }
},
{
binding: 1,
resource: { buffer: resultBuffer }
},
],
} )
bindGroupLayout ichida e’lon qilingan visibility xususiyati, o’sha buffer — sheyderning hisoblash bosqichida ko’rinishi mumkinligini bildiradi. Kodning qolgan qismi menimcha tushunarli.
Resurslar guruhini ham yaratib oldik. Endi navbat hisoblash quvuriga!
Hisoblash quvurini yaratish (ComputePipeline)
Quvur yaratishda ham avval layout kerak bo’ladi. Ammo biz uni to’g’ridan to’g’ri hisoblash quvurini yaratayotgan paytimiz e’lon qilib ketamiz:
const computePipeline = device.createComputePipeline( {
layout: device.createPipelineLayout( {
bindGroupLayouts: [ bindGroupLayout, ]
} ),
compute: {
module, // yuqorida yaratilgan hisoblash moduli
},
} )
Ko’rib turganingizdek, Pipeline yaratish juda oddiy. (unga kerak bo’lgan narsalarni yaratishni aytmaganda 🤭). Unda buyruqlar kodlovchisini yaratishga o’tamiz.
Buyruqlar kodlovchisini yaratish (CommandEncoder)
CommandEncoder’ni qanday tushuntirish haqida juda uzoq o’yladim. Bilasizmi, men uni qanday tasavvur qilaman?
Uzun quvur. Bu quvur RenderPipeline yoki ComputePipeline bo’lishining ahamiyati yo’q. Quvurning narigi tarafida mening 3 yoshli qizim Aisha o’tiripdi. U juda aqlli va rasm chiza oladigan darajada iqtidorli. Ammo men uning bola tilini tushunmayman. Kattalar tilida o’zim tushungancha “Qizim mana bu narsani hisoblab ber”, yoki “Menga rasm chizib ber” deb ayta olaman. (Bu til hozirgi o’rinda WGSL.) Keyin uni baland ovozda o’qiyman. Mening ovoz to’lqinlarim, uning miyasiga o’zi tushunadigan bolalar tilida yetib boradi va u mening topshiriqlarimni bajaradi. Bu joyda CommandEncoder o’sha ovoz to’lqinlarini — bolalar tiliga kodlash (tarjima qilish) uchun kerak.
Texnik tafsilotiga ko’ra CommandEncoder, GPU’da bajariladigan komanda buyrug’larini yaratish, guruhlash va bu buyrug’larni GPU’ga topshirish imkonini beradi. CommandEncoder yordamida siz buferlar va teksturalar o'rtasida ma'lumot ko'chirish, render pass'lar va hisoblash pass'larni boshlash va yakunlash kabi operatsiyalarni amalga oshirishingiz mumkin.
Uni yaratish pipeline yaratishdan ham oson:
const commandEncoder = device.createCommandEncoder()
Demak bizda ComputePipeline quvuri va bu quvur bo’ylab topshiriqlarni yetkazib berish uchun CommandEncoder bor. Endi topshiriq berishni boshlashimiz kerak. Buning uchun commandEncoder bilan hisoblash buyrug’ini yaratib olamiz. Balki buyruqlar quvur orqali o’tgani uchunmi ular — o’tish (Pass) deb nomlanadi:
const passEncoder = commandEncoder.beginComputePass() passEncoder.setPipeline( computePipeline ) passEncoder.setBindGroup( 0, bindGroup ) passEncoder.dispatchWorkgroups( 1 ) passEncoder.end()
Yuqoridagi kodda hisoblash o’tishini (compute pass’ni) amalga oshirdik.
- Topshiriq yaratishni boshladik (beginComputePass).
- U qaysi quvurdan o’tishini belgiladik (setPipeline).
- Buyruq qaysi resurslar guruhiga tegishli ekanini ko’rsatdik (setBindGroup).
- Vazifani bajarish uchun nechta ishchi guruhdan foydalanishi kerakligini bildirdik (dispatchWorkgroups).
- So’ng buyruqni yakunladik (end).
Hisoblash hali amalga oshmadi. Amalga oshganda ham bizga hech narsa bo’lmagandek tuyulardi. Chunki natija GPU ichida qolib ketadi. Shuning uchun biz GPU hisoblab bo’lgan ma’lumotni vaqtinchalik bufferimizga ya’ni stagingBufferga ko’chirishimiz kerak. Bu ishni yana commandEncoder qiladi:
commandEncoder.copyBufferToBuffer( resultBuffer, 0, stagingBuffer, 0, 4 )
copyBufferToBuffer parametrlariga resultBufferning 0 indeksidan boshlab (ya’ni boshidan) stagingBufferning boshiga, 4 bayt hajmdagi ma’lumotni nusxalashi kerakligi yozilgan.
Buyruqlarni GPU’ga yuborish (Submit)
Biz hisoblash pass’ini va nusxalash buyrug’ini e’lon qildik. Endi buyruqlarni yakunlab, uni GPU ga yuborish vaqti keldi!
device.queue.submit( [ commandEncoder.finish() ] )
Ha, bu juda oddiy edi. Nihoyat hisoblashni amalga oshirganingiz bilan tabriklashim mumkin! Ammo sizda hech narsa ko’rinmadimi? Bu juda yaxshi chunki, “Jimlik — oltinga teng!”. Linux operatsion tizim falsafasi.
Chindan ham sizda hech narsa ko’rinmagan bo’lsa demak, kodingizda xatolar yo’q. Aks holda WGSL brauzer konsolingizni turli xatolikka doir xabarlar bilan to’ldirib yuborgan bo’lar edi.
Natijani olish
GPU ba buffer ma’lumotlarimiz qolib ketgani kabi, hozir u stagingBuffer da turibdi. Ammo undan to‘g‘ridan to‘g‘ri foydalanishning iloji yo’q. Buning uchun yana ArrayBuffer va TypedArray’lar yordamga keladi:
await stagingBuffer.mapAsync( GPUMapMode.READ ) const resultArrayBuffer = stagingBuffer.getMappedRange() const result = new Float32Array( resultArrayBuffer )[ 0 ] stagingBuffer.unmap() console.log( result ) // 15
stagingBuffer asinxron tarzda map qilinadi (mapAsync). Yana GPUMapMode.READ bilan CPU ushbu ma’lumotni o’qishi kerakligini aytamiz. Bufferdagi ma’lumotni o’qib uni result o’zgaruvchisiga ko’chiramiz va keyin unmap qilamiz. Tamom.
Barcha kodlarni birlashtirish
const wgsl = `
@group(0) @binding(0) var<storage, read> input: array<f32>;
@group(0) @binding(1) var<storage, read_write> result: f32;
@compute
@workgroup_size(1)
fn sum() {
for (var i = 0u; i < arrayLength(&input); i++) {
result += input[i];
}
}
`
const adapter = await navigator.gpu.requestAdapter()
const device = await adapter.requestDevice()
const module = device.createShaderModule( { code: wgsl } )
const inputData = new Float32Array( [ 1, 2, 3, 4, 5 ] )
const inputBuffer = device.createBuffer( {
size: inputData.byteLength, // 20 = 4 byte * 5 (input'da 5 ta son bor)
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
mappedAtCreation: true,
} )
const arrayBuffer = inputBuffer.getMappedRange()
new Float32Array( arrayBuffer ).set( inputData )
inputBuffer.unmap()
const resultBuffer = device.createBuffer( {
size: 4,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
} )
const stagingBuffer = device.createBuffer( {
size: 4,
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
} )
const bindGroupLayout = device.createBindGroupLayout( {
entries: [
{
binding: 0,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: "read-only-storage" }
},
{
binding: 1,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: "storage" }
},
],
} )
const bindGroup = device.createBindGroup( {
layout: bindGroupLayout,
entries: [
{
binding: 0,
resource: { buffer: inputBuffer }
},
{
binding: 1,
resource: { buffer: resultBuffer }
},
],
} )
const computePipeline = device.createComputePipeline( {
layout: device.createPipelineLayout( {
bindGroupLayouts: [ bindGroupLayout, ]
} ),
compute: {
module, // yuqorida yaratilgan hisoblash moduli
},
} )
const commandEncoder = device.createCommandEncoder()
const passEncoder = commandEncoder.beginComputePass()
passEncoder.setPipeline( computePipeline )
passEncoder.setBindGroup( 0, bindGroup )
passEncoder.dispatchWorkgroups( 1 )
passEncoder.end()
commandEncoder.copyBufferToBuffer( resultBuffer, 0, stagingBuffer, 0, 4 )
device.queue.submit( [ commandEncoder.finish() ] )
await stagingBuffer.mapAsync( GPUMapMode.READ )
const resultArrayBuffer = stagingBuffer.getMappedRange()
const result = new Float32Array( resultArrayBuffer )[ 0 ]
stagingBuffer.unmap()
console.log( result ) // 15
GPU’da hisoblash qanday amalga oshirilishini tushuntirib bera olgan bo’lsam, bundan xursandman. Kodlarim anchagina improvement’ga muhtoj. Albatta men ularni yaxshilashga harakat qilaman. Agar maqolada biror xato yoki kamchilikni aniqlasangiz, iltimos buni bildirishdan tortinmang. Bu maqolani yaxshilashga hissa qo’shishi mumkin.
Rahmat!