mirror of
https://github.com/gfx-rs/wgpu.git
synced 2025-02-27 14:23:40 +00:00
9742f1616c
Use the local (not global) invocation id to decide which invocation should do the initialization, so that every workgroup gets initialized, not just the first.
105 lines
2.8 KiB
Plaintext
105 lines
2.8 KiB
Plaintext
// language: metal2.1
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using metal::uint;
|
|
|
|
struct VertexOutput {
|
|
metal::float4 position;
|
|
float _varying;
|
|
};
|
|
struct FragmentOutput {
|
|
float depth;
|
|
uint sample_mask;
|
|
float color;
|
|
};
|
|
struct type_4 {
|
|
uint inner[1];
|
|
};
|
|
struct Input1_ {
|
|
uint index;
|
|
};
|
|
struct Input2_ {
|
|
uint index;
|
|
};
|
|
|
|
struct vertex_Input {
|
|
uint color [[attribute(10)]];
|
|
};
|
|
struct vertex_Output {
|
|
metal::float4 position [[position, invariant]];
|
|
float _varying [[user(loc1), center_perspective]];
|
|
float _point_size [[point_size]];
|
|
};
|
|
vertex vertex_Output vertex_(
|
|
vertex_Input varyings [[stage_in]]
|
|
, uint vertex_index [[vertex_id]]
|
|
, uint instance_index [[instance_id]]
|
|
) {
|
|
const auto color = varyings.color;
|
|
uint tmp = (vertex_index + instance_index) + color;
|
|
const auto _tmp = VertexOutput {metal::float4(1.0), static_cast<float>(tmp)};
|
|
return vertex_Output { _tmp.position, _tmp._varying, 1.0 };
|
|
}
|
|
|
|
|
|
struct fragment_Input {
|
|
float _varying [[user(loc1), center_perspective]];
|
|
};
|
|
struct fragment_Output {
|
|
float depth [[depth(any)]];
|
|
uint sample_mask [[sample_mask]];
|
|
float color [[color(0)]];
|
|
};
|
|
fragment fragment_Output fragment_(
|
|
fragment_Input varyings_1 [[stage_in]]
|
|
, metal::float4 position [[position]]
|
|
, bool front_facing [[front_facing]]
|
|
, uint sample_index [[sample_id]]
|
|
, uint sample_mask [[sample_mask]]
|
|
) {
|
|
const VertexOutput in = { position, varyings_1._varying };
|
|
uint mask = sample_mask & (1u << sample_index);
|
|
float color_1 = front_facing ? 1.0 : 0.0;
|
|
const auto _tmp = FragmentOutput {in._varying, mask, color_1};
|
|
return fragment_Output { _tmp.depth, _tmp.sample_mask, _tmp.color };
|
|
}
|
|
|
|
|
|
struct compute_Input {
|
|
};
|
|
kernel void compute_(
|
|
metal::uint3 global_id [[thread_position_in_grid]]
|
|
, metal::uint3 local_id [[thread_position_in_threadgroup]]
|
|
, uint local_index [[thread_index_in_threadgroup]]
|
|
, metal::uint3 wg_id [[threadgroup_position_in_grid]]
|
|
, metal::uint3 num_wgs [[threadgroups_per_grid]]
|
|
, threadgroup type_4& output
|
|
) {
|
|
if (metal::all(local_id == metal::uint3(0u))) {
|
|
output = {};
|
|
}
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
output.inner[0] = (((global_id.x + local_id.x) + local_index) + wg_id.x) + num_wgs.x;
|
|
return;
|
|
}
|
|
|
|
|
|
struct vertex_two_structsInput {
|
|
};
|
|
struct vertex_two_structsOutput {
|
|
metal::float4 member_3 [[position, invariant]];
|
|
float _point_size [[point_size]];
|
|
};
|
|
vertex vertex_two_structsOutput vertex_two_structs(
|
|
uint index_1 [[vertex_id]]
|
|
, uint index_2 [[instance_id]]
|
|
) {
|
|
const Input1_ in1_ = { index_1 };
|
|
const Input2_ in2_ = { index_2 };
|
|
uint index = {};
|
|
index = 2u;
|
|
uint _e8 = index;
|
|
return vertex_two_structsOutput { metal::float4(static_cast<float>(in1_.index), static_cast<float>(in2_.index), static_cast<float>(_e8), 0.0), 1.0 };
|
|
}
|