Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

regression: web gpu working incorrectly #24857

Closed
sigmaSd opened this issue Aug 2, 2024 · 8 comments
Closed

regression: web gpu working incorrectly #24857

sigmaSd opened this issue Aug 2, 2024 · 8 comments
Assignees
Labels
needs investigation requires further investigation before determining if it is an issue or not webgpu WebGPU API

Comments

@sigmaSd
Copy link
Contributor

sigmaSd commented Aug 2, 2024

deno run  --import-map https://raw.githubusercontent.com/denoland/webgpu-examples/main/deno.json https://
raw.githubusercontent.com/denoland/webgpu-examples/main/hello-compute/mod.ts

outputs [ "1", "4", "3", "295" ] instead of ["0", "2", "7", "55"]

This works with version 1.43.0 but not 1.44.0 @crowlKats suspects #23684

meta:
Linux 1E-55-11-07-05-2E 6.9.11-200.fc40.x86_64 #1 SMP PREEMPT_DYNAMIC Thu Jul 25 18:17:34 UTC 2024 x86_64 GNU/Linux
Brand Raw: Intel(R) Core(TM) i3-3110M CPU @ 2.40GHz

@crowlKats crowlKats self-assigned this Aug 3, 2024
@crowlKats crowlKats added needs investigation requires further investigation before determining if it is an issue or not webgpu WebGPU API labels Aug 3, 2024
@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 6, 2024

Here is a self contained example

async function addWithWebGPU(a: number, b: number) {
  if (!navigator.gpu) {
    throw new Error("WebGPU not supported on this browser.");
  }

  console.log("Requesting adapter...");
  const adapter = await navigator.gpu.requestAdapter();
  if (!adapter) {
    throw new Error("No appropriate GPUAdapter found.");
  }

  console.log("Requesting device...");
  const device = await adapter.requestDevice();

  console.log("Creating buffer...");
  const bufferSize = 3 * Float32Array.BYTES_PER_ELEMENT;
  const buffer = device.createBuffer({
    size: bufferSize,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC |
      GPUBufferUsage.COPY_DST,
  });

  console.log("Writing input data...");
  const input = new Float32Array([a, b, 0]);
  device.queue.writeBuffer(buffer, 0, input);

  console.log("Creating shader module...");
  const shaderModule = device.createShaderModule({
    code: `
            @group(0) @binding(0) var<storage, read_write> data: array<f32>;

            @compute @workgroup_size(1)
            fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
                data[2] = data[0] + data[1];
            }
        `,
  });

  console.log("Creating compute pipeline...");
  const pipeline = device.createComputePipeline({
    layout: "auto",
    compute: {
      module: shaderModule,
      entryPoint: "main",
    },
  });

  console.log("Creating bind group...");
  const bindGroup = device.createBindGroup({
    layout: pipeline.getBindGroupLayout(0),
    entries: [{ binding: 0, resource: { buffer: buffer } }],
  });

  console.log("Encoding commands...");
  const commandEncoder = device.createCommandEncoder();
  const passEncoder = commandEncoder.beginComputePass();
  passEncoder.setPipeline(pipeline);
  passEncoder.setBindGroup(0, bindGroup);
  passEncoder.dispatchWorkgroups(1);
  passEncoder.end();

  console.log("Creating read buffer...");
  const gpuReadBuffer = device.createBuffer({
    size: bufferSize,
    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
  });

  console.log("Encoding copy command...");
  commandEncoder.copyBufferToBuffer(buffer, 0, gpuReadBuffer, 0, bufferSize);

  console.log("Submitting GPU commands...");
  const gpuCommands = commandEncoder.finish();
  device.queue.submit([gpuCommands]);

  console.log("Mapping read buffer...");
  await gpuReadBuffer.mapAsync(GPUMapMode.READ);
  const resultArrayBuffer = gpuReadBuffer.getMappedRange();
  const result = new Float32Array(resultArrayBuffer);

  console.log("Raw result array:", result);

  // Extract the result before unmapping
  const finalResult = result[2];

  gpuReadBuffer.unmap();

  console.log("Computation complete.");
  return finalResult;
}

// Usage
async function testAddition() {
  try {
    const result = await addWithWebGPU(5, 7);
    console.log(`5 + 7 = ${result}`);
  } catch (error) {
    console.error("An error occurred:", error);
  }
}

// Run the test
testAddition().then(() => console.log("Test completed")).catch((error) =>
  console.error("Test failed:", error)
);

it outputs 0 instead of 12

but this works on firefox nightly 131.0a1

@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 6, 2024

All the other examples work, it seems like the issue is specifically with compute

@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 9, 2024

If I run the rust wgpu compute example, I get

No numbers were provided, defaulting to [1, 2, 3, 4]
MESA-INTEL: warning: Ivy Bridge Vulkan support is incomplete
thread 'main' panicked at /home/mrcool/.cargo/registry/src/index.crates.io-6f17d22bba15001f/wgpu-22.1.0/src/backend/wgpu_core.rs:3411:5:
wgpu error: Validation Error

Caused by:
  In Device::create_compute_pipeline
    Internal error: The selected version doesn't support Features(DYNAMIC_ARRAY_SIZE)

@sigmaSd sigmaSd changed the title regression: web gpu working incorrectly (maybe because of old cpu) regression: web gpu working incorrectly Aug 10, 2024
@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 10, 2024

if I run with XDG_RUNTIME_DIR="" deno run -A a.ts (unsetting XDG_RUNTIME_DIR) it actually works

@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 10, 2024

the issue seem to be specifically with XDG_RUNTIME_DIR="/run/user/1000" if I put another value like "/tmp" it works

@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 10, 2024

the same happens in rust webgpu, so its a webgpu bug

@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 10, 2024

I guess the only thing deno can do better here is surface such errors

@sigmaSd
Copy link
Contributor Author

sigmaSd commented Aug 14, 2024

wgpu now doesn't support incomplete vulkan implementation, that's why it stopped working, it works on xorg because for some reason its picking llvmpipe (vulkan cpu implemtation) as the adapter which workarounds this issue

@sigmaSd sigmaSd closed this as completed Aug 14, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
needs investigation requires further investigation before determining if it is an issue or not webgpu WebGPU API
Projects
None yet
Development

No branches or pull requests

2 participants