WebGPU: Read Pixel via texture_storage_2d from different mip levels

README.md

WebGPU: Read Pixel via texture_storage_2d from different mip levels

view on jsgist

index.css

:root { color-scheme: light dark; }
pre { margin: 0;}

index.html

/*bug-in-github-api-content-can-not-be-empty*/

index.js

const adapter = await navigator.gpu.requestAdapter({
  featureLevel: 'compatibility',
});
const device = await adapter.requestDevice();
device.addEventListener('uncapturederror', e => console.error(e.error.message));

const texture = device.createTexture({
  format: 'rgba8unorm',
  size: [2, 2, 1],
  mipLevelCount: 2,
  usage: GPUTextureUsage.COPY_DST |
         GPUTextureUsage.TEXTURE_BINDING |
         GPUTextureUsage.STORAGE_BINDING,
});

const layers = [0, -2, -1, 21, 22];

const c = {
  black: [0, 0, 0, 0],
  red: [255, 0, 0, 255],
  green: [0, 255, 0, 255],
  blue: [0, 0, 255, 255],
  yellow: [255, 255, 0, 255],
  cyan: [0, 255, 255, 255],
  magenta: [255, 0, 255, 255],
  gray: [128, 128, 128, 255],
  orange: [255, 168, 0, 255],
};
const valueToColor = new Map(Object.entries(c).map(([k, v]) => [v.toString(), k]));

const colors = [
  // mip level 0
  [
    c.red, c.red,
    c.red, c.red,
  ],
  // miplevel 2
  [
    c.green
  ],
];

for (let z = 0; z < texture.mipLevelCount; ++z) {
  log('mip level:', z);
  const width = texture.width >> z;
  const height = texture.width >> z;
  for (let y = 0; y < height; ++y) {
    const row = [];
    for (let x = 0; x < width; ++x) {
      row.push(pixelToColor(colors[z][y * width + x]));
    }
    log('  ', row.join(', '))
  }
}

device.queue.writeTexture(
  {texture},
  new Uint8Array(colors[0].flat()),
  { bytesPerRow: 8, rowsPerImage: 2 },
  [ 2, 2 ],
);
device.queue.writeTexture(
  {texture, mipLevel: 1},
  new Uint8Array(colors[1].flat()),
  { bytesPerRow: 4, rowsPerImage: 1 },
  [ 1, 1 ],
);

const sampler = device.createSampler({
  addressModeW: 'repeat',
});

const code = `
@group(0) @binding(0) var t0: texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var t1: texture_2d<f32>;
@group(0) @binding(2) var<storage, read_write> result: array<vec4f>;

@compute @workgroup_size(1) fn cs() {
  result[0] = textureLoad(t0, vec2u(0));
  result[1] = textureLoad(t1, vec2u(0), 0);
}

`;
log(code);
const module = device.createShaderModule({code});
const pipeline = device.createComputePipeline({
  layout: 'auto',
  compute: { module },
});

const storageBuffer = device.createBuffer({
  size: 512,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});

const resultBuffer = device.createBuffer({
  size: storageBuffer.size,
  usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
});

const encoder = device.createCommandEncoder();
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
for (let i = 0; i < 2; ++i) {
  const bindGroup = device.createBindGroup({
    layout: pipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: texture.createView({baseMipLevel: i, mipLevelCount: 1}) },
      { binding: 1, resource: texture.createView({baseMipLevel: i, mipLevelCount: 1}) },
      { binding: 2, resource: { buffer: storageBuffer, offset: i * 256, size: 32 }},
    ],
  });

  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(1);
}
pass.end();
encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size);
device.queue.submit([encoder.finish()]);

await resultBuffer.mapAsync(GPUMapMode.READ);
const result = new Float32Array(resultBuffer.getMappedRange());
const pixels = Array.from(result).map(v => v * 255 | 0);
for (let i = 0; i < 2; ++i) {
  for (let j = 0; j < 2; ++j) {
    log(`${i} ${j}: ${pixelToColor(pixels.slice(i * 64 + j * 4, i * 64 + j * 4 + 4))}`);
  }
}

function pixelToColor(p) {
  const name = valueToColor.get(p.toString());
  return name ?? p.toString();
}

function log(...args) {
  const elem = document.createElement('pre');
  elem.textContent = args.join(' ');
  document.body.appendChild(elem);
}

jsGist.json

{"name":"WebGPU: Read Pixel via texture_storage_2d from different mip levels","settings":{},"filenames":["index.html","index.css","index.js"]}