Code Monkey home page Code Monkey logo

wgpu-compute-toy's Introduction

wgpu-compute-toy's People

Contributors

cornuammonis avatar davidar avatar jmatsushita avatar munrocket avatar pudnax avatar pwfff avatar slerpyyy avatar stefnotch avatar

Stargazers

 avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar

Watchers

 avatar  avatar  avatar  avatar  avatar

wgpu-compute-toy's Issues

No images found when running some examples

I attempted to test some examples, and most of them ran perfectly. However, I found that some examples require texture images. These examples run successfully when the images are sourced from URLs. However, when attempting to use local images, the examples generate an 'image not found' error. For instance, the '/fabrice/utils.wgsl' example includes a JSON file:

{
    "uniforms": [],
    "textures": [
        {
            "img": "/textures/font0.png"
        },
        {
            "img": "/textures/london.jpg"
        }
    ],
    "float32Enabled": false
}

It appears that the 'font0.png' and 'london.jpg' files are not present in the project's directory structure. How can this issue be resolved?

Additionally, does the 'wgpu-compute-toy' project support cube maps, in addition to 2D textures?

FR: add expression evaluation in defines

Add WIN_WIDTH / WIN_HEIGHT in prelude with ability to write expressions in defines

// usual case
#define GROUP_COUNT_X WIN_WIDTH/16
#define GROUP_COUNT_Y WIN_HEIGHT/16

// example 1
struct Data {
  foo: array<array<u32, GROUP_COUNT_Y>, GROUP_COUNT_X>
}
#storage data Data

// example 2 (this calculated automatically somehow)
#workgroup_count bar GROUP_COUNT_X GROUP_COUNT_Y 1
@compute @workgroup_size(16, 16)
fn bar(@builtin(global_invocation_id) gid: vec3u) {
  //...
}

Shader translation error for stage ShaderStages(COMPUTE): Metal: program_source:79:5: error: too many arguments provided to function-like macro invocation

Hi there,

On MacOS 14.3.1 (23D60), on an M1 Mac I can't run the default shader default.wgsl or any other I tried. The main error seems to be Shader translation error for stage ShaderStages(COMPUTE): Metal: program_source:79:5: error: too many arguments provided to function-like macro invocation. Here is the full error message for cargo run examples/default.wgsl

> cargo run examples/default.wgsl                                                       11:48:12
warning: unused variable: `bind_id`
  --> src/context.rs:81:49
   |
