search

compute-toys / wgpu-compute-toy

Cross-platform compute shader engine
https://compute.toys
MIT License
118 stars 15 forks source link

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

Closed jmatsushita closed 4 months ago

jmatsushita commented 8 months ago

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 https://github.com/gfx-rs/wgpu/issues/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

jmatsushita commented 8 months ago

Oh I see. There's a conflict with the compute-toy prelude's assert function defined here: https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/lib.rs#L415-L419

Renaming that function to assert_toy fixes the error. I'll submit a PR.