archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

fix(namer): escape, rather than strip, non-ASCII ident. characters

Browse Source 
Escape non-ASCII identifier characters with `write!(…, "u{:04x}", …)`,
surrounding with `_` as appropriate. This solves (1) a debugging issue
where stripped characters would otherwise be invisible, and (2) failure
to re-validate that stripped identifiers didn't start with an ASCII
digit.

I've confirmed that this fixes [bug
1978197](https://bugzilla.mozilla.org/show_bug.cgi?id=1978197) on the
Firefox side.
This commit is contained in:
Erich Gubler 2025-07-22 11:08:39 -04:00
parent 00ea850d79
commit 90702156af
37 changed files with 276 additions and 124 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
CHANGELOG.md
View File

@ -237,6 +237,7 @@ By @cwfitzgerald in [#8162](https://github.com/gfx-rs/wgpu/pull/8162).
- [wgsl-in] Allow a trailing comma in `@blend_src(…)` attributes. By @ErichDonGubler in [#8137](https://github.com/gfx-rs/wgpu/pull/8137).
- [wgsl-in] Allow a trailing comma in the list of `case` values inside a `switch`. By @reima in [#8165](https://github.com/gfx-rs/wgpu/pull/8165).
- Escape, rather than strip, identifiers with Unicode. By @ErichDonGubler in [7995](https://github.com/gfx-rs/wgpu/pull/7995).
### Documentation

29
naga/src/proc/namer.rs
View File

@ -118,20 +118,31 @@ impl Namer {
{
Cow::Borrowed(string)
} else {
let mut filtered = string
.chars()
.filter(|&c| c.is_ascii_alphanumeric() || c == '_')
.fold(String::new(), |mut s, c| {
if s.ends_with('_') && c == '_' {
return s;
}
let mut filtered = string.chars().fold(String::new(), |mut s, c| {
let had_underscore_at_end = s.ends_with('_');
if had_underscore_at_end && c == '_' {
return s;
}
if c.is_ascii_alphanumeric() || c == '_' {
s.push(c);
s
});
} else {
use core::fmt::Write as _;
if !s.is_empty() && !had_underscore_at_end {
s.push('_');
}
write!(s, "u{:04x}_", c as u32).unwrap();
}
s
});
let stripped_len = filtered.trim_end_matches(SEPARATOR).len();
filtered.truncate(stripped_len);
if filtered.is_empty() {
filtered.push_str("unnamed");
} else if filtered.starts_with(|c: char| c.is_ascii_digit()) {
unreachable!(
"internal error: invalid identifier starting with ASCII digit {:?}",
filtered.chars().nth(0)
)
}
Cow::Owned(filtered)
};

14
naga/tests/in/wgsl/7995-unicode-idents.wgsl Normal file
View File

@ -0,0 +1,14 @@
// NOTE: This allows us to suppress compaction below, to force the handling of identifiers
// containing Unicode.
@group(0) @binding(0)
var<storage> asdf: f32;
fn compute() -> f32 {
let θ2 = asdf + 9001.0;
return θ2;
}
@compute @workgroup_size(1, 1)
fn main() {
compute();
}

4
naga/tests/out/glsl/spv-do-while.main.Fragment.glsl
View File

@ -4,7 +4,7 @@ precision highp float;
precision highp int;
void fb1_(inout bool cond) {
void f_u0028_b1_u003b(inout bool cond) {
bool loop_init = true;
while(true) {
if (!loop_init) {
@ -22,7 +22,7 @@ void fb1_(inout bool cond) {
void main_1() {
bool param = false;
param = false;
fb1_(param);
f_u0028_b1_u003b(param);
return;
}

21
naga/tests/out/glsl/wgsl-7995-unicode-idents.main.Compute.glsl Normal file
View File

@ -0,0 +1,21 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(std430) readonly buffer type_block_0Compute { float _group_0_binding_0_cs; };
float compute() {
float _e1 = _group_0_binding_0_cs;
float u03b8_2_ = (_e1 + 9001.0);
return u03b8_2_;
}
void main() {
float _e0 = compute();
return;
}

6
naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl
View File

@ -5,11 +5,11 @@ precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
};
@ -50,7 +50,7 @@ void main() {
int new = floatBitsToInt((intBitsToFloat(_e14) + 1.0));
uint _e20 = i;
int _e22 = old;
_atomic_compare_exchange_resultSint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_0_cs[_e20], _e22, new);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e23; _e23.old_value = atomicCompSwap(_group_0_binding_0_cs[_e20], _e22, new);
_e23.exchanged = (_e23.old_value == _e22);
old = _e23.old_value;
exchanged = _e23.exchanged;

6
naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl
View File

@ -5,11 +5,11 @@ precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
};
@ -50,7 +50,7 @@ void main() {
uint new = floatBitsToUint((uintBitsToFloat(_e14) + 1.0));
uint _e20 = i_1;
uint _e22 = old_1;
_atomic_compare_exchange_resultUint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_1_cs[_e20], _e22, new);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e23; _e23.old_value = atomicCompSwap(_group_0_binding_1_cs[_e20], _e22, new);
_e23.exchanged = (_e23.old_value == _e22);
old_1 = _e23.old_value;
exchanged_1 = _e23.exchanged;

20
naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
View File

@ -9,11 +9,11 @@ struct Struct {
uint atomic_scalar;
int atomic_arr[2];
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
};
@ -135,21 +135,21 @@ void main() {
int _e295 = atomicExchange(workgroup_atomic_arr[1], 1);
uint _e299 = atomicExchange(workgroup_struct.atomic_scalar, 1u);
int _e304 = atomicExchange(workgroup_struct.atomic_arr[1], 1);
_atomic_compare_exchange_resultUint4_ _e308; _e308.old_value = atomicCompSwap(_group_0_binding_0_cs, 1u, 2u);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e308; _e308.old_value = atomicCompSwap(_group_0_binding_0_cs, 1u, 2u);
_e308.exchanged = (_e308.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e313; _e313.old_value = atomicCompSwap(_group_0_binding_1_cs[1], 1, 2);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e313; _e313.old_value = atomicCompSwap(_group_0_binding_1_cs[1], 1, 2);
_e313.exchanged = (_e313.old_value == 1);
_atomic_compare_exchange_resultUint4_ _e318; _e318.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e318; _e318.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_scalar, 1u, 2u);
_e318.exchanged = (_e318.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e324; _e324.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_arr[1], 1, 2);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e324; _e324.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_arr[1], 1, 2);
_e324.exchanged = (_e324.old_value == 1);
_atomic_compare_exchange_resultUint4_ _e328; _e328.old_value = atomicCompSwap(workgroup_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e328; _e328.old_value = atomicCompSwap(workgroup_atomic_scalar, 1u, 2u);
_e328.exchanged = (_e328.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e333; _e333.old_value = atomicCompSwap(workgroup_atomic_arr[1], 1, 2);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e333; _e333.old_value = atomicCompSwap(workgroup_atomic_arr[1], 1, 2);
_e333.exchanged = (_e333.old_value == 1);
_atomic_compare_exchange_resultUint4_ _e338; _e338.old_value = atomicCompSwap(workgroup_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e338; _e338.old_value = atomicCompSwap(workgroup_struct.atomic_scalar, 1u, 2u);
_e338.exchanged = (_e338.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
_e344.exchanged = (_e344.old_value == 1);
return;
}

4
naga/tests/out/hlsl/spv-do-while.hlsl
View File

@ -1,4 +1,4 @@
void fb1_(inout bool cond)
void f_u0028_b1_u003b(inout bool cond)
{
uint2 loop_bound = uint2(4294967295u, 4294967295u);
bool loop_init = true;
@ -22,7 +22,7 @@ void main_1()
bool param = (bool)0;
param = false;
fb1_(param);
f_u0028_b1_u003b(param);
return;
}

2
naga/tests/out/hlsl/spv-fetch_depth.hlsl
View File

@ -19,7 +19,7 @@ void function()
}
[numthreads(32, 1, 1)]
void cullfetch_depth()
void cull_u003a_u003a_fetch_depth()
{
function();
}

2
naga/tests/out/hlsl/spv-fetch_depth.ron
View File

@ -5,7 +5,7 @@
],
compute:[
(
entry_point:"cullfetch_depth",
entry_point:"cull_u003a_u003a_fetch_depth",
target_profile:"cs_5_1",
),
],

15
naga/tests/out/hlsl/wgsl-7995-unicode-idents.hlsl Normal file
View File

@ -0,0 +1,15 @@
ByteAddressBuffer asdf : register(t0);
float compute()
{
float _e1 = asfloat(asdf.Load(0));
float u03b8_2_ = (_e1 + 9001.0);
return u03b8_2_;
}
[numthreads(1, 1, 1)]
void main()
{
const float _e0 = compute();
return;
}

12
naga/tests/out/hlsl/wgsl-7995-unicode-idents.ron Normal file
View File

@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"main",
target_profile:"cs_5_1",
),
],
)

8
naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl
View File

@ -5,13 +5,13 @@ struct NagaConstants {
};
ConstantBuffer<NagaConstants> _NagaConstants: register(b0, space1);
struct _atomic_compare_exchange_resultSint8_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e {
int64_t old_value;
bool exchanged;
int _end_pad_0;
};
struct _atomic_compare_exchange_resultUint8_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e {
uint64_t old_value;
bool exchanged;
int _end_pad_0;
@ -63,7 +63,7 @@ void test_atomic_compare_exchange_i64_()
int64_t new_ = (_e14 + 10L);
uint _e19 = i;
int64_t _e21 = old;
_atomic_compare_exchange_resultSint8_ _e22; arr_i64_.InterlockedCompareExchange64(_e19*8, _e21, new_, _e22.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e22; arr_i64_.InterlockedCompareExchange64(_e19*8, _e21, new_, _e22.old_value);
_e22.exchanged = (_e22.old_value == _e21);
old = _e22.old_value;
exchanged = _e22.exchanged;
@ -115,7 +115,7 @@ void test_atomic_compare_exchange_u64_()
uint64_t new_1 = (_e14 + 10uL);
uint _e19 = i_1;
uint64_t _e21 = old_1;
_atomic_compare_exchange_resultUint8_ _e22; arr_u64_.InterlockedCompareExchange64(_e19*8, _e21, new_1, _e22.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e22; arr_u64_.InterlockedCompareExchange64(_e19*8, _e21, new_1, _e22.old_value);
_e22.exchanged = (_e22.old_value == _e21);
old_1 = _e22.old_value;
exchanged_1 = _e22.exchanged;

8
naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl
View File

@ -1,9 +1,9 @@
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
};
@ -54,7 +54,7 @@ void test_atomic_compare_exchange_i32_()
int new_ = asint((asfloat(_e14) + 1.0));
uint _e20 = i;
int _e22 = old;
_atomic_compare_exchange_resultSint4_ _e23; arr_i32_.InterlockedCompareExchange(_e20*4, _e22, new_, _e23.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e23; arr_i32_.InterlockedCompareExchange(_e20*4, _e22, new_, _e23.old_value);
_e23.exchanged = (_e23.old_value == _e22);
old = _e23.old_value;
exchanged = _e23.exchanged;
@ -106,7 +106,7 @@ void test_atomic_compare_exchange_u32_()
uint new_1 = asuint((asfloat(_e14) + 1.0));
uint _e20 = i_1;
uint _e22 = old_1;
_atomic_compare_exchange_resultUint4_ _e23; arr_u32_.InterlockedCompareExchange(_e20*4, _e22, new_1, _e23.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e23; arr_u32_.InterlockedCompareExchange(_e20*4, _e22, new_1, _e23.old_value);
_e23.exchanged = (_e23.old_value == _e22);
old_1 = _e23.old_value;
exchanged_1 = _e23.exchanged;

20
naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
View File

@ -10,13 +10,13 @@ struct Struct {
int64_t atomic_arr[2];
};
struct _atomic_compare_exchange_resultUint8_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e {
uint64_t old_value;
bool exchanged;
int _end_pad_0;
};
struct _atomic_compare_exchange_resultSint8_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e {
int64_t old_value;
bool exchanged;
int _end_pad_0;
@ -126,21 +126,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
int64_t _e279; InterlockedExchange(workgroup_atomic_arr[1], 1L, _e279);
uint64_t _e283; InterlockedExchange(workgroup_struct.atomic_scalar, 1uL, _e283);
int64_t _e288; InterlockedExchange(workgroup_struct.atomic_arr[1], 1L, _e288);
_atomic_compare_exchange_resultUint8_ _e292; storage_atomic_scalar.InterlockedCompareExchange64(0, 1uL, 2uL, _e292.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e292; storage_atomic_scalar.InterlockedCompareExchange64(0, 1uL, 2uL, _e292.old_value);
_e292.exchanged = (_e292.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e297; storage_atomic_arr.InterlockedCompareExchange64(8, 1L, 2L, _e297.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e297; storage_atomic_arr.InterlockedCompareExchange64(8, 1L, 2L, _e297.old_value);
_e297.exchanged = (_e297.old_value == 1L);
_atomic_compare_exchange_resultUint8_ _e302; storage_struct.InterlockedCompareExchange64(0, 1uL, 2uL, _e302.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e302; storage_struct.InterlockedCompareExchange64(0, 1uL, 2uL, _e302.old_value);
_e302.exchanged = (_e302.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e308; storage_struct.InterlockedCompareExchange64(8+8, 1L, 2L, _e308.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e308; storage_struct.InterlockedCompareExchange64(8+8, 1L, 2L, _e308.old_value);
_e308.exchanged = (_e308.old_value == 1L);
_atomic_compare_exchange_resultUint8_ _e312; InterlockedCompareExchange(workgroup_atomic_scalar, 1uL, 2uL, _e312.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e312; InterlockedCompareExchange(workgroup_atomic_scalar, 1uL, 2uL, _e312.old_value);
_e312.exchanged = (_e312.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e317; InterlockedCompareExchange(workgroup_atomic_arr[1], 1L, 2L, _e317.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e317; InterlockedCompareExchange(workgroup_atomic_arr[1], 1L, 2L, _e317.old_value);
_e317.exchanged = (_e317.old_value == 1L);
_atomic_compare_exchange_resultUint8_ _e322; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1uL, 2uL, _e322.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e322; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1uL, 2uL, _e322.old_value);
_e322.exchanged = (_e322.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
_e328.exchanged = (_e328.old_value == 1L);
return;
}

20
naga/tests/out/hlsl/wgsl-atomicOps.hlsl
View File

@ -3,12 +3,12 @@ struct Struct {
int atomic_arr[2];
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
};
@ -117,21 +117,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
int _e295; InterlockedExchange(workgroup_atomic_arr[1], int(1), _e295);
uint _e299; InterlockedExchange(workgroup_struct.atomic_scalar, 1u, _e299);
int _e304; InterlockedExchange(workgroup_struct.atomic_arr[1], int(1), _e304);
_atomic_compare_exchange_resultUint4_ _e308; storage_atomic_scalar.InterlockedCompareExchange(0, 1u, 2u, _e308.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e308; storage_atomic_scalar.InterlockedCompareExchange(0, 1u, 2u, _e308.old_value);
_e308.exchanged = (_e308.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e313; storage_atomic_arr.InterlockedCompareExchange(4, int(1), int(2), _e313.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e313; storage_atomic_arr.InterlockedCompareExchange(4, int(1), int(2), _e313.old_value);
_e313.exchanged = (_e313.old_value == int(1));
_atomic_compare_exchange_resultUint4_ _e318; storage_struct.InterlockedCompareExchange(0, 1u, 2u, _e318.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e318; storage_struct.InterlockedCompareExchange(0, 1u, 2u, _e318.old_value);
_e318.exchanged = (_e318.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e324; storage_struct.InterlockedCompareExchange(4+4, int(1), int(2), _e324.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e324; storage_struct.InterlockedCompareExchange(4+4, int(1), int(2), _e324.old_value);
_e324.exchanged = (_e324.old_value == int(1));
_atomic_compare_exchange_resultUint4_ _e328; InterlockedCompareExchange(workgroup_atomic_scalar, 1u, 2u, _e328.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e328; InterlockedCompareExchange(workgroup_atomic_scalar, 1u, 2u, _e328.old_value);
_e328.exchanged = (_e328.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e333; InterlockedCompareExchange(workgroup_atomic_arr[1], int(1), int(2), _e333.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e333; InterlockedCompareExchange(workgroup_atomic_arr[1], int(1), int(2), _e333.old_value);
_e333.exchanged = (_e333.old_value == int(1));
_atomic_compare_exchange_resultUint4_ _e338; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1u, 2u, _e338.old_value);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e338; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1u, 2u, _e338.old_value);
_e338.exchanged = (_e338.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e344; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], int(1), int(2), _e344.old_value);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e344; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], int(1), int(2), _e344.old_value);
_e344.exchanged = (_e344.old_value == int(1));
return;
}

4
naga/tests/out/msl/spv-do-while.msl
View File

@ -5,7 +5,7 @@
using metal::uint;
void fb1_(
void f_u0028_b1_u003b(
thread bool& cond
) {
uint2 loop_bound = uint2(4294967295u);
@ -29,7 +29,7 @@ void main_1(
) {
bool param = {};
param = false;
fb1_(param);
f_u0028_b1_u003b(param);
return;
}

2
naga/tests/out/msl/spv-fetch_depth.msl
View File

@ -22,7 +22,7 @@ void function(
return;
}
kernel void cullfetch_depth(
kernel void cull_u003a_u003a_fetch_depth(
device type_2& global [[user(fake0)]]
, device type_4 const& global_1 [[user(fake0)]]
, metal::depth2d<float, metal::access::sample> global_2 [[user(fake0)]]

21
naga/tests/out/msl/wgsl-7995-unicode-idents.msl Normal file
View File

@ -0,0 +1,21 @@
// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
float compute(
device float const& asdf
) {
float _e1 = asdf;
float u03b8_2_ = _e1 + 9001.0;
return u03b8_2_;
}
kernel void main_(
device float const& asdf [[user(fake0)]]
) {
float _e0 = compute(asdf);
return;
}

24
naga/tests/out/msl/wgsl-atomicCompareExchange.msl
View File

@ -10,19 +10,19 @@ struct type_2 {
struct type_4 {
metal::atomic_uint inner[128];
};
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
char _pad2[3];
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
char _pad2[3];
};
template <typename A>
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
int cmp,
int v
@ -31,10 +31,10 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
int cmp,
int v
@ -43,11 +43,11 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
uint cmp,
uint v
@ -56,10 +56,10 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
uint cmp,
uint v
@ -68,7 +68,7 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
}
constant uint SIZE = 128u;
@ -112,7 +112,7 @@ kernel void test_atomic_compare_exchange_i32_(
int new_ = as_type<int>(as_type<float>(_e14) + 1.0);
uint _e20 = i;
int _e22 = old;
_atomic_compare_exchange_resultSint4_ _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_i32_.inner[_e20], _e22, new_);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_i32_.inner[_e20], _e22, new_);
old = _e23.old_value;
exchanged = _e23.exchanged;
}
@ -163,7 +163,7 @@ kernel void test_atomic_compare_exchange_u32_(
uint new_1 = as_type<uint>(as_type<float>(_e14) + 1.0);
uint _e20 = i_1;
uint _e22 = old_1;
_atomic_compare_exchange_resultUint4_ _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_u32_.inner[_e20], _e22, new_1);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_u32_.inner[_e20], _e22, new_1);
old_1 = _e23.old_value;
exchanged_1 = _e23.exchanged;
}

36
naga/tests/out/msl/wgsl-atomicOps.msl
View File

@ -11,19 +11,19 @@ struct Struct {
metal::atomic_uint atomic_scalar;
type_2 atomic_arr;
};
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
char _pad2[3];
};
struct _atomic_compare_exchange_resultSint4_ {
struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
int old_value;
bool exchanged;
char _pad2[3];
};
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
uint cmp,
uint v
@ -32,10 +32,10 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
uint cmp,
uint v
@ -44,11 +44,11 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
int cmp,
int v
@ -57,10 +57,10 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
int cmp,
int v
@ -69,7 +69,7 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
}
struct cs_mainInput {
@ -182,13 +182,13 @@ kernel void cs_main(
int _e295 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
uint _e299 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
int _e304 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
_atomic_compare_exchange_resultUint4_ _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_resultUint4_ _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_resultUint4_ _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_resultUint4_ _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
return;
}

12
naga/tests/out/msl/wgsl-overrides-atomicCompareExchangeWeak.msl
View File

@ -4,14 +4,14 @@
using metal::uint;
struct _atomic_compare_exchange_resultUint4_ {
struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
uint old_value;
bool exchanged;
char _pad2[3];
};
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
uint cmp,
uint v
@ -20,10 +20,10 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
uint cmp,
uint v
@ -32,7 +32,7 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
}
constant int o = 2;
@ -44,6 +44,6 @@ kernel void f(
metal::atomic_store_explicit(&a, 0, metal::memory_order_relaxed);
}
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
_atomic_compare_exchange_resultUint4_ _e5 = naga_atomic_compare_exchange_weak_explicit(&a, 2u, 1u);
_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e5 = naga_atomic_compare_exchange_weak_explicit(&a, 2u, 1u);
return;
}

43
naga/tests/out/spv/wgsl-7995-unicode-idents.spvasm Normal file
View File

@ -0,0 +1,43 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 24
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %19 "main"
OpExecutionMode %19 LocalSize 1 1 1
OpDecorate %4 NonWritable
OpDecorate %4 DescriptorSet 0
OpDecorate %4 Binding 0
OpDecorate %5 Block
OpMemberDecorate %5 0 Offset 0
%2 = OpTypeVoid
%3 = OpTypeFloat 32
%5 = OpTypeStruct %3
%6 = OpTypePointer StorageBuffer %5
%4 = OpVariable %6 StorageBuffer
%9 = OpTypeFunction %3
%10 = OpTypePointer StorageBuffer %3
%12 = OpTypeInt 32 0
%11 = OpConstant %12 0
%14 = OpConstant %3 9001
%20 = OpTypeFunction %2
%8 = OpFunction %3 None %9
%7 = OpLabel
%13 = OpAccessChain %10 %4 %11
OpBranch %15
%15 = OpLabel
%16 = OpLoad %3 %13
%17 = OpFAdd %3 %16 %14
OpReturnValue %17
OpFunctionEnd
%19 = OpFunction %2 None %20
%18 = OpLabel
%21 = OpAccessChain %10 %4 %11
OpBranch %22
%22 = OpLabel
%23 = OpFunctionCall %3 %8
OpReturn
OpFunctionEnd

2
naga/tests/out/wgsl/spv-atomic_compare_exchange.wgsl
View File

@ -61,6 +61,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_compare_exchange() {
fn stage_u003a_u003a_test_atomic_compare_exchange() {
function();
}

2
naga/tests/out/wgsl/spv-atomic_exchange.wgsl
View File

@ -75,6 +75,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_exchange() {
fn stage_u003a_u003a_test_atomic_exchange() {
function();
}

2
naga/tests/out/wgsl/spv-atomic_i_add_sub.wgsl
View File

@ -21,6 +21,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_i_add_sub() {
fn stage_u003a_u003a_test_atomic_i_add_sub() {
function();
}

2
naga/tests/out/wgsl/spv-atomic_i_decrement.wgsl
View File

@ -32,6 +32,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_i_decrement() {
fn stage_u003a_u003a_test_atomic_i_decrement() {
function();
}

2
naga/tests/out/wgsl/spv-atomic_i_increment.wgsl
View File

@ -37,6 +37,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_i_increment() {
fn stage_u003a_u003a_test_atomic_i_increment() {
function();
}

2
naga/tests/out/wgsl/spv-atomic_load_and_store.wgsl
View File

@ -67,6 +67,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn stagetest_atomic_load_and_store() {
fn stage_u003a_u003a_test_atomic_load_and_store() {
function();
}

14
naga/tests/out/wgsl/spv-binding-arrays.runtime.wgsl
View File

@ -1,5 +1,5 @@
var<private> inputtexture_coordinates_1: vec2<f32>;
var<private> inputtexture_index_1: u32;
var<private> input_u002e_texture_coordinates_1: vec2<f32>;
var<private> input_u002e_texture_index_1: u32;
@group(0) @binding(0)
var textures: binding_array<texture_2d<f32>>;
@group(0) @binding(1)
@ -7,17 +7,17 @@ var linear_sampler: sampler;
var<private> entryPointParam_main: vec4<f32>;
fn main_1() {
let _e5 = inputtexture_coordinates_1;
let _e6 = inputtexture_index_1;
let _e5 = input_u002e_texture_coordinates_1;
let _e6 = input_u002e_texture_index_1;
let _e8 = textureSample(textures[_e6], linear_sampler, _e5);
entryPointParam_main = _e8;
return;
}
@fragment
fn main(@location(0) inputtexture_coordinates: vec2<f32>, @location(1) @interpolate(flat) inputtexture_index: u32) -> @location(0) vec4<f32> {
inputtexture_coordinates_1 = inputtexture_coordinates;
inputtexture_index_1 = inputtexture_index;
fn main(@location(0) input_u002e_texture_coordinates: vec2<f32>, @location(1) @interpolate(flat) input_u002e_texture_index: u32) -> @location(0) vec4<f32> {
input_u002e_texture_coordinates_1 = input_u002e_texture_coordinates;
input_u002e_texture_index_1 = input_u002e_texture_index;
main_1();
let _e5 = entryPointParam_main;
return _e5;

4
naga/tests/out/wgsl/spv-builtin-accessed-outside-entrypoint.wgsl
View File

@ -8,7 +8,7 @@ struct gl_PerVertex {
var<private> unnamed: gl_PerVertex = gl_PerVertex(vec4<f32>(0f, 0f, 0f, 1f), 1f, array<f32, 1>(), array<f32, 1>());
var<private> gl_VertexIndex_1: i32;
fn builtin_usage() {
fn builtin_usage_u0028_() {
let _e9 = gl_VertexIndex_1;
let _e12 = gl_VertexIndex_1;
unnamed.gl_Position = vec4<f32>(select(1f, -4f, (_e9 == 0i)), select(-1f, 4f, (_e12 == 2i)), 0f, 1f);
@ -16,7 +16,7 @@ fn builtin_usage() {
}
fn main_1() {
builtin_usage();
builtin_usage_u0028_();
return;
}

4
naga/tests/out/wgsl/spv-do-while.wgsl
View File

@ -1,4 +1,4 @@
fn fb1_(cond: ptr<function, bool>) {
fn f_u0028_b1_u003b(cond: ptr<function, bool>) {
loop {
continue;
continuing {
@ -13,7 +13,7 @@ fn main_1() {
var param: bool;
param = false;
fb1_((&param));
f_u0028_b1_u003b((&param));
return;
}

2
naga/tests/out/wgsl/spv-fetch_depth.wgsl
View File

@ -21,6 +21,6 @@ fn function() {
}
@compute @workgroup_size(32, 1, 1)
fn cullfetch_depth() {
fn cull_u003a_u003a_fetch_depth() {
function();
}

8
naga/tests/out/wgsl/spv-gather-cmp.wgsl
View File

@ -1,4 +1,4 @@
var<private> inputtexture_coordinates_1: vec2<f32>;
var<private> input_u002e_texture_coordinates_1: vec2<f32>;
@group(0) @binding(0)
var texture: texture_depth_2d;
@group(0) @binding(1)
@ -6,15 +6,15 @@ var depth_sampler: sampler_comparison;
var<private> entryPointParam_main: vec4<f32>;
fn main_1() {
let _e5 = inputtexture_coordinates_1;
let _e5 = input_u002e_texture_coordinates_1;
let _e6 = textureGatherCompare(texture, depth_sampler, _e5, 0.5f);
entryPointParam_main = _e6;
return;
}
@fragment
fn main(@location(0) inputtexture_coordinates: vec2<f32>) -> @location(0) vec4<f32> {
inputtexture_coordinates_1 = inputtexture_coordinates;
fn main(@location(0) input_u002e_texture_coordinates: vec2<f32>) -> @location(0) vec4<f32> {
input_u002e_texture_coordinates_1 = input_u002e_texture_coordinates;
main_1();
let _e3 = entryPointParam_main;
return _e3;

8
naga/tests/out/wgsl/spv-gather.wgsl
View File

@ -1,4 +1,4 @@
var<private> inputtexture_coordinates_1: vec2<f32>;
var<private> input_u002e_texture_coordinates_1: vec2<f32>;
@group(0) @binding(0)
var texture: texture_2d<f32>;
@group(0) @binding(1)
@ -6,15 +6,15 @@ var linear_sampler: sampler;
var<private> entryPointParam_main: vec4<f32>;
fn main_1() {
let _e4 = inputtexture_coordinates_1;
let _e4 = input_u002e_texture_coordinates_1;
let _e5 = textureGather(1, texture, linear_sampler, _e4);
entryPointParam_main = _e5;
return;
}
@fragment
fn main(@location(0) inputtexture_coordinates: vec2<f32>) -> @location(0) vec4<f32> {
inputtexture_coordinates_1 = inputtexture_coordinates;
fn main(@location(0) input_u002e_texture_coordinates: vec2<f32>) -> @location(0) vec4<f32> {
input_u002e_texture_coordinates_1 = input_u002e_texture_coordinates;
main_1();
let _e3 = entryPointParam_main;
return _e3;

14
naga/tests/out/wgsl/wgsl-7995-unicode-idents.wgsl Normal file
View File

@ -0,0 +1,14 @@
@group(0) @binding(0)
var<storage> asdf: f32;
fn compute() -> f32 {
let _e1 = asdf;
let u03b8_2_ = (_e1 + 9001f);
return u03b8_2_;
}
@compute @workgroup_size(1, 1, 1)
fn main() {
let _e0 = compute();
return;
}