81 | pub async fn init_wgpu(width: u32, height: u32, bind_id: &str) -> Result<WgpuContext, String> {
   |                                                 ^^^^^^^ help: if this is intentional, prefix it with an underscore: `_bind_id`
   |
   = note: `#[warn(unused_variables)]` on by default

warning: unused variable: `timestamps`
   --> src/lib.rs:346:25
    |
346 |                     let timestamps: &[u64] = bytemuck::cast_slice(&data[ASSERTS_SIZE..]);
    |                         ^^^^^^^^^^ help: if this is intentional, prefix it with an underscore: `_timestamps`

warning: a method with this name may be added to the standard library in the future
   --> src/bind.rs:610:14
    |
610 |             .intersperse("\n".to_string())
    |              ^^^^^^^^^^^
    |
    = warning: once this associated item is added to the standard library, the ambiguity may cause an error or change in behavior!
    = note: for more information, see issue #48919 <https://github.com/rust-lang/rust/issues/48919>
    = help: call with fully qualified syntax `itertools::Itertools::intersperse(...)` to keep using the current method
    = note: `#[warn(unstable_name_collisions)]` on by default

warning: field `name` is never read
  --> src/lib.rs:55:5
   |
54 | struct ComputePipeline {
   |        --------------- field in this struct
55 |     name: String,
   |     ^^^^
   |
   = note: `#[warn(dead_code)]` on by default

warning: field `on_success_cb` is never read
  --> src/lib.rs:75:5
   |
64 | pub struct WgpuToyRenderer {
   |            --------------- field in this struct
...
75 |     on_success_cb: SuccessCallback,
   |     ^^^^^^^^^^^^^

warning: function `set_panic_hook` is never used
 --> src/utils.rs:9:8
  |
9 | pub fn set_panic_hook() {
  |        ^^^^^^^^^^^^^^

warning: `wgputoy` (lib) generated 6 warnings (run `cargo fix --lib -p wgputoy` to apply 2 suggestions)
warning: unreachable pattern
   --> src/bin/toy.rs:161:25
    |
161 |                         _ => ()
    |                         ^
    |
    = note: `#[warn(unreachable_patterns)]` on by default

warning: `wgputoy` (bin "toy") generated 1 warning
    Finished dev [unoptimized + debuginfo] target(s) in 0.34s
     Running `target/debug/toy examples/default.wgsl`
[2024-03-09T10:49:23Z INFO  wgpu_core::instance] Adapter Metal AdapterInfo { name: "Apple M1 Pro", vendor: 0, device: 0, device_type: IntegratedGpu, driver: "", driver_info: "", backend: Metal }
[2024-03-09T10:49:23Z INFO  wgputoy::context] adapter.limits = Limits {
        max_texture_dimension_1d: 16384,
        max_texture_dimension_2d: 16384,
        max_texture_dimension_3d: 2048,
        max_texture_array_layers: 2048,
        max_bind_groups: 8,
        max_bindings_per_bind_group: 65535,
        max_dynamic_uniform_buffers_per_pipeline_layout: 8,
        max_dynamic_storage_buffers_per_pipeline_layout: 4,
        max_sampled_textures_per_shader_stage: 128,
        max_samplers_per_shader_stage: 16,
        max_storage_buffers_per_shader_stage: 31,
        max_storage_textures_per_shader_stage: 128,
        max_uniform_buffers_per_shader_stage: 31,
        max_uniform_buffer_binding_size: 4294967295,
        max_storage_buffer_binding_size: 4294967295,
        max_vertex_buffers: 16,
        max_buffer_size: 17179869184,
        max_vertex_attributes: 31,
        max_vertex_buffer_array_stride: 2048,
        min_uniform_buffer_offset_alignment: 256,
        min_storage_buffer_offset_alignment: 256,
        max_inter_stage_shader_components: 124,
        max_compute_workgroup_storage_size: 32768,
        max_compute_invocations_per_workgroup: 1024,
        max_compute_workgroup_size_x: 1024,
        max_compute_workgroup_size_y: 1024,
        max_compute_workgroup_size_z: 1024,
        max_compute_workgroups_per_dimension: 65535,
        max_push_constant_size: 4096,
        max_non_sampler_bindings: 4294967295,
    }
[2024-03-09T10:49:23Z INFO  wgputoy::bind] Creating bindings
@compute @workgroup_size(16, 16)
fn main_image(@builtin(global_invocation_id) id: uint3) {
    // Viewport resolution (in pixels)
    let screen_size = uint2(textureDimensions(screen));

    // Prevent overdraw for workgroups on the edge of the viewport
    if (id.x >= screen_size.x || id.y >= screen_size.y) { return; }

    // Pixel coordinates (centre of pixel, origin at bottom left)
    let fragCoord = float2(float(id.x) + .5, float(screen_size.y - id.y) - .5);

    // Normalised pixel coordinates (from 0 to 1)
    let uv = fragCoord / float2(screen_size);

    // Time varying pixel colour
    var col = .5 + .5 * cos(time.elapsed + uv.xyx + float3(0.,2.,4.));

    // Convert from gamma-encoded to linear colour space
    col = pow(col, float3(2.2));

    // Output to screen (linear colour space)
    textureStore(screen, int2(id.xy), float4(col, 1.));
}

[2024-03-09T10:49:23Z INFO  wgputoy] Entry points: ["main_image"]
[2024-03-09T10:49:23Z WARN  wgpu_hal::metal::device] Naga generated shader:
    // language: metal2.4
    #include <metal_stdlib>
    #include <simd/simd.h>

    using metal::uint;
    struct DefaultConstructible {
        template<typename T>
        operator T() && {
            return T {};
        }
    };

    struct _mslBufferSizes {
        uint size5;
    };

    typedef int int_;
    typedef uint uint_;
    typedef float float_;
    typedef metal::int2 int2_;
    typedef metal::int3 int3_;
    typedef metal::int4 int4_;
    typedef metal::uint2 uint2_;
    typedef metal::uint3 uint3_;
    typedef metal::uint4 uint4_;
    typedef metal::float2 float2_;
    typedef metal::float3 float3_;
    typedef metal::float4 float4_;
    typedef metal::bool2 bool2_;
    typedef metal::bool3 bool3_;
    typedef metal::bool4 bool4_;
    typedef metal::float2x2 float2x2_;
    typedef metal::float2x3 float2x3_;
    typedef metal::float2x4 float2x4_;
    typedef metal::float3x2 float3x2_;
    typedef metal::float3x3 float3x3_;
    typedef metal::float3x4 float3x4_;
    typedef metal::float4x2 float4x2_;
    typedef metal::float4x3 float4x3_;
    typedef metal::float4x4 float4x4_;
    struct Time {
        uint_ frame;
        float_ elapsed;
        float_ delta;
    };
    struct Mouse {
        uint2_ pos;
        int_ click;
    };
    struct DispatchInfo {
        uint_ id;
    };
    struct Custom {
        float_ _dummy;
    };
    struct type_1 {
        uint inner[1];
    };
    struct Data {
        type_1 _dummy;
    };
    struct type_3 {
        metal::uint4 inner[2];
    };
    typedef metal::atomic_uint type_5[1];

    bool keyDown(
        uint_ keycode,
        constant type_3& _keyboard
    ) {
        uint_ _e3 = keycode / 128u;
        uint_ _e8 = (keycode % 128u) / 32u;
        uint _e10 = uint(_e8) < 4 && uint(_e3) < 2 ? _keyboard.inner[_e3][_e8] : DefaultConstructible();
        return ((_e10 >> (keycode % 32u)) & 1u) == 1u;
    }

    void assert(
        int_ index,
        bool success,
        device type_5& _assert_counts,
        constant _mslBufferSizes& _buffer_sizes
    ) {
        if (!(success)) {
            uint _e6 = uint(index) < 1 + (_buffer_sizes.size5 - 0 - 4) / 4 ? metal::atomic_fetch_add_explicit(&_assert_counts[index], 1u, metal::memory_order_relaxed) : DefaultConstructible();
            return;
        } else {
            return;
        }
    }

    void passStore(
        int_ pass_index,
        int2_ coord,
        float4_ value,
        metal::texture2d_array<float, metal::access::write> pass_out
    ) {
        pass_out.write(value, metal::uint2(coord), pass_index);
        return;
    }

    float4_ passLoad(
        int_ pass_index_1,
        int2_ coord_1,
        int_ lod,
        metal::texture2d_array<float, metal::access::sample> pass_in
    ) {
        metal::float4 _e4 = (uint(lod) < pass_in.get_num_mip_levels() && uint(pass_index_1) < pass_in.get_array_size() && metal::all(metal::uint2(coord_1) < metal::uint2(pass_in.get_width(lod), pass_in.get_height(lod))) ? pass_in.read(metal::uint2(coord_1), pass_index_1, lod): DefaultConstructible());
        return _e4;
    }

    float4_ passSampleLevelBilinearRepeat(
        int_ pass_index_2,
        float2_ uv,
        float_ lod_1,
        metal::texture2d_array<float, metal::access::sample> pass_in,
        metal::sampler bilinear
    ) {
        metal::float4 _e6 = pass_in.sample(bilinear, metal::fract(uv), pass_index_2, metal::level(lod_1));
        return _e6;
    }

    struct main_imageInput {
    };
    kernel void main_image(
      uint3_ id [[thread_position_in_grid]]
    , constant Time& time [[buffer(2)]]
    , metal::texture2d<float, metal::access::write> screen [[texture(0)]]
    ) {
        metal::float3 col = {};
        metal::uint2 screen_size = static_cast<metal::uint2>(metal::uint2(screen.get_width(), screen.get_height()));
        if ((id.x >= screen_size.x) || (id.y >= screen_size.y)) {
            return;
        }
        float2_ fragCoord = float2_(static_cast<float>(id.x) + 0.5, static_cast<float>(screen_size.y - id.y) - 0.5);
        float2_ uv_1 = fragCoord / static_cast<metal::float2>(screen_size);
        float_ _e26 = time.elapsed;
        col = metal::float3(0.5) + (0.5 * metal::cos((metal::float3(_e26) + uv_1.xyx) + float3_(0.0, 2.0, 4.0)));
        metal::float3 _e42 = col;
        col = metal::pow(_e42, metal::float3(2.2));
        metal::float3 _e49 = col;
        screen.write(float4_(_e49, 1.0), metal::uint2(static_cast<metal::int2>(id.xy)));
        return;
    }

[2024-03-09T10:49:23Z ERROR wgpu::backend::wgpu_core] Shader translation error for stage ShaderStages(COMPUTE): Metal: program_source:79:5: error: too many arguments provided to function-like macro invocation
        bool success,
        ^
    /System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.116/include/metal/metal_assert:19:9: note: macro 'assert' defined here
    #define assert(condition) ((void) 0)
            ^
    program_source:77:6: error: program scope variable must reside in constant address space
    void assert(
         ^
    program_source:83:5: error: expected expression
        if (!(success)) {
        ^
    program_source:89:2: error: expected ';' after top level declarator
    }
     ^
     ;

[2024-03-09T10:49:23Z ERROR wgpu::backend::wgpu_core] Please report it to https://github.com/gfx-rs/wgpu
[2024-03-09T10:49:23Z ERROR wgputoy] Validation Error

    Caused by:
        In Device::create_compute_pipeline
        Internal error: Metal: program_source:79:5: error: too many arguments provided to function-like macro invocation
        bool success,
        ^
    /System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.116/include/metal/metal_assert:19:9: note: macro 'assert' defined here
    #define assert(condition) ((void) 0)
            ^
    program_source:77:6: error: program scope variable must reside in constant address space
    void assert(
         ^
    program_source:83:5: error: expected expression
        if (!(success)) {
        ^
    program_source:89:2: error: expected ';' after top level declarator
    }
     ^
     ;


thread 'main' panicked at src/pp.rs:29:9:
0:0: Validation Error

Caused by:
    In Device::create_compute_pipeline
    Internal error: Metal: program_source:79:5: error: too many arguments provided to function-like macro invocation
    bool success,
    ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.116/include/metal/metal_assert:19:9: note: macro 'assert' defined here
#define assert(condition) ((void) 0)
        ^
program_source:77:6: error: program scope variable must reside in constant address space
void assert(
     ^
program_source:83:5: error: expected expression
    if (!(success)) {
    ^
program_source:89:2: error: expected ';' after top level declarator
}
 ^
 ;


note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace
[2024-03-09T10:49:23Z INFO  wgputoy::bind] Destroying bindings

There seems to be already the same issue posted upstream gfx-rs/wgpu#5347 which points out:

It also seems wrong that we're reporting a validation error with a cause of an internal error. It seems that this is…just an internal error.

Cheers,

Jun

Recommend Projects

  • React photo React

    A declarative, efficient, and flexible JavaScript library for building user interfaces.

  • Vue.js photo Vue.js

    🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.

  • Typescript photo Typescript

    TypeScript is a superset of JavaScript that compiles to clean JavaScript output.

  • TensorFlow photo TensorFlow

    An Open Source Machine Learning Framework for Everyone

  • Django photo Django

    The Web framework for perfectionists with deadlines.

  • D3 photo D3

    Bring data to life with SVG, Canvas and HTML. 📊📈🎉

Recommend Topics

  • javascript

    JavaScript (JS) is a lightweight interpreted programming language with first-class functions.

  • web

    Some thing interesting about web. New door for the world.

  • server

    A server is a program made to process requests and deliver data to clients.

  • Machine learning

    Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.

  • Game

    Some thing interesting about game, make everyone happy.

Recommend Org

  • Facebook photo Facebook

    We are working to build community through open source technology. NB: members must have two-factor auth.

  • Microsoft photo Microsoft

    Open source projects and samples from Microsoft.

  • Google photo Google

    Google ❤️ Open Source for everyone.

  • D3 photo D3

    Data-Driven Documents codes.