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

[naga] Ensure globals in wgsl snapshot tests are reachable from an entry point

Browse Source 
Convert tabs to spaces in access.wgsl
This commit is contained in:
Andy Leiserson 2025-05-15 15:28:38 -07:00 committed by Jim Blandy
parent fd6f16f598
commit 25636e274b
43 changed files with 3609 additions and 2912 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

71
naga/tests/in/wgsl/abstract-types-var.wgsl
View File

@ -80,6 +80,76 @@ var<private> iavipai = array(vec3(1));
var<private> iavfpai = array(vec3(1));
var<private> iavfpaf = array(vec3(1.0));
fn globals() {
_ = xvipaiai;
_ = xvupaiai;
_ = xvfpaiai;
_ = xvfpafaf;
_ = xvfpaiaf;
_ = xvupuai;
_ = xvupaiu;
_ = xvuuai;
_ = xvuaiu;
_ = xvip____;
_ = xvup____;
_ = xvfp____;
_ = xmfp____;
_ = xmfpaiaiaiai;
_ = xmfpafaiaiai;
_ = xmfpaiafaiai;
_ = xmfpaiaiafai;
_ = xmfpaiaiaiaf;
_ = xvispai;
_ = xvfspaf;
_ = xvis_ai;
_ = xvus_ai;
_ = xvfs_ai;
_ = xvfs_af;
_ = xafafaf;
_ = xafaiai;
_ = xaipaiai;
_ = xaupaiai;
_ = xafpaiaf;
_ = xafpafai;
_ = xafpafaf;
_ = xavipai;
_ = xavfpai;
_ = xavfpaf;
// Construction with splats
_ = xvisai;
_ = xvusai;
_ = xvfsai;
_ = xvfsaf;
_ = ivispai;
_ = ivfspaf;
_ = ivis_ai;
_ = ivus_ai;
_ = ivfs_ai;
_ = ivfs_af;
_ = iafafaf;
_ = iafaiai;
_ = iaipaiai;
_ = iafpafaf;
_ = iafpaiaf;
_ = iafpafai;
_ = iavipai;
_ = iavfpai;
_ = iavfpaf;
}
fn all_constant_arguments() {
var xvipaiai: vec2<i32> = vec2(42, 43);
var xvupaiai: vec2<u32> = vec2(44, 45);
@ -278,6 +348,7 @@ fn mixed_constant_and_runtime_arguments() {
@compute @workgroup_size(1)
fn main() {
globals();
all_constant_arguments();
mixed_constant_and_runtime_arguments();
}

83
naga/tests/in/wgsl/access.wgsl
View File

@ -6,26 +6,26 @@ struct GlobalConst {
c: i32,
}
// tests msl padding insertion for global constants
var<private> global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0);
var<private> msl_padding_global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0);
struct AlignedWrapper {
@align(8) value: i32
@align(8) value: i32
}
struct Bar {
_matrix: mat4x3<f32>,
matrix_array: array<mat2x2<f32>, 2>,
atom: atomic<i32>,
atom_arr: array<atomic<i32>, 10>,
arr: array<vec2<u32>, 2>,
data: array<AlignedWrapper>,
_matrix: mat4x3<f32>,
matrix_array: array<mat2x2<f32>, 2>,
atom: atomic<i32>,
atom_arr: array<atomic<i32>, 10>,
arr: array<vec2<u32>, 2>,
data: array<AlignedWrapper>,
}
@group(0) @binding(0)
var<storage,read_write> bar: Bar;
struct Baz {
m: mat3x2<f32>,
m: mat3x2<f32>,
}
@group(0) @binding(1)
@ -35,11 +35,11 @@ var<uniform> baz: Baz;
var<storage,read_write> qux: vec2<i32>;
fn test_matrix_within_struct_accesses() {
var idx = 1;
var idx = 1;
idx--;
// loads
// loads
let l0 = baz.m;
let l1 = baz.m[0];
let l2 = baz.m[idx];
@ -52,7 +52,7 @@ fn test_matrix_within_struct_accesses() {
idx++;
// stores
// stores
t.m = mat3x2<f32>(vec2<f32>(6.0), vec2<f32>(5.0), vec2<f32>(4.0));
t.m[0] = vec2<f32>(9.0);
t.m[idx] = vec2<f32>(90.0);
@ -63,18 +63,18 @@ fn test_matrix_within_struct_accesses() {
}
struct MatCx2InArray {
am: array<mat4x2<f32>, 2>,
am: array<mat4x2<f32>, 2>,
}
@group(0) @binding(3)
var<uniform> nested_mat_cx2: MatCx2InArray;
fn test_matrix_within_array_within_struct_accesses() {
var idx = 1;
var idx = 1;
idx--;
// loads
// loads
let l0 = nested_mat_cx2.am;
let l1 = nested_mat_cx2.am[0];
let l2 = nested_mat_cx2.am[0][0];
@ -88,7 +88,7 @@ fn test_matrix_within_array_within_struct_accesses() {
idx++;
// stores
// stores
t.am = array<mat4x2<f32>, 2>();
t.am[0] = mat4x2<f32>(vec2<f32>(8.0), vec2<f32>(7.0), vec2<f32>(6.0), vec2<f32>(5.0));
t.am[0][0] = vec2<f32>(9.0);
@ -114,41 +114,42 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let baz: f32 = foo;
foo = 1.0;
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
_ = msl_padding_global_const;
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
// test storage loads
let _matrix = bar._matrix;
let arr = bar.arr;
let index = 3u;
let b = bar._matrix[index].x;
let a = bar.data[arrayLength(&bar.data) - 2u].value;
let c = qux;
let _matrix = bar._matrix;
let arr = bar.arr;
let index = 3u;
let b = bar._matrix[index].x;
let a = bar.data[arrayLength(&bar.data) - 2u].value;
let c = qux;
// test pointer types
let data_pointer: ptr<storage, i32, read_write> = &bar.data[0].value;
let foo_value = read_from_private(&foo);
// test pointer types
let data_pointer: ptr<storage, i32, read_write> = &bar.data[0].value;
let foo_value = read_from_private(&foo);
// test array indexing
var c2 = array<i32, 5>(a, i32(b), 3, 4, 5);
c2[vi + 1u] = 42;
let value = c2[vi];
// test array indexing
var c2 = array<i32, 5>(a, i32(b), 3, 4, 5);
c2[vi + 1u] = 42;
let value = c2[vi];
test_arr_as_arg(array<array<f32, 10>, 5>());
test_arr_as_arg(array<array<f32, 10>, 5>());
return vec4<f32>(_matrix * vec4<f32>(vec4<i32>(value)), 2.0);
return vec4<f32>(_matrix * vec4<f32>(vec4<i32>(value)), 2.0);
}
@fragment
fn foo_frag() -> @location(0) vec4<f32> {
// test storage stores
bar._matrix[1].z = 1.0;
bar._matrix = mat4x3<f32>(vec3<f32>(0.0), vec3<f32>(1.0), vec3<f32>(2.0), vec3<f32>(3.0));
bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
bar.data[1].value = 1;
qux = vec2<i32>();
// test storage stores
bar._matrix[1].z = 1.0;
bar._matrix = mat4x3<f32>(vec3<f32>(0.0), vec3<f32>(1.0), vec3<f32>(2.0), vec3<f32>(3.0));
bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
bar.data[1].value = 1;
qux = vec2<i32>();
return vec4<f32>(0.0);
return vec4<f32>(0.0);
}
fn assign_through_ptr_fn(p: ptr<function, u32>) {
@ -163,7 +164,7 @@ fn assign_through_ptr() {
var val = 33u;
assign_through_ptr_fn(&val);
var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
assign_array_through_ptr_fn(&arr);
}

5
naga/tests/in/wgsl/debug-symbol-large-source.wgsl
View File

@ -7418,6 +7418,11 @@ fn color23(p: vec2<f32>) -> vec3<f32> {
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
_ = t_diffuse;
_ = s_diffuse;
_ = t_normal;
_ = s_normal;
color23(vec2(1, 2));
var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));

7
naga/tests/in/wgsl/debug-symbol-terrain.wgsl
View File

@ -186,7 +186,7 @@ fn gen_terrain_fragment(in: GenVertexOutput) -> GenFragmentOutput {
let v = terrain_vertex(p, gen_data.min_max_height);
var vert_component: f32 = 0.;
switch comp_index {
case 0u: { vert_component = v.position.x; }
case 1u: { vert_component = v.position.y; }
@ -275,6 +275,11 @@ fn color23(p: vec2<f32>) -> vec3<f32> {
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
_ = t_diffuse;
_ = s_diffuse;
_ = t_normal;
_ = s_normal;
color23(vec2(1, 2));
var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));

3
naga/tests/in/wgsl/f16.wgsl
View File

@ -51,6 +51,9 @@ var<storage, read_write> output: UniformCompatible;
var<storage, read_write> output_arrays: StorageCompatible;
fn f16_function(x: f16) -> f16 {
_ = private_variable;
var l: LayoutTest;
var val: f16 = f16(constant_variable);
// A number too big for f16
val += 1h - 33333h;

1
naga/tests/in/wgsl/f64.wgsl
View File

@ -2,6 +2,7 @@ var<private> v: f64 = 1lf;
const k: f64 = 2.0lf;
fn f(x: f64) -> f64 {
_ = v;
let y: f64 = 3e1lf + 4.0e2lf;
var z = y + f64(5);
var w = -1.0lf;

1
naga/tests/in/wgsl/image.wgsl
View File

@ -29,6 +29,7 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
let value5 = textureLoad(image_array_src, itc, local_id.z, i32(local_id.z) + 1);
let value6 = textureLoad(image_array_src, itc, i32(local_id.z), i32(local_id.z) + 1);
let value7 = textureLoad(image_1d_src, i32(local_id.x), i32(local_id.z));
let value8 = textureLoad(image_dup_src, i32(local_id.x));
// loads with uvec2 coords.
let value1u = textureLoad(image_mipmapped_src, vec2<u32>(itc), i32(local_id.z));
let value2u = textureLoad(image_multisampled_src, vec2<u32>(itc), i32(local_id.z));

1
naga/tests/in/wgsl/int64.wgsl
View File

@ -39,6 +39,7 @@ var<storage, read_write> output: UniformCompatible;
var<storage, read_write> output_arrays: StorageCompatible;
fn int64_function(x: i64) -> i64 {
_ = private_variable;
var val: i64 = i64(constant_variable);
// A number too big for i32
val += 31li - 1002003004005006li + -0x7fffffffffffffffli;

5
naga/tests/in/wgsl/overrides.wgsl
View File

@ -27,4 +27,9 @@ fn main() {
var gain_x_100 = gain_x_10 * 10.;
store_override = gain;
_ = specular_param;
_ = width;
_ = inferred_f32;
_ = auto_conversion;
}

141
naga/tests/out/analysis/wgsl-access.info.ron
View File

@ -3606,7 +3606,7 @@
may_kill: false,
sampling_set: [],
global_uses: [
(""),
("READ"),
("READ | QUERY"),
("READ"),
("READ"),
@ -3673,6 +3673,27 @@
requirements: (""),
),
ref_count: 1,
assignable_global: Some(0),
ty: Value(Pointer(
base: 3,
space: Private,
)),
),
(
uniformity: (
non_uniform_result: Some(5),
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(7),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Value(Pointer(
base: 14,
@ -3683,7 +3704,7 @@
),
(
uniformity: (
non_uniform_result: Some(5),
non_uniform_result: Some(7),
requirements: (""),
),
ref_count: 1,
@ -3697,7 +3718,7 @@
),
(
uniformity: (
non_uniform_result: Some(5),
non_uniform_result: Some(7),
requirements: (""),
),
ref_count: 1,
@ -3706,7 +3727,7 @@
),
(
uniformity: (
non_uniform_result: Some(8),
non_uniform_result: Some(10),
requirements: (""),
),
ref_count: 1,
@ -3720,7 +3741,7 @@
),
(
uniformity: (
non_uniform_result: Some(8),
non_uniform_result: Some(10),
requirements: (""),
),
ref_count: 1,
@ -3734,7 +3755,7 @@
),
(
uniformity: (
non_uniform_result: Some(8),
non_uniform_result: Some(10),
requirements: (""),
),
ref_count: 0,
@ -3755,7 +3776,7 @@
),
(
uniformity: (
non_uniform_result: Some(12),
non_uniform_result: Some(14),
requirements: (""),
),
ref_count: 1,
@ -3769,7 +3790,7 @@
),
(
uniformity: (
non_uniform_result: Some(12),
non_uniform_result: Some(14),
requirements: (""),
),
ref_count: 1,
@ -3783,7 +3804,7 @@
),
(
uniformity: (
non_uniform_result: Some(12),
non_uniform_result: Some(14),
requirements: (""),
),
ref_count: 1,
@ -3801,7 +3822,7 @@
),
(
uniformity: (
non_uniform_result: Some(12),
non_uniform_result: Some(14),
requirements: (""),
),
ref_count: 1,
@ -3819,7 +3840,7 @@
),
(
uniformity: (
non_uniform_result: Some(12),
non_uniform_result: Some(14),
requirements: (""),
),
ref_count: 1,
@ -3829,34 +3850,6 @@
width: 4,
))),
),
(
uniformity: (
non_uniform_result: Some(17),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Value(Pointer(
base: 14,
space: Storage(
access: ("LOAD | STORE"),
),
)),
),
(
uniformity: (
non_uniform_result: Some(17),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Value(Pointer(
base: 13,
space: Storage(
access: ("LOAD | STORE"),
),
)),
),
(
uniformity: (
non_uniform_result: Some(19),
@ -3887,7 +3880,35 @@
),
(
uniformity: (
non_uniform_result: Some(19),
non_uniform_result: Some(21),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Value(Pointer(
base: 14,
space: Storage(
access: ("LOAD | STORE"),
),
)),
),
(
uniformity: (
non_uniform_result: Some(21),
requirements: (""),
),
ref_count: 1,
assignable_global: Some(1),
ty: Value(Pointer(
base: 13,
space: Storage(
access: ("LOAD | STORE"),
),
)),
),
(
uniformity: (
non_uniform_result: Some(21),
requirements: (""),
),
ref_count: 1,
@ -3911,7 +3932,7 @@
),
(
uniformity: (
non_uniform_result: Some(19),
non_uniform_result: Some(21),
requirements: (""),
),
ref_count: 1,
@ -3923,7 +3944,7 @@
),
(
uniformity: (
non_uniform_result: Some(17),
non_uniform_result: Some(19),
requirements: (""),
),
ref_count: 1,
@ -3937,7 +3958,7 @@
),
(
uniformity: (
non_uniform_result: Some(17),
non_uniform_result: Some(19),
requirements: (""),
),
ref_count: 1,
@ -3951,7 +3972,7 @@
),
(
uniformity: (
non_uniform_result: Some(17),
non_uniform_result: Some(19),
requirements: (""),
),
ref_count: 1,
@ -3960,7 +3981,7 @@
),
(
uniformity: (
non_uniform_result: Some(27),
non_uniform_result: Some(29),
requirements: (""),
),
ref_count: 1,
@ -3974,7 +3995,7 @@
),
(
uniformity: (
non_uniform_result: Some(27),
non_uniform_result: Some(29),
requirements: (""),
),
ref_count: 0,
@ -3983,7 +4004,7 @@
),
(
uniformity: (
non_uniform_result: Some(29),
non_uniform_result: Some(31),
requirements: (""),
),
ref_count: 1,
@ -3997,7 +4018,7 @@
),
(
uniformity: (
non_uniform_result: Some(29),
non_uniform_result: Some(31),
requirements: (""),
),
ref_count: 1,
@ -4011,7 +4032,7 @@
),
(
uniformity: (
non_uniform_result: Some(29),
non_uniform_result: Some(31),
requirements: (""),
),
ref_count: 1,
@ -4025,7 +4046,7 @@
),
(
uniformity: (
non_uniform_result: Some(29),
non_uniform_result: Some(31),
requirements: (""),
),
ref_count: 0,
@ -4048,7 +4069,7 @@
),
(
uniformity: (
non_uniform_result: Some(12),
non_uniform_result: Some(14),
requirements: (""),
),
ref_count: 1,
@ -4096,7 +4117,7 @@
),
(
uniformity: (
non_uniform_result: Some(17),
non_uniform_result: Some(19),
requirements: (""),
),
ref_count: 1,
@ -4105,7 +4126,7 @@
),
(
uniformity: (
non_uniform_result: Some(39),
non_uniform_result: Some(41),
requirements: (""),
),
ref_count: 3,
@ -4138,7 +4159,7 @@
),
(
uniformity: (
non_uniform_result: Some(39),
non_uniform_result: Some(41),
requirements: (""),
),
ref_count: 1,
@ -4162,7 +4183,7 @@
),
(
uniformity: (
non_uniform_result: Some(39),
non_uniform_result: Some(41),
requirements: (""),
),
ref_count: 1,
@ -4174,7 +4195,7 @@
),
(
uniformity: (
non_uniform_result: Some(39),
non_uniform_result: Some(41),
requirements: (""),
),
ref_count: 1,
@ -4201,7 +4222,7 @@
),
(
uniformity: (
non_uniform_result: Some(39),
non_uniform_result: Some(41),
requirements: (""),
),
ref_count: 1,
@ -4216,7 +4237,7 @@
),
(
uniformity: (
non_uniform_result: Some(39),
non_uniform_result: Some(41),
requirements: (""),
),
ref_count: 1,
@ -4231,7 +4252,7 @@
),
(
uniformity: (
non_uniform_result: Some(5),
non_uniform_result: Some(7),
requirements: (""),
),
ref_count: 1,
@ -4258,7 +4279,7 @@
),
(
uniformity: (
non_uniform_result: Some(5),
non_uniform_result: Some(7),
requirements: (""),
),
ref_count: 1,

36
naga/tests/out/analysis/wgsl-overrides.info.ron
View File

@ -161,6 +161,42 @@
assignable_global: None,
ty: Handle(1),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Handle(1),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Handle(1),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Handle(1),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 0,
assignable_global: None,
ty: Handle(2),
),
],
sampling: [],
dual_source_blending: false,

164
naga/tests/out/glsl/wgsl-abstract-types-var.main.Compute.glsl
View File

@ -5,6 +5,169 @@ precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
ivec2 xvipaiai_1 = ivec2(42, 43);
uvec2 xvupaiai_1 = uvec2(44u, 45u);
vec2 xvfpaiai_1 = vec2(46.0, 47.0);
vec2 xvfpafaf_1 = vec2(48.0, 49.0);
vec2 xvfpaiaf_1 = vec2(48.0, 49.0);
uvec2 xvupuai_2 = uvec2(42u, 43u);
uvec2 xvupaiu_2 = uvec2(42u, 43u);
uvec2 xvuuai_2 = uvec2(42u, 43u);
uvec2 xvuaiu_2 = uvec2(42u, 43u);
ivec2 xvip_1 = ivec2(0, 0);
uvec2 xvup_1 = uvec2(0u, 0u);
vec2 xvfp_1 = vec2(0.0, 0.0);
mat2x2 xmfp_1 = mat2x2(vec2(0.0, 0.0), vec2(0.0, 0.0));
mat2x2 xmfpaiaiaiai_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
mat2x2 xmfpafaiaiai_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
mat2x2 xmfpaiafaiai_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
mat2x2 xmfpaiaiafai_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
mat2x2 xmfpaiaiaiaf_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0));
ivec2 xvispai_1 = ivec2(1);
vec2 xvfspaf_1 = vec2(1.0);
ivec2 xvis_ai_1 = ivec2(1);
uvec2 xvus_ai_1 = uvec2(1u);
vec2 xvfs_ai_1 = vec2(1.0);
vec2 xvfs_af_1 = vec2(1.0);
float xafafaf_1[2] = float[2](1.0, 2.0);
float xafaiai_1[2] = float[2](1.0, 2.0);
int xaipaiai_1[2] = int[2](1, 2);
uint xaupaiai[2] = uint[2](1u, 2u);
float xafpaiaf_1[2] = float[2](1.0, 2.0);
float xafpafai_1[2] = float[2](1.0, 2.0);
float xafpafaf_1[2] = float[2](1.0, 2.0);
ivec3 xavipai_1[1] = ivec3[1](ivec3(1));
vec3 xavfpai_1[1] = vec3[1](vec3(1.0));
vec3 xavfpaf_1[1] = vec3[1](vec3(1.0));
ivec2 xvisai_1 = ivec2(1);
uvec2 xvusai_1 = uvec2(1u);
vec2 xvfsai_1 = vec2(1.0);
vec2 xvfsaf_1 = vec2(1.0);
ivec2 ivispai = ivec2(1);
vec2 ivfspaf = vec2(1.0);
ivec2 ivis_ai = ivec2(1);
uvec2 ivus_ai = uvec2(1u);
vec2 ivfs_ai = vec2(1.0);
vec2 ivfs_af = vec2(1.0);
float iafafaf[2] = float[2](1.0, 2.0);
float iafaiai[2] = float[2](1.0, 2.0);
int iaipaiai_1[2] = int[2](1, 2);
float iafpafaf_1[2] = float[2](1.0, 2.0);
float iafpaiaf_1[2] = float[2](1.0, 2.0);
float iafpafai_1[2] = float[2](1.0, 2.0);
ivec3 iavipai[1] = ivec3[1](ivec3(1));
ivec3 iavfpai[1] = ivec3[1](ivec3(1));
vec3 iavfpaf[1] = vec3[1](vec3(1.0));
void globals() {
ivec2 phony = xvipaiai_1;
uvec2 phony_1 = xvupaiai_1;
vec2 phony_2 = xvfpaiai_1;
vec2 phony_3 = xvfpafaf_1;
vec2 phony_4 = xvfpaiaf_1;
uvec2 phony_5 = xvupuai_2;
uvec2 phony_6 = xvupaiu_2;
uvec2 phony_7 = xvuuai_2;
uvec2 phony_8 = xvuaiu_2;
ivec2 phony_9 = xvip_1;
uvec2 phony_10 = xvup_1;
vec2 phony_11 = xvfp_1;
mat2x2 phony_12 = xmfp_1;
mat2x2 phony_13 = xmfpaiaiaiai_1;
mat2x2 phony_14 = xmfpafaiaiai_1;
mat2x2 phony_15 = xmfpaiafaiai_1;
mat2x2 phony_16 = xmfpaiaiafai_1;
mat2x2 phony_17 = xmfpaiaiaiaf_1;
ivec2 phony_18 = xvispai_1;
vec2 phony_19 = xvfspaf_1;
ivec2 phony_20 = xvis_ai_1;
uvec2 phony_21 = xvus_ai_1;
vec2 phony_22 = xvfs_ai_1;
vec2 phony_23 = xvfs_af_1;
float phony_24[2] = xafafaf_1;
float phony_25[2] = xafaiai_1;
int phony_26[2] = xaipaiai_1;
uint phony_27[2] = xaupaiai;
float phony_28[2] = xafpaiaf_1;
float phony_29[2] = xafpafai_1;
float phony_30[2] = xafpafaf_1;
ivec3 phony_31[1] = xavipai_1;
vec3 phony_32[1] = xavfpai_1;
vec3 phony_33[1] = xavfpaf_1;
ivec2 phony_34 = xvisai_1;
uvec2 phony_35 = xvusai_1;
vec2 phony_36 = xvfsai_1;
vec2 phony_37 = xvfsaf_1;
ivec2 phony_38 = ivispai;
vec2 phony_39 = ivfspaf;
ivec2 phony_40 = ivis_ai;
uvec2 phony_41 = ivus_ai;
vec2 phony_42 = ivfs_ai;
vec2 phony_43 = ivfs_af;
float phony_44[2] = iafafaf;
float phony_45[2] = iafaiai;
int phony_46[2] = iaipaiai_1;
float phony_47[2] = iafpafaf_1;
float phony_48[2] = iafpaiaf_1;
float phony_49[2] = iafpafai_1;
ivec3 phony_50[1] = iavipai;
ivec3 phony_51[1] = iavfpai;
vec3 phony_52[1] = iavfpaf;
return;
}
void all_constant_arguments() {
ivec2 xvipaiai = ivec2(42, 43);
@ -239,6 +402,7 @@ void mixed_constant_and_runtime_arguments() {
}
void main() {
globals();
all_constant_arguments();
mixed_constant_and_runtime_arguments();
return;

7
naga/tests/out/glsl/wgsl-access.foo_vert.Vertex.glsl
View File

@ -30,6 +30,8 @@ struct Outer {
Inner om_nom_nom;
uint thing;
};
GlobalConst msl_padding_global_const = GlobalConst(0u, uvec3(0u, 0u, 0u), 0);
layout(std430) buffer Bar_block_0Vertex {
mat4x3 _matrix;
mat2x2 matrix_array[2];
@ -216,6 +218,7 @@ void main() {
int c2_[5] = int[5](0, 0, 0, 0, 0);
float baz_1 = foo;
foo = 1.0;
GlobalConst phony = msl_padding_global_const;
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
mat4x3 _matrix = _group_0_binding_0_vs._matrix;
@ -223,11 +226,11 @@ void main() {
float b = _group_0_binding_0_vs._matrix[3u][0];
int a_2 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
ivec2 c = _group_0_binding_2_vs;
float _e33 = read_from_private(foo);
float _e35 = read_from_private(foo);
c2_ = int[5](a_2, int(b), 3, 4, 5);
c2_[(vi + 1u)] = 42;
int value_1 = c2_[vi];
float _e47 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
float _e49 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
gl_Position = vec4((_matrix * vec4(ivec4(value_1))), 2.0);
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;

3
naga/tests/out/glsl/wgsl-f64.main.Compute.glsl
View File

@ -4,10 +4,13 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
const double k = 2.0LF;
double v = 1.0LF;
double f(double x) {
double z = 0.0;
double w = -1.0LF;
double phony = v;
double y = (30.0LF + 400.0LF);
z = (y + 5.0LF);
return (((x + y) + k) + 5.0LF);

3
naga/tests/out/glsl/wgsl-image.main.Compute.glsl
View File

@ -10,6 +10,8 @@ layout(rgba8ui) readonly uniform uimage2D _group_0_binding_1_cs;
uniform usampler2DArray _group_0_binding_5_cs;
layout(r32ui) readonly uniform uimage1D _group_0_binding_6_cs;
uniform usampler1D _group_0_binding_7_cs;
layout(r32ui) writeonly uniform uimage1D _group_0_binding_2_cs;
@ -27,6 +29,7 @@ void main() {
uvec4 value5_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, local_id.z), (int(local_id.z) + 1));
uvec4 value6_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1));
uvec4 value7_ = texelFetch(_group_0_binding_7_cs, int(local_id.x), int(local_id.z));
uvec4 value8_ = imageLoad(_group_0_binding_6_cs, int(local_id.x));
uvec4 value1u = texelFetch(_group_0_binding_0_cs, ivec2(uvec2(itc)), int(local_id.z));
uvec4 value2u = texelFetch(_group_0_binding_3_cs, ivec2(uvec2(itc)), int(local_id.z));
uvec4 value3u = texelFetch(_group_0_binding_3_cs, ivec2(uvec2(itc)), int(uint(local_id.z)));

7
naga/tests/out/hlsl/wgsl-access.hlsl
View File

@ -93,7 +93,7 @@ GlobalConst ConstructGlobalConst(uint arg0, uint3 arg1, int arg2) {
return ret;
}
static GlobalConst global_const = ConstructGlobalConst(0u, uint3(0u, 0u, 0u), int(0));
static GlobalConst msl_padding_global_const = ConstructGlobalConst(0u, uint3(0u, 0u, 0u), int(0));
RWByteAddressBuffer bar : register(u0);
cbuffer baz : register(b1) { Baz baz; }
RWByteAddressBuffer qux : register(u2);
@ -388,6 +388,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
float baz_1 = foo;
foo = 1.0;
GlobalConst phony = msl_padding_global_const;
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
@ -395,11 +396,11 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
float b = asfloat(bar.Load(0+3u*16+0));
int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
int2 c = asint(qux.Load2(0));
const float _e33 = read_from_private(foo);
const float _e35 = read_from_private(foo);
c2_ = Constructarray5_int_(a_2, naga_f2i32(b), int(3), int(4), int(5));
c2_[min(uint((vi + 1u)), 4u)] = int(42);
int value_1 = c2_[min(uint(vi), 4u)];
const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__());
const float _e49 = test_arr_as_arg(ZeroValuearray5_array10_float__());
return float4(mul(float4((value_1).xxxx), _matrix), 2.0);
}

184
naga/tests/out/hlsl/wgsl-f16.hlsl
View File

@ -137,215 +137,217 @@ ret_Constructarray2_half_ Constructarray2_half_(half arg0, half arg1) {
half f16_function(half x)
{
LayoutTest l = (LayoutTest)0;
half val = 15.203125h;
half _e4 = val;
val = (_e4 + -33344.0h);
half phony = private_variable;
half _e6 = val;
half _e9 = val;
val = (_e9 + (_e6 + 5.0h));
float _e13 = input_uniform.val_f32_;
half _e14 = val;
half _e18 = val;
val = (_e18 + half((_e13 + float(_e14))));
half _e22 = input_uniform.val_f16_;
half _e25 = val;
val = (_e25 + (_e22).xxx.z);
val = (_e6 + -33344.0h);
half _e8 = val;
half _e11 = val;
val = (_e11 + (_e8 + 5.0h));
float _e15 = input_uniform.val_f32_;
half _e16 = val;
half _e20 = val;
val = (_e20 + half((_e15 + float(_e16))));
half _e24 = input_uniform.val_f16_;
half _e27 = val;
val = (_e27 + (_e24).xxx.z);
output.Store(4, asuint(int(65504)));
output.Store(4, asuint(int(-65504)));
output.Store(0, asuint(65504u));
output.Store(0, asuint(0u));
output.Store(8, asuint(65504.0));
output.Store(8, asuint(-65504.0));
half _e49 = input_uniform.val_f16_;
half _e52 = input_storage.Load<half>(12);
output.Store(12, (_e49 + _e52));
half2 _e58 = input_uniform.val_f16_2_;
half2 _e61 = input_storage.Load<half2>(16);
output.Store(16, (_e58 + _e61));
half3 _e67 = input_uniform.val_f16_3_;
half3 _e70 = input_storage.Load<half3>(24);
output.Store(24, (_e67 + _e70));
half4 _e76 = input_uniform.val_f16_4_;
half4 _e79 = input_storage.Load<half4>(32);
output.Store(32, (_e76 + _e79));
half2x2 _e85 = GetMatval_mat2x2_OnUniformCompatible(input_uniform);
half2x2 _e88 = half2x2(input_storage.Load<half2>(44+0), input_storage.Load<half2>(44+4));
half _e51 = input_uniform.val_f16_;
half _e54 = input_storage.Load<half>(12);
output.Store(12, (_e51 + _e54));
half2 _e60 = input_uniform.val_f16_2_;
half2 _e63 = input_storage.Load<half2>(16);
output.Store(16, (_e60 + _e63));
half3 _e69 = input_uniform.val_f16_3_;
half3 _e72 = input_storage.Load<half3>(24);
output.Store(24, (_e69 + _e72));
half4 _e78 = input_uniform.val_f16_4_;
half4 _e81 = input_storage.Load<half4>(32);
output.Store(32, (_e78 + _e81));
half2x2 _e87 = GetMatval_mat2x2_OnUniformCompatible(input_uniform);
half2x2 _e90 = half2x2(input_storage.Load<half2>(44+0), input_storage.Load<half2>(44+4));
{
half2x2 _value2 = (_e85 + _e88);
half2x2 _value2 = (_e87 + _e90);
output.Store(44+0, _value2[0]);
output.Store(44+4, _value2[1]);
}
half2x3 _e94 = input_uniform.val_mat2x3_;
half2x3 _e97 = half2x3(input_storage.Load<half3>(56+0), input_storage.Load<half3>(56+8));
half2x3 _e96 = input_uniform.val_mat2x3_;
half2x3 _e99 = half2x3(input_storage.Load<half3>(56+0), input_storage.Load<half3>(56+8));
{
half2x3 _value2 = (_e94 + _e97);
half2x3 _value2 = (_e96 + _e99);
output.Store(56+0, _value2[0]);
output.Store(56+8, _value2[1]);
}
half2x4 _e103 = input_uniform.val_mat2x4_;
half2x4 _e106 = half2x4(input_storage.Load<half4>(72+0), input_storage.Load<half4>(72+8));
half2x4 _e105 = input_uniform.val_mat2x4_;
half2x4 _e108 = half2x4(input_storage.Load<half4>(72+0), input_storage.Load<half4>(72+8));
{
half2x4 _value2 = (_e103 + _e106);
half2x4 _value2 = (_e105 + _e108);
output.Store(72+0, _value2[0]);
output.Store(72+8, _value2[1]);
}
half3x2 _e112 = GetMatval_mat3x2_OnUniformCompatible(input_uniform);
half3x2 _e115 = half3x2(input_storage.Load<half2>(88+0), input_storage.Load<half2>(88+4), input_storage.Load<half2>(88+8));
half3x2 _e114 = GetMatval_mat3x2_OnUniformCompatible(input_uniform);
half3x2 _e117 = half3x2(input_storage.Load<half2>(88+0), input_storage.Load<half2>(88+4), input_storage.Load<half2>(88+8));
{
half3x2 _value2 = (_e112 + _e115);
half3x2 _value2 = (_e114 + _e117);
output.Store(88+0, _value2[0]);
output.Store(88+4, _value2[1]);
output.Store(88+8, _value2[2]);
}
half3x3 _e121 = input_uniform.val_mat3x3_;
half3x3 _e124 = half3x3(input_storage.Load<half3>(104+0), input_storage.Load<half3>(104+8), input_storage.Load<half3>(104+16));
half3x3 _e123 = input_uniform.val_mat3x3_;
half3x3 _e126 = half3x3(input_storage.Load<half3>(104+0), input_storage.Load<half3>(104+8), input_storage.Load<half3>(104+16));
{
half3x3 _value2 = (_e121 + _e124);
half3x3 _value2 = (_e123 + _e126);
output.Store(104+0, _value2[0]);
output.Store(104+8, _value2[1]);
output.Store(104+16, _value2[2]);
}
half3x4 _e130 = input_uniform.val_mat3x4_;
half3x4 _e133 = half3x4(input_storage.Load<half4>(128+0), input_storage.Load<half4>(128+8), input_storage.Load<half4>(128+16));
half3x4 _e132 = input_uniform.val_mat3x4_;
half3x4 _e135 = half3x4(input_storage.Load<half4>(128+0), input_storage.Load<half4>(128+8), input_storage.Load<half4>(128+16));
{
half3x4 _value2 = (_e130 + _e133);
half3x4 _value2 = (_e132 + _e135);
output.Store(128+0, _value2[0]);
output.Store(128+8, _value2[1]);
output.Store(128+16, _value2[2]);
}
half4x2 _e139 = GetMatval_mat4x2_OnUniformCompatible(input_uniform);
half4x2 _e142 = half4x2(input_storage.Load<half2>(152+0), input_storage.Load<half2>(152+4), input_storage.Load<half2>(152+8), input_storage.Load<half2>(152+12));
half4x2 _e141 = GetMatval_mat4x2_OnUniformCompatible(input_uniform);
half4x2 _e144 = half4x2(input_storage.Load<half2>(152+0), input_storage.Load<half2>(152+4), input_storage.Load<half2>(152+8), input_storage.Load<half2>(152+12));
{
half4x2 _value2 = (_e139 + _e142);
half4x2 _value2 = (_e141 + _e144);
output.Store(152+0, _value2[0]);
output.Store(152+4, _value2[1]);
output.Store(152+8, _value2[2]);
output.Store(152+12, _value2[3]);
}
half4x3 _e148 = input_uniform.val_mat4x3_;
half4x3 _e151 = half4x3(input_storage.Load<half3>(168+0), input_storage.Load<half3>(168+8), input_storage.Load<half3>(168+16), input_storage.Load<half3>(168+24));
half4x3 _e150 = input_uniform.val_mat4x3_;
half4x3 _e153 = half4x3(input_storage.Load<half3>(168+0), input_storage.Load<half3>(168+8), input_storage.Load<half3>(168+16), input_storage.Load<half3>(168+24));
{
half4x3 _value2 = (_e148 + _e151);
half4x3 _value2 = (_e150 + _e153);
output.Store(168+0, _value2[0]);
output.Store(168+8, _value2[1]);
output.Store(168+16, _value2[2]);
output.Store(168+24, _value2[3]);
}
half4x4 _e157 = input_uniform.val_mat4x4_;
half4x4 _e160 = half4x4(input_storage.Load<half4>(200+0), input_storage.Load<half4>(200+8), input_storage.Load<half4>(200+16), input_storage.Load<half4>(200+24));
half4x4 _e159 = input_uniform.val_mat4x4_;
half4x4 _e162 = half4x4(input_storage.Load<half4>(200+0), input_storage.Load<half4>(200+8), input_storage.Load<half4>(200+16), input_storage.Load<half4>(200+24));
{
half4x4 _value2 = (_e157 + _e160);
half4x4 _value2 = (_e159 + _e162);
output.Store(200+0, _value2[0]);
output.Store(200+8, _value2[1]);
output.Store(200+16, _value2[2]);
output.Store(200+24, _value2[3]);
}
half _e166[2] = Constructarray2_half_(input_arrays.Load<half>(0+0), input_arrays.Load<half>(0+2));
half _e168[2] = Constructarray2_half_(input_arrays.Load<half>(0+0), input_arrays.Load<half>(0+2));
{
half _value2[2] = _e166;
half _value2[2] = _e168;
output_arrays.Store(0+0, _value2[0]);
output_arrays.Store(0+2, _value2[1]);
}
half _e167 = val;
half _e169 = val;
val = (_e169 + abs(_e167));
half _e171 = val;
half _e172 = val;
val = (_e171 + abs(_e169));
half _e173 = val;
half _e174 = val;
half _e175 = val;
val = (_e175 + clamp(_e171, _e172, _e173));
half _e177 = val;
val = (_e177 + clamp(_e173, _e174, _e175));
half _e179 = val;
half _e182 = val;
val = (_e182 + dot((_e177).xx, (_e179).xx));
half _e181 = val;
half _e184 = val;
half _e185 = val;
val = (_e184 + dot((_e179).xx, (_e181).xx));
half _e186 = val;
half _e187 = val;
val = (_e187 + max(_e184, _e185));
half _e189 = val;
half _e190 = val;
val = (_e189 + max(_e186, _e187));
half _e191 = val;
half _e192 = val;
val = (_e192 + min(_e189, _e190));
half _e194 = val;
val = (_e194 + min(_e191, _e192));
half _e196 = val;
val = (_e196 + sign(_e194));
half _e199 = val;
val = (_e199 + 1.0h);
half2 _e203 = input_uniform.val_f16_2_;
float2 float_vec2_ = float2(_e203);
half _e198 = val;
val = (_e198 + sign(_e196));
half _e201 = val;
val = (_e201 + 1.0h);
half2 _e205 = input_uniform.val_f16_2_;
float2 float_vec2_ = float2(_e205);
output.Store(16, half2(float_vec2_));
half3 _e210 = input_uniform.val_f16_3_;
float3 float_vec3_ = float3(_e210);
half3 _e212 = input_uniform.val_f16_3_;
float3 float_vec3_ = float3(_e212);
output.Store(24, half3(float_vec3_));
half4 _e217 = input_uniform.val_f16_4_;
float4 float_vec4_ = float4(_e217);
half4 _e219 = input_uniform.val_f16_4_;
float4 float_vec4_ = float4(_e219);
output.Store(32, half4(float_vec4_));
half2x2 _e226 = GetMatval_mat2x2_OnUniformCompatible(input_uniform);
half2x2 _e228 = GetMatval_mat2x2_OnUniformCompatible(input_uniform);
{
half2x2 _value2 = half2x2(float2x2(_e226));
half2x2 _value2 = half2x2(float2x2(_e228));
output.Store(44+0, _value2[0]);
output.Store(44+4, _value2[1]);
}
half2x3 _e233 = input_uniform.val_mat2x3_;
half2x3 _e235 = input_uniform.val_mat2x3_;
{
half2x3 _value2 = half2x3(float2x3(_e233));
half2x3 _value2 = half2x3(float2x3(_e235));
output.Store(56+0, _value2[0]);
output.Store(56+8, _value2[1]);
}
half2x4 _e240 = input_uniform.val_mat2x4_;
half2x4 _e242 = input_uniform.val_mat2x4_;
{
half2x4 _value2 = half2x4(float2x4(_e240));
half2x4 _value2 = half2x4(float2x4(_e242));
output.Store(72+0, _value2[0]);
output.Store(72+8, _value2[1]);
}
half3x2 _e247 = GetMatval_mat3x2_OnUniformCompatible(input_uniform);
half3x2 _e249 = GetMatval_mat3x2_OnUniformCompatible(input_uniform);
{
half3x2 _value2 = half3x2(float3x2(_e247));
half3x2 _value2 = half3x2(float3x2(_e249));
output.Store(88+0, _value2[0]);
output.Store(88+4, _value2[1]);
output.Store(88+8, _value2[2]);
}
half3x3 _e254 = input_uniform.val_mat3x3_;
half3x3 _e256 = input_uniform.val_mat3x3_;
{
half3x3 _value2 = half3x3(float3x3(_e254));
half3x3 _value2 = half3x3(float3x3(_e256));
output.Store(104+0, _value2[0]);
output.Store(104+8, _value2[1]);
output.Store(104+16, _value2[2]);
}
half3x4 _e261 = input_uniform.val_mat3x4_;
half3x4 _e263 = input_uniform.val_mat3x4_;
{
half3x4 _value2 = half3x4(float3x4(_e261));
half3x4 _value2 = half3x4(float3x4(_e263));
output.Store(128+0, _value2[0]);
output.Store(128+8, _value2[1]);
output.Store(128+16, _value2[2]);
}
half4x2 _e268 = GetMatval_mat4x2_OnUniformCompatible(input_uniform);
half4x2 _e270 = GetMatval_mat4x2_OnUniformCompatible(input_uniform);
{
half4x2 _value2 = half4x2(float4x2(_e268));
half4x2 _value2 = half4x2(float4x2(_e270));
output.Store(152+0, _value2[0]);
output.Store(152+4, _value2[1]);
output.Store(152+8, _value2[2]);
output.Store(152+12, _value2[3]);
}
half4x3 _e275 = input_uniform.val_mat4x3_;
half4x3 _e277 = input_uniform.val_mat4x3_;
{
half4x3 _value2 = half4x3(float4x3(_e275));
half4x3 _value2 = half4x3(float4x3(_e277));
output.Store(168+0, _value2[0]);
output.Store(168+8, _value2[1]);
output.Store(168+16, _value2[2]);
output.Store(168+24, _value2[3]);
}
half4x4 _e282 = input_uniform.val_mat4x4_;
half4x4 _e284 = input_uniform.val_mat4x4_;
{
half4x4 _value2 = half4x4(float4x4(_e282));
half4x4 _value2 = half4x4(float4x4(_e284));
output.Store(200+0, _value2[0]);
output.Store(200+8, _value2[1]);
output.Store(200+16, _value2[2]);
output.Store(200+24, _value2[3]);
}
half _e285 = val;
return _e285;
half _e287 = val;
return _e287;
}
[numthreads(1, 1, 1)]

1
naga/tests/out/hlsl/wgsl-f64.hlsl
View File

@ -7,6 +7,7 @@ double f(double x)
double z = (double)0;
double w = -1.0L;
double phony = v;
double y = (30.0L + 400.0L);
z = (y + 5.0L);
return (((x + y) + k) + 5.0L);

2
naga/tests/out/hlsl/wgsl-image.hlsl
View File

@ -29,6 +29,7 @@ int2 naga_mod(int2 lhs, int2 rhs) {
return lhs - (lhs / divisor) * divisor;
}
uint4 LoadedStorageValueFromuint(uint arg) {uint4 ret = uint4(arg, 0u, 0u, 1u);return ret;}
uint2 NagaRWDimensions2D(RWTexture2D<uint4> tex)
{
uint4 ret;
@ -49,6 +50,7 @@ void main(uint3 local_id : SV_GroupThreadID)
uint4 value5_ = image_array_src.Load(int4(itc, local_id.z, asint(asuint(int(local_id.z)) + asuint(int(1)))));
uint4 value6_ = image_array_src.Load(int4(itc, int(local_id.z), asint(asuint(int(local_id.z)) + asuint(int(1)))));
uint4 value7_ = image_1d_src.Load(int2(int(local_id.x), int(local_id.z)));
uint4 value8_ = LoadedStorageValueFromuint(image_dup_src.Load(int(local_id.x)));
uint4 value1u = image_mipmapped_src.Load(int3(uint2(itc), int(local_id.z)));
uint4 value2u = image_multisampled_src.Load(uint2(itc), int(local_id.z));
uint4 value3u = image_multisampled_src.Load(uint2(itc), uint(local_id.z));

119
naga/tests/out/hlsl/wgsl-int64.hlsl
View File

@ -67,83 +67,84 @@ int64_t int64_function(int64_t x)
{
int64_t val = 20L;
int64_t _e8 = val;
val = (_e8 + ((31L - 1002003004005006L) + -9223372036854775807L));
int64_t phony = private_variable;
int64_t _e10 = val;
int64_t _e13 = val;
val = (_e13 + (_e10 + 5L));
uint _e17 = input_uniform.val_u32_;
int64_t _e18 = val;
int64_t _e22 = val;
val = (_e22 + int64_t((_e17 + uint(_e18))));
int _e26 = input_uniform.val_i32_;
int64_t _e27 = val;
int64_t _e31 = val;
val = (_e31 + int64_t(asint(asuint(_e26) + asuint(int(_e27)))));
float _e35 = input_uniform.val_f32_;
int64_t _e36 = val;
int64_t _e40 = val;
val = (_e40 + naga_f2i64((_e35 + float(_e36))));
int64_t _e44 = input_uniform.val_i64_;
int64_t _e47 = val;
val = (_e47 + (_e44).xxx.z);
uint64_t _e51 = input_uniform.val_u64_;
int64_t _e53 = val;
val = (_e53 + _e51);
uint64_t2 _e57 = input_uniform.val_u64_2_;
int64_t _e60 = val;
val = (_e60 + _e57.y);
uint64_t3 _e64 = input_uniform.val_u64_3_;
int64_t _e67 = val;
val = (_e67 + _e64.z);
uint64_t4 _e71 = input_uniform.val_u64_4_;
int64_t _e74 = val;
val = (_e74 + _e71.w);
int64_t _e77 = val;
val = (_e77 + (-9223372036854775807L - 1L));
int64_t _e83 = input_uniform.val_i64_;
int64_t _e86 = input_storage.Load<int64_t>(128);
output.Store(128, (_e83 + _e86));
int64_t2 _e92 = input_uniform.val_i64_2_;
int64_t2 _e95 = input_storage.Load<int64_t2>(144);
output.Store(144, (_e92 + _e95));
int64_t3 _e101 = input_uniform.val_i64_3_;
int64_t3 _e104 = input_storage.Load<int64_t3>(160);
output.Store(160, (_e101 + _e104));
int64_t4 _e110 = input_uniform.val_i64_4_;
int64_t4 _e113 = input_storage.Load<int64_t4>(192);
output.Store(192, (_e110 + _e113));
int64_t _e119[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
val = (_e10 + ((31L - 1002003004005006L) + -9223372036854775807L));
int64_t _e12 = val;
int64_t _e15 = val;
val = (_e15 + (_e12 + 5L));
uint _e19 = input_uniform.val_u32_;
int64_t _e20 = val;
int64_t _e24 = val;
val = (_e24 + int64_t((_e19 + uint(_e20))));
int _e28 = input_uniform.val_i32_;
int64_t _e29 = val;
int64_t _e33 = val;
val = (_e33 + int64_t(asint(asuint(_e28) + asuint(int(_e29)))));
float _e37 = input_uniform.val_f32_;
int64_t _e38 = val;
int64_t _e42 = val;
val = (_e42 + naga_f2i64((_e37 + float(_e38))));
int64_t _e46 = input_uniform.val_i64_;
int64_t _e49 = val;
val = (_e49 + (_e46).xxx.z);
uint64_t _e53 = input_uniform.val_u64_;
int64_t _e55 = val;
val = (_e55 + _e53);
uint64_t2 _e59 = input_uniform.val_u64_2_;
int64_t _e62 = val;
val = (_e62 + _e59.y);
uint64_t3 _e66 = input_uniform.val_u64_3_;
int64_t _e69 = val;
val = (_e69 + _e66.z);
uint64_t4 _e73 = input_uniform.val_u64_4_;
int64_t _e76 = val;
val = (_e76 + _e73.w);
int64_t _e79 = val;
val = (_e79 + (-9223372036854775807L - 1L));
int64_t _e85 = input_uniform.val_i64_;
int64_t _e88 = input_storage.Load<int64_t>(128);
output.Store(128, (_e85 + _e88));
int64_t2 _e94 = input_uniform.val_i64_2_;
int64_t2 _e97 = input_storage.Load<int64_t2>(144);
output.Store(144, (_e94 + _e97));
int64_t3 _e103 = input_uniform.val_i64_3_;
int64_t3 _e106 = input_storage.Load<int64_t3>(160);
output.Store(160, (_e103 + _e106));
int64_t4 _e112 = input_uniform.val_i64_4_;
int64_t4 _e115 = input_storage.Load<int64_t4>(192);
output.Store(192, (_e112 + _e115));
int64_t _e121[2] = Constructarray2_int64_t_(input_arrays.Load<int64_t>(16+0), input_arrays.Load<int64_t>(16+8));
{
int64_t _value2[2] = _e119;
int64_t _value2[2] = _e121;
output_arrays.Store(16+0, _value2[0]);
output_arrays.Store(16+8, _value2[1]);
}
int64_t _e120 = val;
int64_t _e122 = val;
val = (_e122 + abs(_e120));
int64_t _e124 = val;
int64_t _e125 = val;
val = (_e124 + abs(_e122));
int64_t _e126 = val;
int64_t _e127 = val;
int64_t _e128 = val;
val = (_e128 + clamp(_e124, _e125, _e126));
int64_t _e130 = val;
val = (_e130 + clamp(_e126, _e127, _e128));
int64_t _e132 = val;
int64_t _e135 = val;
val = (_e135 + dot((_e130).xx, (_e132).xx));
int64_t _e134 = val;
int64_t _e137 = val;
int64_t _e138 = val;
val = (_e137 + dot((_e132).xx, (_e134).xx));
int64_t _e139 = val;
int64_t _e140 = val;
val = (_e140 + max(_e137, _e138));
int64_t _e142 = val;
int64_t _e143 = val;
val = (_e142 + max(_e139, _e140));
int64_t _e144 = val;
int64_t _e145 = val;
val = (_e145 + min(_e142, _e143));
int64_t _e147 = val;
val = (_e147 + min(_e144, _e145));
int64_t _e149 = val;
val = (_e149 + sign(_e147));
int64_t _e151 = val;
return _e151;
val = (_e151 + sign(_e149));
int64_t _e153 = val;
return _e153;
}
uint64_t naga_f2u64(float value) {

189
naga/tests/out/ir/wgsl-access.compact.ron
View File

@ -427,7 +427,7 @@
overrides: [],
global_variables: [
(
name: Some("global_const"),
name: Some("msl_padding_global_const"),
space: Private,
binding: None,
ty: 3,
@ -2383,87 +2383,91 @@
pointer: 2,
),
Literal(F32(1.0)),
GlobalVariable(0),
Load(
pointer: 5,
),
GlobalVariable(1),
AccessIndex(
base: 5,
base: 7,
index: 0,
),
Load(
pointer: 6,
pointer: 8,
),
GlobalVariable(1),
AccessIndex(
base: 8,
base: 10,
index: 4,
),
Load(
pointer: 9,
pointer: 11,
),
Literal(U32(3)),
GlobalVariable(1),
AccessIndex(
base: 12,
index: 0,
),
Access(
base: 13,
index: 11,
),
AccessIndex(
base: 14,
index: 0,
),
Load(
pointer: 15,
Access(
base: 15,
index: 13,
),
GlobalVariable(1),
AccessIndex(
base: 17,
index: 5,
base: 16,
index: 0,
),
Load(
pointer: 17,
),
GlobalVariable(1),
AccessIndex(
base: 19,
index: 5,
),
ArrayLength(20),
GlobalVariable(1),
AccessIndex(
base: 21,
index: 5,
),
ArrayLength(22),
Literal(U32(2)),
Binary(
op: Subtract,
left: 21,
right: 22,
left: 23,
right: 24,
),
Access(
base: 18,
index: 23,
base: 20,
index: 25,
),
AccessIndex(
base: 24,
base: 26,
index: 0,
),
Load(
pointer: 25,
),
GlobalVariable(3),
Load(
pointer: 27,
),
GlobalVariable(3),
Load(
pointer: 29,
),
GlobalVariable(1),
AccessIndex(
base: 29,
base: 31,
index: 5,
),
AccessIndex(
base: 30,
base: 32,
index: 0,
),
AccessIndex(
base: 31,
base: 33,
index: 0,
),
CallResult(2),
As(
expr: 16,
expr: 18,
kind: Sint,
convert: Some(4),
),
@ -2473,11 +2477,11 @@
Compose(
ty: 25,
components: [
26,
34,
35,
28,
36,
37,
38,
39,
],
),
LocalVariable(1),
@ -2485,57 +2489,58 @@
Binary(
op: Add,
left: 0,
right: 40,
right: 42,
),
Access(
base: 39,
index: 41,
base: 41,
index: 43,
),
Literal(I32(42)),
Access(
base: 39,
base: 41,
index: 0,
),
Load(
pointer: 44,
pointer: 46,
),
ZeroValue(23),
CallResult(3),
Splat(
size: Quad,
value: 45,
value: 47,
),
As(
expr: 48,
expr: 50,
kind: Float,
convert: Some(4),
),
Binary(
op: Multiply,
left: 7,
right: 49,
left: 9,
right: 51,
),
Literal(F32(2.0)),
Compose(
ty: 24,
components: [
50,
51,
52,
53,
],
),
],
named_expressions: {
0: "vi",
3: "baz",
7: "_matrix",
10: "arr",
11: "index",
16: "b",
26: "a",
28: "c",
32: "data_pointer",
33: "foo_value",
45: "value",
6: "phony",
9: "_matrix",
12: "arr",
13: "index",
18: "b",
28: "a",
30: "c",
34: "data_pointer",
35: "foo_value",
47: "value",
},
body: [
Emit((
@ -2546,6 +2551,10 @@
pointer: 2,
value: 4,
),
Emit((
start: 6,
end: 7,
)),
Call(
function: 0,
arguments: [],
@ -2557,31 +2566,27 @@
result: None,
),
Emit((
start: 6,
end: 8,
start: 8,
end: 10,
)),
Emit((
start: 9,
end: 11,
start: 11,
end: 13,
)),
Emit((
start: 13,
end: 17,
)),
Emit((
start: 18,
start: 15,
end: 19,
)),
Emit((
start: 20,
end: 22,
end: 21,
)),
Emit((
start: 23,
end: 27,
start: 22,
end: 24,
)),
Emit((
start: 28,
start: 25,
end: 29,
)),
Emit((
@ -2589,57 +2594,61 @@
end: 31,
)),
Emit((
start: 31,
start: 32,
end: 33,
)),
Emit((
start: 33,
end: 35,
)),
Call(
function: 2,
arguments: [
2,
],
result: Some(33),
result: Some(35),
),
Emit((
start: 34,
end: 35,
start: 36,
end: 37,
)),
Emit((
start: 38,
end: 39,
start: 40,
end: 41,
)),
Store(
pointer: 39,
value: 38,
pointer: 41,
value: 40,
),
Emit((
start: 41,
end: 43,
start: 43,
end: 45,
)),
Store(
pointer: 42,
value: 43,
pointer: 44,
value: 45,
),
Emit((
start: 44,
end: 46,
start: 46,
end: 48,
)),
Call(
function: 3,
arguments: [
46,
48,
],
result: Some(47),
result: Some(49),
),
Emit((
start: 48,
end: 51,
)),
Emit((
start: 52,
start: 50,
end: 53,
)),
Emit((
start: 54,
end: 55,
)),
Return(
value: Some(52),
value: Some(54),
),
],
diagnostic_filter_leaf: None,

189
naga/tests/out/ir/wgsl-access.ron
View File

@ -427,7 +427,7 @@
overrides: [],
global_variables: [
(
name: Some("global_const"),
name: Some("msl_padding_global_const"),
space: Private,
binding: None,
ty: 3,
@ -2383,87 +2383,91 @@
pointer: 2,
),
Literal(F32(1.0)),
GlobalVariable(0),
Load(
pointer: 5,
),
GlobalVariable(1),
AccessIndex(
base: 5,
base: 7,
index: 0,
),
Load(
pointer: 6,
pointer: 8,
),
GlobalVariable(1),
AccessIndex(
base: 8,
base: 10,
index: 4,
),
Load(
pointer: 9,
pointer: 11,
),
Literal(U32(3)),
GlobalVariable(1),
AccessIndex(
base: 12,
index: 0,
),
Access(
base: 13,
index: 11,
),
AccessIndex(
base: 14,
index: 0,
),
Load(
pointer: 15,
Access(
base: 15,
index: 13,
),
GlobalVariable(1),
AccessIndex(
base: 17,
index: 5,
base: 16,
index: 0,
),
Load(
pointer: 17,
),
GlobalVariable(1),
AccessIndex(
base: 19,
index: 5,
),
ArrayLength(20),
GlobalVariable(1),
AccessIndex(
base: 21,
index: 5,
),
ArrayLength(22),
Literal(U32(2)),
Binary(
op: Subtract,
left: 21,
right: 22,
left: 23,
right: 24,
),
Access(
base: 18,
index: 23,
base: 20,
index: 25,
),
AccessIndex(
base: 24,
base: 26,
index: 0,
),
Load(
pointer: 25,
),
GlobalVariable(3),
Load(
pointer: 27,
),
GlobalVariable(3),
Load(
pointer: 29,
),
GlobalVariable(1),
AccessIndex(
base: 29,
base: 31,
index: 5,
),
AccessIndex(
base: 30,
base: 32,
index: 0,
),
AccessIndex(
base: 31,
base: 33,
index: 0,
),
CallResult(2),
As(
expr: 16,
expr: 18,
kind: Sint,
convert: Some(4),
),
@ -2473,11 +2477,11 @@
Compose(
ty: 25,
components: [
26,
34,
35,
28,
36,
37,
38,
39,
],
),
LocalVariable(1),
@ -2485,57 +2489,58 @@
Binary(
op: Add,
left: 0,
right: 40,
right: 42,
),
Access(
base: 39,
index: 41,
base: 41,
index: 43,
),
Literal(I32(42)),
Access(
base: 39,
base: 41,
index: 0,
),
Load(
pointer: 44,
pointer: 46,
),
ZeroValue(23),
CallResult(3),
Splat(
size: Quad,
value: 45,
value: 47,
),
As(
expr: 48,
expr: 50,
kind: Float,
convert: Some(4),
),
Binary(
op: Multiply,
left: 7,
right: 49,
left: 9,
right: 51,
),
Literal(F32(2.0)),
Compose(
ty: 24,
components: [
50,
51,
52,
53,
],
),
],
named_expressions: {
0: "vi",
3: "baz",
7: "_matrix",
10: "arr",
11: "index",
16: "b",
26: "a",
28: "c",
32: "data_pointer",
33: "foo_value",
45: "value",
6: "phony",
9: "_matrix",
12: "arr",
13: "index",
18: "b",
28: "a",
30: "c",
34: "data_pointer",
35: "foo_value",
47: "value",
},
body: [
Emit((
@ -2546,6 +2551,10 @@
pointer: 2,
value: 4,
),
Emit((
start: 6,
end: 7,
)),
Call(
function: 0,
arguments: [],
@ -2557,31 +2566,27 @@
result: None,
),
Emit((
start: 6,
end: 8,
start: 8,
end: 10,
)),
Emit((
start: 9,
end: 11,
start: 11,
end: 13,
)),
Emit((
start: 13,
end: 17,
)),
Emit((
start: 18,
start: 15,
end: 19,
)),
Emit((
start: 20,
end: 22,
end: 21,
)),
Emit((
start: 23,
end: 27,
start: 22,
end: 24,
)),
Emit((
start: 28,
start: 25,
end: 29,
)),
Emit((
@ -2589,57 +2594,61 @@
end: 31,
)),
Emit((
start: 31,
start: 32,
end: 33,
)),
Emit((
start: 33,
end: 35,
)),
Call(
function: 2,
arguments: [
2,
],
result: Some(33),
result: Some(35),
),
Emit((
start: 34,
end: 35,
start: 36,
end: 37,
)),
Emit((
start: 38,
end: 39,
start: 40,
end: 41,
)),
Store(
pointer: 39,
value: 38,
pointer: 41,
value: 40,
),
Emit((
start: 41,
end: 43,
start: 43,
end: 45,
)),
Store(
pointer: 42,
value: 43,
pointer: 44,
value: 45,
),
Emit((
start: 44,
end: 46,
start: 46,
end: 48,
)),
Call(
function: 3,
arguments: [
46,
48,
],
result: Some(47),
result: Some(49),
),
Emit((
start: 48,
end: 51,
)),
Emit((
start: 52,
start: 50,
end: 53,
)),
Emit((
start: 54,
end: 55,
)),
Return(
value: Some(52),
value: Some(54),
),
],
diagnostic_filter_leaf: None,

8
naga/tests/out/ir/wgsl-overrides.compact.ron
View File

@ -172,9 +172,17 @@
LocalVariable(2),
GlobalVariable(1),
Override(2),
Override(1),
Override(3),
Override(6),
Override(7),
],
named_expressions: {
4: "a",
13: "phony",
14: "phony",
15: "phony",
16: "phony",
},
body: [
Emit((

8
naga/tests/out/ir/wgsl-overrides.ron
View File

@ -172,9 +172,17 @@
LocalVariable(2),
GlobalVariable(1),
Override(2),
Override(1),
Override(3),
Override(6),
Override(7),
],
named_expressions: {
4: "a",
13: "phony",
14: "phony",
15: "phony",
16: "phony",
},
body: [
Emit((

165
naga/tests/out/msl/wgsl-abstract-types-var.msl
View File

@ -20,6 +20,117 @@ struct type_13 {
metal::float3 inner[1];
};
void globals(
thread metal::int2& xvipaiai_1,
thread metal::uint2& xvupaiai_1,
thread metal::float2& xvfpaiai_1,
thread metal::float2& xvfpafaf_1,
thread metal::float2& xvfpaiaf_1,
thread metal::uint2& xvupuai_2,
thread metal::uint2& xvupaiu_2,
thread metal::uint2& xvuuai_2,
thread metal::uint2& xvuaiu_2,
thread metal::int2& xvip_1,
thread metal::uint2& xvup_1,
thread metal::float2& xvfp_1,
thread metal::float2x2& xmfp_1,
thread metal::float2x2& xmfpaiaiaiai_1,
thread metal::float2x2& xmfpafaiaiai_1,
thread metal::float2x2& xmfpaiafaiai_1,
thread metal::float2x2& xmfpaiaiafai_1,
thread metal::float2x2& xmfpaiaiaiaf_1,
thread metal::int2& xvispai_1,
thread metal::float2& xvfspaf_1,
thread metal::int2& xvis_ai_1,
thread metal::uint2& xvus_ai_1,
thread metal::float2& xvfs_ai_1,
thread metal::float2& xvfs_af_1,
thread type_7& xafafaf_1,
thread type_7& xafaiai_1,
thread type_8& xaipaiai_1,
thread type_9& xaupaiai,
thread type_7& xafpaiaf_1,
thread type_7& xafpafai_1,
thread type_7& xafpafaf_1,
thread type_11& xavipai_1,
thread type_13& xavfpai_1,
thread type_13& xavfpaf_1,
thread metal::int2& xvisai_1,
thread metal::uint2& xvusai_1,
thread metal::float2& xvfsai_1,
thread metal::float2& xvfsaf_1,
thread metal::int2& ivispai,
thread metal::float2& ivfspaf,
thread metal::int2& ivis_ai,
thread metal::uint2& ivus_ai,
thread metal::float2& ivfs_ai,
thread metal::float2& ivfs_af,
thread type_7& iafafaf,
thread type_7& iafaiai,
thread type_8& iaipaiai_1,
thread type_7& iafpafaf_1,
thread type_7& iafpaiaf_1,
thread type_7& iafpafai_1,
thread type_11& iavipai,
thread type_11& iavfpai,
thread type_13& iavfpaf
) {
metal::int2 phony = xvipaiai_1;
metal::uint2 phony_1 = xvupaiai_1;
metal::float2 phony_2 = xvfpaiai_1;
metal::float2 phony_3 = xvfpafaf_1;
metal::float2 phony_4 = xvfpaiaf_1;
metal::uint2 phony_5 = xvupuai_2;
metal::uint2 phony_6 = xvupaiu_2;
metal::uint2 phony_7 = xvuuai_2;
metal::uint2 phony_8 = xvuaiu_2;
metal::int2 phony_9 = xvip_1;
metal::uint2 phony_10 = xvup_1;
metal::float2 phony_11 = xvfp_1;
metal::float2x2 phony_12 = xmfp_1;
metal::float2x2 phony_13 = xmfpaiaiaiai_1;
metal::float2x2 phony_14 = xmfpafaiaiai_1;
metal::float2x2 phony_15 = xmfpaiafaiai_1;
metal::float2x2 phony_16 = xmfpaiaiafai_1;
metal::float2x2 phony_17 = xmfpaiaiaiaf_1;
metal::int2 phony_18 = xvispai_1;
metal::float2 phony_19 = xvfspaf_1;
metal::int2 phony_20 = xvis_ai_1;
metal::uint2 phony_21 = xvus_ai_1;
metal::float2 phony_22 = xvfs_ai_1;
metal::float2 phony_23 = xvfs_af_1;
type_7 phony_24 = xafafaf_1;
type_7 phony_25 = xafaiai_1;
type_8 phony_26 = xaipaiai_1;
type_9 phony_27 = xaupaiai;
type_7 phony_28 = xafpaiaf_1;
type_7 phony_29 = xafpafai_1;
type_7 phony_30 = xafpafaf_1;
type_11 phony_31 = xavipai_1;
type_13 phony_32 = xavfpai_1;
type_13 phony_33 = xavfpaf_1;
metal::int2 phony_34 = xvisai_1;
metal::uint2 phony_35 = xvusai_1;
metal::float2 phony_36 = xvfsai_1;
metal::float2 phony_37 = xvfsaf_1;
metal::int2 phony_38 = ivispai;
metal::float2 phony_39 = ivfspaf;
metal::int2 phony_40 = ivis_ai;
metal::uint2 phony_41 = ivus_ai;
metal::float2 phony_42 = ivfs_ai;
metal::float2 phony_43 = ivfs_af;
type_7 phony_44 = iafafaf;
type_7 phony_45 = iafaiai;
type_8 phony_46 = iaipaiai_1;
type_7 phony_47 = iafpafaf_1;
type_7 phony_48 = iafpaiaf_1;
type_7 phony_49 = iafpafai_1;
type_11 phony_50 = iavipai;
type_11 phony_51 = iavfpai;
type_13 phony_52 = iavfpaf;
return;
}
void all_constant_arguments(
) {
metal::int2 xvipaiai = metal::int2(42, 43);
@ -256,6 +367,60 @@ void mixed_constant_and_runtime_arguments(
kernel void main_(
) {
metal::int2 xvipaiai_1 = metal::int2(42, 43);
metal::uint2 xvupaiai_1 = metal::uint2(44u, 45u);
metal::float2 xvfpaiai_1 = metal::float2(46.0, 47.0);
metal::float2 xvfpafaf_1 = metal::float2(48.0, 49.0);
metal::float2 xvfpaiaf_1 = metal::float2(48.0, 49.0);
metal::uint2 xvupuai_2 = metal::uint2(42u, 43u);
metal::uint2 xvupaiu_2 = metal::uint2(42u, 43u);
metal::uint2 xvuuai_2 = metal::uint2(42u, 43u);
metal::uint2 xvuaiu_2 = metal::uint2(42u, 43u);
metal::int2 xvip_1 = metal::int2(0, 0);
metal::uint2 xvup_1 = metal::uint2(0u, 0u);
metal::float2 xvfp_1 = metal::float2(0.0, 0.0);
metal::float2x2 xmfp_1 = metal::float2x2(metal::float2(0.0, 0.0), metal::float2(0.0, 0.0));
metal::float2x2 xmfpaiaiaiai_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
metal::float2x2 xmfpafaiaiai_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
metal::float2x2 xmfpaiafaiai_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
metal::float2x2 xmfpaiaiafai_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
metal::float2x2 xmfpaiaiaiaf_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
metal::int2 xvispai_1 = metal::int2(1);
metal::float2 xvfspaf_1 = metal::float2(1.0);
metal::int2 xvis_ai_1 = metal::int2(1);
metal::uint2 xvus_ai_1 = metal::uint2(1u);
metal::float2 xvfs_ai_1 = metal::float2(1.0);
metal::float2 xvfs_af_1 = metal::float2(1.0);
type_7 xafafaf_1 = type_7 {1.0, 2.0};
type_7 xafaiai_1 = type_7 {1.0, 2.0};
type_8 xaipaiai_1 = type_8 {1, 2};
type_9 xaupaiai = type_9 {1u, 2u};
type_7 xafpaiaf_1 = type_7 {1.0, 2.0};
type_7 xafpafai_1 = type_7 {1.0, 2.0};
type_7 xafpafaf_1 = type_7 {1.0, 2.0};
type_11 xavipai_1 = type_11 {metal::int3(1)};
type_13 xavfpai_1 = type_13 {metal::float3(1.0)};
type_13 xavfpaf_1 = type_13 {metal::float3(1.0)};
metal::int2 xvisai_1 = metal::int2(1);
metal::uint2 xvusai_1 = metal::uint2(1u);
metal::float2 xvfsai_1 = metal::float2(1.0);
metal::float2 xvfsaf_1 = metal::float2(1.0);
metal::int2 ivispai = metal::int2(1);
metal::float2 ivfspaf = metal::float2(1.0);
metal::int2 ivis_ai = metal::int2(1);
metal::uint2 ivus_ai = metal::uint2(1u);
metal::float2 ivfs_ai = metal::float2(1.0);
metal::float2 ivfs_af = metal::float2(1.0);
type_7 iafafaf = type_7 {1.0, 2.0};
type_7 iafaiai = type_7 {1.0, 2.0};
type_8 iaipaiai_1 = type_8 {1, 2};
type_7 iafpafaf_1 = type_7 {1.0, 2.0};
type_7 iafpaiaf_1 = type_7 {1.0, 2.0};
type_7 iafpafai_1 = type_7 {1.0, 2.0};
type_11 iavipai = type_11 {metal::int3(1)};
type_11 iavfpai = type_11 {metal::int3(1)};
type_13 iavfpaf = type_13 {metal::float3(1.0)};
globals(xvipaiai_1, xvupaiai_1, xvfpaiai_1, xvfpafaf_1, xvfpaiaf_1, xvupuai_2, xvupaiu_2, xvuuai_2, xvuaiu_2, xvip_1, xvup_1, xvfp_1, xmfp_1, xmfpaiaiaiai_1, xmfpafaiaiai_1, xmfpaiafaiai_1, xmfpaiaiafai_1, xmfpaiaiaiaf_1, xvispai_1, xvfspaf_1, xvis_ai_1, xvus_ai_1, xvfs_ai_1, xvfs_af_1, xafafaf_1, xafaiai_1, xaipaiai_1, xaupaiai, xafpaiaf_1, xafpafai_1, xafpafaf_1, xavipai_1, xavfpai_1, xavfpaf_1, xvisai_1, xvusai_1, xvfsai_1, xvfsaf_1, ivispai, ivfspaf, ivis_ai, ivus_ai, ivfs_ai, ivfs_af, iafafaf, iafaiai, iaipaiai_1, iafpafaf_1, iafpaiaf_1, iafpafai_1, iavipai, iavfpai, iavfpaf);
all_constant_arguments();
mixed_constant_and_runtime_arguments();
return;

6
naga/tests/out/msl/wgsl-access.msl
View File

@ -285,10 +285,12 @@ vertex foo_vertOutput foo_vert(
, constant MatCx2InArray& nested_mat_cx2_ [[buffer(3)]]
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
GlobalConst msl_padding_global_const = GlobalConst {0u, {}, metal::uint3(0u, 0u, 0u), 0};
float foo = 0.0;
type_20 c2_ = {};
float baz_1 = foo;
foo = 1.0;
GlobalConst phony = msl_padding_global_const;
test_matrix_within_struct_accesses(baz);
test_matrix_within_array_within_struct_accesses(nested_mat_cx2_);
metal::float4x3 _matrix = bar._matrix;
@ -296,11 +298,11 @@ vertex foo_vertOutput foo_vert(
float b = bar._matrix[3u].x;
int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
metal::int2 c = qux;
float _e33 = read_from_private(foo);
float _e35 = read_from_private(foo);
c2_ = type_20 {a_2, naga_f2i32(b), 3, 4, 5};
c2_.inner[vi + 1u] = 42;
int value_1 = c2_.inner[vi];
float _e47 = test_arr_as_arg(type_18 {});
float _e49 = test_arr_as_arg(type_18 {});
return foo_vertOutput { metal::float4(_matrix * static_cast<metal::float4>(metal::int4(value_1)), 2.0) };
}

188
naga/tests/out/msl/wgsl-f16.msl
View File

@ -48,126 +48,129 @@ constant half constant_variable = 15.203125h;
half f16_function(
half x,
thread half& private_variable,
constant UniformCompatible& input_uniform,
device UniformCompatible const& input_storage,
device StorageCompatible const& input_arrays,
device UniformCompatible& output,
device StorageCompatible& output_arrays
) {
LayoutTest l = {};
half val = 15.203125h;
half _e4 = val;
val = _e4 + -33344.0h;
half phony = private_variable;
half _e6 = val;
half _e9 = val;
val = _e9 + (_e6 + 5.0h);
float _e13 = input_uniform.val_f32_;
half _e14 = val;
half _e18 = val;
val = _e18 + static_cast<half>(_e13 + static_cast<float>(_e14));
half _e22 = input_uniform.val_f16_;
half _e25 = val;
val = _e25 + metal::half3(_e22).z;
val = _e6 + -33344.0h;
half _e8 = val;
half _e11 = val;
val = _e11 + (_e8 + 5.0h);
float _e15 = input_uniform.val_f32_;
half _e16 = val;
half _e20 = val;
val = _e20 + static_cast<half>(_e15 + static_cast<float>(_e16));
half _e24 = input_uniform.val_f16_;
half _e27 = val;
val = _e27 + metal::half3(_e24).z;
output.val_i32_ = 65504;
output.val_i32_ = -65504;
output.val_u32_ = 65504u;
output.val_u32_ = 0u;
output.val_f32_ = 65504.0;
output.val_f32_ = -65504.0;
half _e49 = input_uniform.val_f16_;
half _e52 = input_storage.val_f16_;
output.val_f16_ = _e49 + _e52;
metal::half2 _e58 = input_uniform.val_f16_2_;
metal::half2 _e61 = input_storage.val_f16_2_;
output.val_f16_2_ = _e58 + _e61;
metal::half3 _e67 = input_uniform.val_f16_3_;
metal::half3 _e70 = input_storage.val_f16_3_;
output.val_f16_3_ = _e67 + _e70;
metal::half4 _e76 = input_uniform.val_f16_4_;
metal::half4 _e79 = input_storage.val_f16_4_;
output.val_f16_4_ = _e76 + _e79;
metal::half2x2 _e85 = input_uniform.val_mat2x2_;
metal::half2x2 _e88 = input_storage.val_mat2x2_;
output.val_mat2x2_ = _e85 + _e88;
metal::half2x3 _e94 = input_uniform.val_mat2x3_;
metal::half2x3 _e97 = input_storage.val_mat2x3_;
output.val_mat2x3_ = _e94 + _e97;
metal::half2x4 _e103 = input_uniform.val_mat2x4_;
metal::half2x4 _e106 = input_storage.val_mat2x4_;
output.val_mat2x4_ = _e103 + _e106;
metal::half3x2 _e112 = input_uniform.val_mat3x2_;
metal::half3x2 _e115 = input_storage.val_mat3x2_;
output.val_mat3x2_ = _e112 + _e115;
metal::half3x3 _e121 = input_uniform.val_mat3x3_;
metal::half3x3 _e124 = input_storage.val_mat3x3_;
output.val_mat3x3_ = _e121 + _e124;
metal::half3x4 _e130 = input_uniform.val_mat3x4_;
metal::half3x4 _e133 = input_storage.val_mat3x4_;
output.val_mat3x4_ = _e130 + _e133;
metal::half4x2 _e139 = input_uniform.val_mat4x2_;
metal::half4x2 _e142 = input_storage.val_mat4x2_;
output.val_mat4x2_ = _e139 + _e142;
metal::half4x3 _e148 = input_uniform.val_mat4x3_;
metal::half4x3 _e151 = input_storage.val_mat4x3_;
output.val_mat4x3_ = _e148 + _e151;
metal::half4x4 _e157 = input_uniform.val_mat4x4_;
metal::half4x4 _e160 = input_storage.val_mat4x4_;
output.val_mat4x4_ = _e157 + _e160;
type_16 _e166 = input_arrays.val_f16_array_2_;
output_arrays.val_f16_array_2_ = _e166;
half _e167 = val;
half _e51 = input_uniform.val_f16_;
half _e54 = input_storage.val_f16_;
output.val_f16_ = _e51 + _e54;
metal::half2 _e60 = input_uniform.val_f16_2_;
metal::half2 _e63 = input_storage.val_f16_2_;
output.val_f16_2_ = _e60 + _e63;
metal::half3 _e69 = input_uniform.val_f16_3_;
metal::half3 _e72 = input_storage.val_f16_3_;
output.val_f16_3_ = _e69 + _e72;
metal::half4 _e78 = input_uniform.val_f16_4_;
metal::half4 _e81 = input_storage.val_f16_4_;
output.val_f16_4_ = _e78 + _e81;
metal::half2x2 _e87 = input_uniform.val_mat2x2_;
metal::half2x2 _e90 = input_storage.val_mat2x2_;
output.val_mat2x2_ = _e87 + _e90;
metal::half2x3 _e96 = input_uniform.val_mat2x3_;
metal::half2x3 _e99 = input_storage.val_mat2x3_;
output.val_mat2x3_ = _e96 + _e99;
metal::half2x4 _e105 = input_uniform.val_mat2x4_;
metal::half2x4 _e108 = input_storage.val_mat2x4_;
output.val_mat2x4_ = _e105 + _e108;
metal::half3x2 _e114 = input_uniform.val_mat3x2_;
metal::half3x2 _e117 = input_storage.val_mat3x2_;
output.val_mat3x2_ = _e114 + _e117;
metal::half3x3 _e123 = input_uniform.val_mat3x3_;
metal::half3x3 _e126 = input_storage.val_mat3x3_;
output.val_mat3x3_ = _e123 + _e126;
metal::half3x4 _e132 = input_uniform.val_mat3x4_;
metal::half3x4 _e135 = input_storage.val_mat3x4_;
output.val_mat3x4_ = _e132 + _e135;
metal::half4x2 _e141 = input_uniform.val_mat4x2_;
metal::half4x2 _e144 = input_storage.val_mat4x2_;
output.val_mat4x2_ = _e141 + _e144;
metal::half4x3 _e150 = input_uniform.val_mat4x3_;
metal::half4x3 _e153 = input_storage.val_mat4x3_;
output.val_mat4x3_ = _e150 + _e153;
metal::half4x4 _e159 = input_uniform.val_mat4x4_;
metal::half4x4 _e162 = input_storage.val_mat4x4_;
output.val_mat4x4_ = _e159 + _e162;
type_16 _e168 = input_arrays.val_f16_array_2_;
output_arrays.val_f16_array_2_ = _e168;
half _e169 = val;
val = _e169 + metal::abs(_e167);
half _e171 = val;
half _e172 = val;
val = _e171 + metal::abs(_e169);
half _e173 = val;
half _e174 = val;
half _e175 = val;
val = _e175 + metal::clamp(_e171, _e172, _e173);
half _e177 = val;
val = _e177 + metal::clamp(_e173, _e174, _e175);
half _e179 = val;
half _e182 = val;
val = _e182 + metal::dot(metal::half2(_e177), metal::half2(_e179));
half _e181 = val;
half _e184 = val;
half _e185 = val;
val = _e184 + metal::dot(metal::half2(_e179), metal::half2(_e181));
half _e186 = val;
half _e187 = val;
val = _e187 + metal::max(_e184, _e185);
half _e189 = val;
half _e190 = val;
val = _e189 + metal::max(_e186, _e187);
half _e191 = val;
half _e192 = val;
val = _e192 + metal::min(_e189, _e190);
half _e194 = val;
val = _e194 + metal::min(_e191, _e192);
half _e196 = val;
val = _e196 + metal::sign(_e194);
half _e199 = val;
val = _e199 + 1.0h;
metal::half2 _e203 = input_uniform.val_f16_2_;
metal::float2 float_vec2_ = static_cast<metal::float2>(_e203);
half _e198 = val;
val = _e198 + metal::sign(_e196);
half _e201 = val;
val = _e201 + 1.0h;
metal::half2 _e205 = input_uniform.val_f16_2_;
metal::float2 float_vec2_ = static_cast<metal::float2>(_e205);
output.val_f16_2_ = static_cast<metal::half2>(float_vec2_);
metal::half3 _e210 = input_uniform.val_f16_3_;
metal::float3 float_vec3_ = static_cast<metal::float3>(_e210);
metal::half3 _e212 = input_uniform.val_f16_3_;
metal::float3 float_vec3_ = static_cast<metal::float3>(_e212);
output.val_f16_3_ = static_cast<metal::half3>(float_vec3_);
metal::half4 _e217 = input_uniform.val_f16_4_;
metal::float4 float_vec4_ = static_cast<metal::float4>(_e217);
metal::half4 _e219 = input_uniform.val_f16_4_;
metal::float4 float_vec4_ = static_cast<metal::float4>(_e219);
output.val_f16_4_ = static_cast<metal::half4>(float_vec4_);
metal::half2x2 _e226 = input_uniform.val_mat2x2_;
output.val_mat2x2_ = metal::half2x2(metal::float2x2(_e226));
metal::half2x3 _e233 = input_uniform.val_mat2x3_;
output.val_mat2x3_ = metal::half2x3(metal::float2x3(_e233));
metal::half2x4 _e240 = input_uniform.val_mat2x4_;
output.val_mat2x4_ = metal::half2x4(metal::float2x4(_e240));
metal::half3x2 _e247 = input_uniform.val_mat3x2_;
output.val_mat3x2_ = metal::half3x2(metal::float3x2(_e247));
metal::half3x3 _e254 = input_uniform.val_mat3x3_;
output.val_mat3x3_ = metal::half3x3(metal::float3x3(_e254));
metal::half3x4 _e261 = input_uniform.val_mat3x4_;
output.val_mat3x4_ = metal::half3x4(metal::float3x4(_e261));
metal::half4x2 _e268 = input_uniform.val_mat4x2_;
output.val_mat4x2_ = metal::half4x2(metal::float4x2(_e268));
metal::half4x3 _e275 = input_uniform.val_mat4x3_;
output.val_mat4x3_ = metal::half4x3(metal::float4x3(_e275));
metal::half4x4 _e282 = input_uniform.val_mat4x4_;
output.val_mat4x4_ = metal::half4x4(metal::float4x4(_e282));
half _e285 = val;
return _e285;
metal::half2x2 _e228 = input_uniform.val_mat2x2_;
output.val_mat2x2_ = metal::half2x2(metal::float2x2(_e228));
metal::half2x3 _e235 = input_uniform.val_mat2x3_;
output.val_mat2x3_ = metal::half2x3(metal::float2x3(_e235));
metal::half2x4 _e242 = input_uniform.val_mat2x4_;
output.val_mat2x4_ = metal::half2x4(metal::float2x4(_e242));
metal::half3x2 _e249 = input_uniform.val_mat3x2_;
output.val_mat3x2_ = metal::half3x2(metal::float3x2(_e249));
metal::half3x3 _e256 = input_uniform.val_mat3x3_;
output.val_mat3x3_ = metal::half3x3(metal::float3x3(_e256));
metal::half3x4 _e263 = input_uniform.val_mat3x4_;
output.val_mat3x4_ = metal::half3x4(metal::float3x4(_e263));
metal::half4x2 _e270 = input_uniform.val_mat4x2_;
output.val_mat4x2_ = metal::half4x2(metal::float4x2(_e270));
metal::half4x3 _e277 = input_uniform.val_mat4x3_;
output.val_mat4x3_ = metal::half4x3(metal::float4x3(_e277));
metal::half4x4 _e284 = input_uniform.val_mat4x4_;
output.val_mat4x4_ = metal::half4x4(metal::float4x4(_e284));
half _e287 = val;
return _e287;
}
kernel void main_(
@ -177,7 +180,8 @@ kernel void main_(
, device UniformCompatible& output [[user(fake0)]]
, device StorageCompatible& output_arrays [[user(fake0)]]
) {
half _e3 = f16_function(2.0h, input_uniform, input_storage, input_arrays, output, output_arrays);
half private_variable = 1.0h;
half _e3 = f16_function(2.0h, private_variable, input_uniform, input_storage, input_arrays, output, output_arrays);
output.final_value = _e3;
return;
}

2
naga/tests/out/msl/wgsl-image.msl
View File

@ -18,6 +18,7 @@ kernel void main_(
, metal::texture2d_ms<uint, metal::access::read> image_multisampled_src [[user(fake0)]]
, metal::texture2d<uint, metal::access::read> image_storage_src [[user(fake0)]]
, metal::texture2d_array<uint, metal::access::sample> image_array_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::read> image_dup_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::sample> image_1d_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
) {
@ -31,6 +32,7 @@ kernel void main_(
metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), local_id.z, as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
metal::uint4 value6_ = image_array_src.read(metal::uint2(itc), static_cast<int>(local_id.z), as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
metal::uint4 value7_ = image_1d_src.read(uint(static_cast<int>(local_id.x)));
metal::uint4 value8_ = image_dup_src.read(uint(static_cast<int>(local_id.x)));
metal::uint4 value1u = image_mipmapped_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z));
metal::uint4 value2u = image_multisampled_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z));
metal::uint4 value3u = image_multisampled_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<uint>(local_id.z));

125
naga/tests/out/msl/wgsl-int64.msl
View File

@ -44,6 +44,7 @@ long naga_abs(long val) {
long int64_function(
long x,
thread long& private_variable,
constant UniformCompatible& input_uniform,
device UniformCompatible const& input_storage,
device StorageCompatible const& input_arrays,
@ -51,81 +52,82 @@ long int64_function(
device StorageCompatible& output_arrays
) {
long val = 20L;
long _e8 = val;
val = as_type<long>(as_type<ulong>(_e8) + as_type<ulong>(as_type<long>(as_type<ulong>(as_type<long>(as_type<ulong>(31L) - as_type<ulong>(1002003004005006L))) + as_type<ulong>(-9223372036854775807L))));
long phony = private_variable;
long _e10 = val;
long _e13 = val;
val = as_type<long>(as_type<ulong>(_e13) + as_type<ulong>(as_type<long>(as_type<ulong>(_e10) + as_type<ulong>(5L))));
uint _e17 = input_uniform.val_u32_;
long _e18 = val;
long _e22 = val;
val = as_type<long>(as_type<ulong>(_e22) + as_type<ulong>(static_cast<long>(_e17 + static_cast<uint>(_e18))));
int _e26 = input_uniform.val_i32_;
long _e27 = val;
long _e31 = val;
val = as_type<long>(as_type<ulong>(_e31) + as_type<ulong>(static_cast<long>(as_type<int>(as_type<uint>(_e26) + as_type<uint>(static_cast<int>(_e27))))));
float _e35 = input_uniform.val_f32_;
long _e36 = val;
long _e40 = val;
val = as_type<long>(as_type<ulong>(_e40) + as_type<ulong>(naga_f2i64(_e35 + static_cast<float>(_e36))));
long _e44 = input_uniform.val_i64_;
long _e47 = val;
val = as_type<long>(as_type<ulong>(_e47) + as_type<ulong>(metal::long3(_e44).z));
ulong _e51 = input_uniform.val_u64_;
long _e53 = val;
val = as_type<long>(as_type<ulong>(_e53) + as_type<ulong>(as_type<long>(_e51)));
metal::ulong2 _e57 = input_uniform.val_u64_2_;
long _e60 = val;
val = as_type<long>(as_type<ulong>(_e60) + as_type<ulong>(as_type<metal::long2>(_e57).y));
metal::ulong3 _e64 = input_uniform.val_u64_3_;
long _e67 = val;
val = as_type<long>(as_type<ulong>(_e67) + as_type<ulong>(as_type<metal::long3>(_e64).z));
metal::ulong4 _e71 = input_uniform.val_u64_4_;
long _e74 = val;
val = as_type<long>(as_type<ulong>(_e74) + as_type<ulong>(as_type<metal::long4>(_e71).w));
long _e77 = val;
val = as_type<long>(as_type<ulong>(_e77) + as_type<ulong>((-9223372036854775807L - 1L)));
long _e83 = input_uniform.val_i64_;
long _e86 = input_storage.val_i64_;
output.val_i64_ = as_type<long>(as_type<ulong>(_e83) + as_type<ulong>(_e86));
metal::long2 _e92 = input_uniform.val_i64_2_;
metal::long2 _e95 = input_storage.val_i64_2_;
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e92) + as_type<metal::ulong2>(_e95));
metal::long3 _e101 = input_uniform.val_i64_3_;
metal::long3 _e104 = input_storage.val_i64_3_;
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e101) + as_type<metal::ulong3>(_e104));
metal::long4 _e110 = input_uniform.val_i64_4_;
metal::long4 _e113 = input_storage.val_i64_4_;
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e110) + as_type<metal::ulong4>(_e113));
type_12 _e119 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e119;
long _e120 = val;
val = as_type<long>(as_type<ulong>(_e10) + as_type<ulong>(as_type<long>(as_type<ulong>(as_type<long>(as_type<ulong>(31L) - as_type<ulong>(1002003004005006L))) + as_type<ulong>(-9223372036854775807L))));
long _e12 = val;
long _e15 = val;
val = as_type<long>(as_type<ulong>(_e15) + as_type<ulong>(as_type<long>(as_type<ulong>(_e12) + as_type<ulong>(5L))));
uint _e19 = input_uniform.val_u32_;
long _e20 = val;
long _e24 = val;
val = as_type<long>(as_type<ulong>(_e24) + as_type<ulong>(static_cast<long>(_e19 + static_cast<uint>(_e20))));
int _e28 = input_uniform.val_i32_;
long _e29 = val;
long _e33 = val;
val = as_type<long>(as_type<ulong>(_e33) + as_type<ulong>(static_cast<long>(as_type<int>(as_type<uint>(_e28) + as_type<uint>(static_cast<int>(_e29))))));
float _e37 = input_uniform.val_f32_;
long _e38 = val;
long _e42 = val;
val = as_type<long>(as_type<ulong>(_e42) + as_type<ulong>(naga_f2i64(_e37 + static_cast<float>(_e38))));
long _e46 = input_uniform.val_i64_;
long _e49 = val;
val = as_type<long>(as_type<ulong>(_e49) + as_type<ulong>(metal::long3(_e46).z));
ulong _e53 = input_uniform.val_u64_;
long _e55 = val;
val = as_type<long>(as_type<ulong>(_e55) + as_type<ulong>(as_type<long>(_e53)));
metal::ulong2 _e59 = input_uniform.val_u64_2_;
long _e62 = val;
val = as_type<long>(as_type<ulong>(_e62) + as_type<ulong>(as_type<metal::long2>(_e59).y));
metal::ulong3 _e66 = input_uniform.val_u64_3_;
long _e69 = val;
val = as_type<long>(as_type<ulong>(_e69) + as_type<ulong>(as_type<metal::long3>(_e66).z));
metal::ulong4 _e73 = input_uniform.val_u64_4_;
long _e76 = val;
val = as_type<long>(as_type<ulong>(_e76) + as_type<ulong>(as_type<metal::long4>(_e73).w));
long _e79 = val;
val = as_type<long>(as_type<ulong>(_e79) + as_type<ulong>((-9223372036854775807L - 1L)));
long _e85 = input_uniform.val_i64_;
long _e88 = input_storage.val_i64_;
output.val_i64_ = as_type<long>(as_type<ulong>(_e85) + as_type<ulong>(_e88));
metal::long2 _e94 = input_uniform.val_i64_2_;
metal::long2 _e97 = input_storage.val_i64_2_;
output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e94) + as_type<metal::ulong2>(_e97));
metal::long3 _e103 = input_uniform.val_i64_3_;
metal::long3 _e106 = input_storage.val_i64_3_;
output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e103) + as_type<metal::ulong3>(_e106));
metal::long4 _e112 = input_uniform.val_i64_4_;
metal::long4 _e115 = input_storage.val_i64_4_;
output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e112) + as_type<metal::ulong4>(_e115));
type_12 _e121 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e121;
long _e122 = val;
val = as_type<long>(as_type<ulong>(_e122) + as_type<ulong>(naga_abs(_e120)));
long _e124 = val;
long _e125 = val;
val = as_type<long>(as_type<ulong>(_e124) + as_type<ulong>(naga_abs(_e122)));
long _e126 = val;
long _e127 = val;
long _e128 = val;
val = as_type<long>(as_type<ulong>(_e128) + as_type<ulong>(metal::clamp(_e124, _e125, _e126)));
long _e130 = val;
metal::long2 _e131 = metal::long2(_e130);
val = as_type<long>(as_type<ulong>(_e130) + as_type<ulong>(metal::clamp(_e126, _e127, _e128)));
long _e132 = val;
metal::long2 _e133 = metal::long2(_e132);
long _e135 = val;
val = as_type<long>(as_type<ulong>(_e135) + as_type<ulong>(( + _e131.x * _e133.x + _e131.y * _e133.y)));
long _e134 = val;
metal::long2 _e135 = metal::long2(_e134);
long _e137 = val;
long _e138 = val;
val = as_type<long>(as_type<ulong>(_e137) + as_type<ulong>(( + _e133.x * _e135.x + _e133.y * _e135.y)));
long _e139 = val;
long _e140 = val;
val = as_type<long>(as_type<ulong>(_e140) + as_type<ulong>(metal::max(_e137, _e138)));
long _e142 = val;
long _e143 = val;
val = as_type<long>(as_type<ulong>(_e142) + as_type<ulong>(metal::max(_e139, _e140)));
long _e144 = val;
long _e145 = val;
val = as_type<long>(as_type<ulong>(_e145) + as_type<ulong>(metal::min(_e142, _e143)));
long _e147 = val;
val = as_type<long>(as_type<ulong>(_e147) + as_type<ulong>(metal::min(_e144, _e145)));
long _e149 = val;
val = as_type<long>(as_type<ulong>(_e149) + as_type<ulong>(metal::select(metal::select(long(-1), long(1), (_e147 > 0)), long(0), (_e147 == 0))));
long _e151 = val;
return _e151;
val = as_type<long>(as_type<ulong>(_e151) + as_type<ulong>(metal::select(metal::select(long(-1), long(1), (_e149 > 0)), long(0), (_e149 == 0))));
long _e153 = val;
return _e153;
}
ulong naga_f2u64(float value) {
@ -220,8 +222,9 @@ kernel void main_(
, device UniformCompatible& output [[user(fake0)]]
, device StorageCompatible& output_arrays [[user(fake0)]]
) {
long private_variable = 1L;
ulong _e3 = uint64_function(67uL, input_uniform, input_storage, input_arrays, output, output_arrays);
long _e5 = int64_function(60L, input_uniform, input_storage, input_arrays, output, output_arrays);
long _e5 = int64_function(60L, private_variable, input_uniform, input_storage, input_arrays, output, output_arrays);
output.final_value = _e3 + as_type<ulong>(_e5);
return;
}

728
naga/tests/out/spv/wgsl-abstract-types-var.spvasm
View File

@ -1,12 +1,12 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 355
; Bound: 412
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %351 "main"
OpExecutionMode %351 LocalSize 1 1 1
OpEntryPoint GLCompute %407 "main"
OpExecutionMode %407 LocalSize 1 1 1
OpDecorate %10 ArrayStride 4
OpDecorate %12 ArrayStride 4
OpDecorate %13 ArrayStride 4
@ -133,342 +133,402 @@ OpDecorate %18 ArrayStride 16
%120 = OpVariable %99 Private %57
%121 = OpVariable %101 Private %59
%124 = OpTypeFunction %2
%126 = OpTypePointer Function %4
%128 = OpTypePointer Function %6
%130 = OpTypePointer Function %8
%141 = OpTypePointer Function %9
%158 = OpTypePointer Function %10
%163 = OpTypePointer Function %12
%171 = OpTypePointer Function %15
%173 = OpTypePointer Function %18
%187 = OpTypePointer Function %5
%188 = OpConstantNull %5
%190 = OpTypePointer Function %3
%191 = OpConstantNull %3
%193 = OpTypePointer Function %7
%194 = OpConstantNull %7
%196 = OpConstantNull %6
%198 = OpConstantNull %6
%200 = OpConstantNull %8
%202 = OpConstantNull %8
%204 = OpConstantNull %6
%206 = OpConstantNull %6
%208 = OpConstantNull %9
%210 = OpConstantNull %9
%212 = OpConstantNull %9
%214 = OpConstantNull %9
%216 = OpConstantNull %10
%218 = OpConstantNull %10
%220 = OpConstantNull %10
%222 = OpConstantNull %10
%224 = OpConstantNull %12
%226 = OpConstantNull %12
%228 = OpConstantNull %10
%230 = OpConstantNull %10
%232 = OpConstantNull %10
%234 = OpConstantNull %10
%236 = OpConstantNull %12
%238 = OpConstantNull %12
%240 = OpConstantNull %4
%242 = OpConstantNull %6
%244 = OpConstantNull %8
%182 = OpTypePointer Function %4
%184 = OpTypePointer Function %6
%186 = OpTypePointer Function %8
%197 = OpTypePointer Function %9
%214 = OpTypePointer Function %10
%219 = OpTypePointer Function %12
%227 = OpTypePointer Function %15
%229 = OpTypePointer Function %18
%243 = OpTypePointer Function %5
%244 = OpConstantNull %5
%246 = OpTypePointer Function %3
%247 = OpConstantNull %3
%249 = OpTypePointer Function %7
%250 = OpConstantNull %7
%252 = OpConstantNull %6
%254 = OpConstantNull %6
%256 = OpConstantNull %8
%258 = OpConstantNull %8
%260 = OpConstantNull %6
%262 = OpConstantNull %6
%264 = OpConstantNull %9
%266 = OpConstantNull %9
%268 = OpConstantNull %9
%270 = OpConstantNull %9
%272 = OpConstantNull %10
%274 = OpConstantNull %10
%276 = OpConstantNull %10
%278 = OpConstantNull %10
%280 = OpConstantNull %12
%282 = OpConstantNull %12
%284 = OpConstantNull %10
%286 = OpConstantNull %10
%288 = OpConstantNull %10
%290 = OpConstantNull %10
%292 = OpConstantNull %12
%294 = OpConstantNull %12
%296 = OpConstantNull %4
%298 = OpConstantNull %6
%300 = OpConstantNull %8
%123 = OpFunction %2 None %124
%122 = OpLabel
%181 = OpVariable %158 Function %52
%178 = OpVariable %130 Function %50
%175 = OpVariable %126 Function %49
%170 = OpVariable %171 Function %57
%167 = OpVariable %158 Function %52
%164 = OpVariable %163 Function %54
%160 = OpVariable %158 Function %52
%156 = OpVariable %130 Function %50
%153 = OpVariable %126 Function %49
%150 = OpVariable %141 Function %47
%147 = OpVariable %141 Function %47
%144 = OpVariable %141 Function %47
%140 = OpVariable %141 Function %40
%137 = OpVariable %126 Function %35
%134 = OpVariable %128 Function %33
%131 = OpVariable %130 Function %30
%125 = OpVariable %126 Function %21
%182 = OpVariable %158 Function %52
%179 = OpVariable %163 Function %54
%176 = OpVariable %128 Function %51
%172 = OpVariable %173 Function %59
%168 = OpVariable %158 Function %52
%165 = OpVariable %163 Function %54
%161 = OpVariable %158 Function %52
%157 = OpVariable %158 Function %52
%154 = OpVariable %128 Function %51
%151 = OpVariable %126 Function %49
%148 = OpVariable %141 Function %47
%145 = OpVariable %141 Function %47
%142 = OpVariable %141 Function %47
%138 = OpVariable %128 Function %37
%135 = OpVariable %128 Function %33
%132 = OpVariable %130 Function %30
%127 = OpVariable %128 Function %24
%180 = OpVariable %158 Function %52
%177 = OpVariable %130 Function %50
%174 = OpVariable %173 Function %59
%169 = OpVariable %158 Function %52
%166 = OpVariable %158 Function %52
%162 = OpVariable %163 Function %54
%159 = OpVariable %158 Function %52
%155 = OpVariable %130 Function %50
%152 = OpVariable %130 Function %50
%149 = OpVariable %141 Function %47
%146 = OpVariable %141 Function %47
%143 = OpVariable %141 Function %47
%139 = OpVariable %130 Function %39
%136 = OpVariable %128 Function %33
%133 = OpVariable %128 Function %33
%129 = OpVariable %130 Function %27
OpBranch %183
%183 = OpLabel
OpStore %125 %21
OpStore %127 %24
OpStore %129 %27
OpStore %131 %30
OpStore %132 %30
OpStore %133 %33
OpStore %134 %33
OpStore %135 %33
OpStore %136 %33
OpStore %137 %35
OpStore %138 %37
OpStore %139 %39
OpStore %140 %40
OpStore %142 %47
OpStore %143 %47
OpStore %144 %47
OpStore %145 %47
OpStore %146 %47
OpStore %147 %47
OpStore %148 %47
OpStore %149 %47
OpStore %150 %47
OpStore %151 %49
OpStore %152 %50
OpStore %153 %49
OpStore %154 %51
OpStore %155 %50
OpStore %156 %50
OpStore %157 %52
OpStore %159 %52
OpStore %160 %52
OpStore %161 %52
OpStore %162 %54
OpStore %164 %54
OpStore %165 %54
OpStore %166 %52
OpStore %167 %52
OpStore %168 %52
OpStore %169 %52
OpStore %170 %57
OpStore %172 %59
OpStore %174 %59
OpStore %175 %49
OpStore %176 %51
OpStore %177 %50
OpStore %178 %50
OpStore %179 %54
OpStore %180 %52
OpStore %181 %52
OpStore %182 %52
OpBranch %125
%125 = OpLabel
%126 = OpLoad %4 %60
%127 = OpLoad %6 %62
%128 = OpLoad %8 %64
%129 = OpLoad %8 %66
%130 = OpLoad %8 %67
%131 = OpLoad %6 %68
%132 = OpLoad %6 %69
%133 = OpLoad %6 %70
%134 = OpLoad %6 %71
%135 = OpLoad %4 %72
%136 = OpLoad %6 %73
%137 = OpLoad %8 %74
%138 = OpLoad %9 %75
%139 = OpLoad %9 %77
%140 = OpLoad %9 %78
%141 = OpLoad %9 %79
%142 = OpLoad %9 %80
%143 = OpLoad %9 %81
%144 = OpLoad %4 %82
%145 = OpLoad %8 %83
%146 = OpLoad %4 %84
%147 = OpLoad %6 %85
%148 = OpLoad %8 %86
%149 = OpLoad %8 %87
%150 = OpLoad %10 %88
%151 = OpLoad %10 %90
%152 = OpLoad %12 %91
%153 = OpLoad %13 %93
%154 = OpLoad %10 %95
%155 = OpLoad %10 %96
%156 = OpLoad %10 %97
%157 = OpLoad %15 %98
%158 = OpLoad %18 %100
%159 = OpLoad %18 %102
%160 = OpLoad %4 %103
%161 = OpLoad %6 %104
%162 = OpLoad %8 %105
%163 = OpLoad %8 %106
%164 = OpLoad %4 %107
%165 = OpLoad %8 %108
%166 = OpLoad %4 %109
%167 = OpLoad %6 %110
%168 = OpLoad %8 %111
%169 = OpLoad %8 %112
%170 = OpLoad %10 %113
%171 = OpLoad %10 %114
%172 = OpLoad %12 %115
%173 = OpLoad %10 %116
%174 = OpLoad %10 %117
%175 = OpLoad %10 %118
%176 = OpLoad %15 %119
%177 = OpLoad %15 %120
%178 = OpLoad %18 %121
OpReturn
OpFunctionEnd
%185 = OpFunction %2 None %124
%184 = OpLabel
%239 = OpVariable %126 Function %240
%233 = OpVariable %158 Function %234
%227 = OpVariable %158 Function %228
%221 = OpVariable %158 Function %222
%215 = OpVariable %158 Function %216
%209 = OpVariable %141 Function %210
%203 = OpVariable %128 Function %204
%197 = OpVariable %128 Function %198
%189 = OpVariable %190 Function %191
%243 = OpVariable %130 Function %244
%237 = OpVariable %163 Function %238
%231 = OpVariable %158 Function %232
%225 = OpVariable %163 Function %226
%219 = OpVariable %158 Function %220
%213 = OpVariable %141 Function %214
%207 = OpVariable %141 Function %208
%201 = OpVariable %130 Function %202
%195 = OpVariable %128 Function %196
%186 = OpVariable %187 Function %188
%241 = OpVariable %128 Function %242
%235 = OpVariable %163 Function %236
%229 = OpVariable %158 Function %230
%223 = OpVariable %163 Function %224
%217 = OpVariable %158 Function %218
%211 = OpVariable %141 Function %212
%205 = OpVariable %128 Function %206
%199 = OpVariable %130 Function %200
%192 = OpVariable %193 Function %194
OpBranch %245
%245 = OpLabel
%246 = OpLoad %5 %186
%247 = OpCompositeConstruct %6 %246 %32
OpStore %195 %247
%248 = OpLoad %5 %186
%249 = OpCompositeConstruct %6 %31 %248
OpStore %197 %249
%250 = OpLoad %7 %192
%251 = OpCompositeConstruct %8 %250 %26
OpStore %199 %251
%252 = OpLoad %7 %192
%253 = OpCompositeConstruct %8 %252 %29
OpStore %201 %253
%254 = OpLoad %5 %186
%255 = OpCompositeConstruct %6 %254 %32
OpStore %203 %255
%256 = OpLoad %5 %186
%257 = OpCompositeConstruct %6 %31 %256
OpStore %205 %257
%258 = OpLoad %7 %192
%259 = OpCompositeConstruct %8 %258 %42
%260 = OpCompositeConstruct %9 %259 %46
OpStore %207 %260
%261 = OpLoad %7 %192
%262 = OpCompositeConstruct %8 %41 %261
%263 = OpCompositeConstruct %9 %262 %46
OpStore %209 %263
%264 = OpLoad %7 %192
%265 = OpCompositeConstruct %8 %264 %45
%266 = OpCompositeConstruct %9 %43 %265
OpStore %211 %266
%267 = OpLoad %7 %192
%268 = OpCompositeConstruct %8 %44 %267
%269 = OpCompositeConstruct %9 %43 %268
OpStore %213 %269
%270 = OpLoad %7 %192
%271 = OpCompositeConstruct %10 %270 %42
OpStore %215 %271
%272 = OpLoad %7 %192
%273 = OpCompositeConstruct %10 %41 %272
OpStore %217 %273
%274 = OpLoad %7 %192
%275 = OpCompositeConstruct %10 %274 %42
OpStore %219 %275
%276 = OpLoad %7 %192
%277 = OpCompositeConstruct %10 %41 %276
OpStore %221 %277
%278 = OpLoad %3 %189
%279 = OpCompositeConstruct %12 %278 %53
OpStore %223 %279
%280 = OpLoad %3 %189
%281 = OpCompositeConstruct %12 %48 %280
OpStore %225 %281
%282 = OpLoad %7 %192
%283 = OpCompositeConstruct %10 %282 %42
OpStore %227 %283
%284 = OpLoad %7 %192
%285 = OpCompositeConstruct %10 %41 %284
OpStore %229 %285
%286 = OpLoad %7 %192
%287 = OpCompositeConstruct %10 %286 %42
OpStore %231 %287
%288 = OpLoad %7 %192
%289 = OpCompositeConstruct %10 %41 %288
OpStore %233 %289
%290 = OpLoad %3 %189
%291 = OpCompositeConstruct %12 %290 %53
OpStore %235 %291
%292 = OpLoad %3 %189
%293 = OpCompositeConstruct %12 %48 %292
OpStore %237 %293
%294 = OpLoad %3 %189
%295 = OpCompositeConstruct %4 %294 %294
OpStore %239 %295
%296 = OpLoad %5 %186
%297 = OpCompositeConstruct %6 %296 %296
OpStore %241 %297
%298 = OpLoad %7 %192
%299 = OpCompositeConstruct %8 %298 %298
OpStore %243 %299
%300 = OpLoad %5 %186
%301 = OpCompositeConstruct %6 %300 %32
OpStore %195 %301
%302 = OpLoad %5 %186
%303 = OpCompositeConstruct %6 %31 %302
OpStore %197 %303
%304 = OpLoad %5 %186
%305 = OpCompositeConstruct %6 %304 %32
OpStore %203 %305
%306 = OpLoad %5 %186
%307 = OpCompositeConstruct %6 %31 %306
OpStore %205 %307
%308 = OpLoad %7 %192
%309 = OpCompositeConstruct %8 %308 %42
%310 = OpCompositeConstruct %9 %309 %46
OpStore %207 %310
%311 = OpLoad %7 %192
%312 = OpCompositeConstruct %8 %41 %311
%313 = OpCompositeConstruct %9 %312 %46
OpStore %209 %313
%314 = OpLoad %7 %192
%315 = OpCompositeConstruct %8 %314 %45
%316 = OpCompositeConstruct %9 %43 %315
OpStore %211 %316
%317 = OpLoad %7 %192
%318 = OpCompositeConstruct %8 %44 %317
%319 = OpCompositeConstruct %9 %43 %318
OpStore %213 %319
%320 = OpLoad %7 %192
%321 = OpCompositeConstruct %10 %320 %42
OpStore %215 %321
%322 = OpLoad %7 %192
%323 = OpCompositeConstruct %10 %41 %322
OpStore %217 %323
%324 = OpLoad %7 %192
%325 = OpCompositeConstruct %10 %324 %42
OpStore %219 %325
%326 = OpLoad %7 %192
%327 = OpCompositeConstruct %10 %41 %326
OpStore %221 %327
%328 = OpLoad %3 %189
%329 = OpCompositeConstruct %12 %328 %53
OpStore %223 %329
%330 = OpLoad %3 %189
%331 = OpCompositeConstruct %12 %48 %330
OpStore %225 %331
%332 = OpLoad %7 %192
%333 = OpCompositeConstruct %10 %332 %42
OpStore %227 %333
%334 = OpLoad %7 %192
%335 = OpCompositeConstruct %10 %41 %334
OpStore %229 %335
%336 = OpLoad %7 %192
%337 = OpCompositeConstruct %10 %336 %42
OpStore %231 %337
%338 = OpLoad %7 %192
%339 = OpCompositeConstruct %10 %41 %338
OpStore %233 %339
%340 = OpLoad %3 %189
%341 = OpCompositeConstruct %12 %340 %53
OpStore %235 %341
%342 = OpLoad %3 %189
%343 = OpCompositeConstruct %12 %48 %342
OpStore %237 %343
%344 = OpLoad %3 %189
%345 = OpCompositeConstruct %4 %344 %344
OpStore %239 %345
%346 = OpLoad %5 %186
%347 = OpCompositeConstruct %6 %346 %346
OpStore %241 %347
%348 = OpLoad %7 %192
%349 = OpCompositeConstruct %8 %348 %348
OpStore %243 %349
%180 = OpFunction %2 None %124
%179 = OpLabel
%237 = OpVariable %214 Function %52
%234 = OpVariable %186 Function %50
%231 = OpVariable %182 Function %49
%226 = OpVariable %227 Function %57
%223 = OpVariable %214 Function %52
%220 = OpVariable %219 Function %54
%216 = OpVariable %214 Function %52
%212 = OpVariable %186 Function %50
%209 = OpVariable %182 Function %49
%206 = OpVariable %197 Function %47
%203 = OpVariable %197 Function %47
%200 = OpVariable %197 Function %47
%196 = OpVariable %197 Function %40
%193 = OpVariable %182 Function %35
%190 = OpVariable %184 Function %33
%187 = OpVariable %186 Function %30
%181 = OpVariable %182 Function %21
%238 = OpVariable %214 Function %52
%235 = OpVariable %219 Function %54
%232 = OpVariable %184 Function %51
%228 = OpVariable %229 Function %59
%224 = OpVariable %214 Function %52
%221 = OpVariable %219 Function %54
%217 = OpVariable %214 Function %52
%213 = OpVariable %214 Function %52
%210 = OpVariable %184 Function %51
%207 = OpVariable %182 Function %49
%204 = OpVariable %197 Function %47
%201 = OpVariable %197 Function %47
%198 = OpVariable %197 Function %47
%194 = OpVariable %184 Function %37
%191 = OpVariable %184 Function %33
%188 = OpVariable %186 Function %30
%183 = OpVariable %184 Function %24
%236 = OpVariable %214 Function %52
%233 = OpVariable %186 Function %50
%230 = OpVariable %229 Function %59
%225 = OpVariable %214 Function %52
%222 = OpVariable %214 Function %52
%218 = OpVariable %219 Function %54
%215 = OpVariable %214 Function %52
%211 = OpVariable %186 Function %50
%208 = OpVariable %186 Function %50
%205 = OpVariable %197 Function %47
%202 = OpVariable %197 Function %47
%199 = OpVariable %197 Function %47
%195 = OpVariable %186 Function %39
%192 = OpVariable %184 Function %33
%189 = OpVariable %184 Function %33
%185 = OpVariable %186 Function %27
OpBranch %239
%239 = OpLabel
OpStore %181 %21
OpStore %183 %24
OpStore %185 %27
OpStore %187 %30
OpStore %188 %30
OpStore %189 %33
OpStore %190 %33
OpStore %191 %33
OpStore %192 %33
OpStore %193 %35
OpStore %194 %37
OpStore %195 %39
OpStore %196 %40
OpStore %198 %47
OpStore %199 %47
OpStore %200 %47
OpStore %201 %47
OpStore %202 %47
OpStore %203 %47
OpStore %204 %47
OpStore %205 %47
OpStore %206 %47
OpStore %207 %49
OpStore %208 %50
OpStore %209 %49
OpStore %210 %51
OpStore %211 %50
OpStore %212 %50
OpStore %213 %52
OpStore %215 %52
OpStore %216 %52
OpStore %217 %52
OpStore %218 %54
OpStore %220 %54
OpStore %221 %54
OpStore %222 %52
OpStore %223 %52
OpStore %224 %52
OpStore %225 %52
OpStore %226 %57
OpStore %228 %59
OpStore %230 %59
OpStore %231 %49
OpStore %232 %51
OpStore %233 %50
OpStore %234 %50
OpStore %235 %54
OpStore %236 %52
OpStore %237 %52
OpStore %238 %52
OpReturn
OpFunctionEnd
%351 = OpFunction %2 None %124
%350 = OpLabel
OpBranch %352
%352 = OpLabel
%353 = OpFunctionCall %2 %123
%354 = OpFunctionCall %2 %185
%241 = OpFunction %2 None %124
%240 = OpLabel
%295 = OpVariable %182 Function %296
%289 = OpVariable %214 Function %290
%283 = OpVariable %214 Function %284
%277 = OpVariable %214 Function %278
%271 = OpVariable %214 Function %272
%265 = OpVariable %197 Function %266
%259 = OpVariable %184 Function %260
%253 = OpVariable %184 Function %254
%245 = OpVariable %246 Function %247
%299 = OpVariable %186 Function %300
%293 = OpVariable %219 Function %294
%287 = OpVariable %214 Function %288
%281 = OpVariable %219 Function %282
%275 = OpVariable %214 Function %276
%269 = OpVariable %197 Function %270
%263 = OpVariable %197 Function %264
%257 = OpVariable %186 Function %258
%251 = OpVariable %184 Function %252
%242 = OpVariable %243 Function %244
%297 = OpVariable %184 Function %298
%291 = OpVariable %219 Function %292
%285 = OpVariable %214 Function %286
%279 = OpVariable %219 Function %280
%273 = OpVariable %214 Function %274
%267 = OpVariable %197 Function %268
%261 = OpVariable %184 Function %262
%255 = OpVariable %186 Function %256
%248 = OpVariable %249 Function %250
OpBranch %301
%301 = OpLabel
%302 = OpLoad %5 %242
%303 = OpCompositeConstruct %6 %302 %32
OpStore %251 %303
%304 = OpLoad %5 %242
%305 = OpCompositeConstruct %6 %31 %304
OpStore %253 %305
%306 = OpLoad %7 %248
%307 = OpCompositeConstruct %8 %306 %26
OpStore %255 %307
%308 = OpLoad %7 %248
%309 = OpCompositeConstruct %8 %308 %29
OpStore %257 %309
%310 = OpLoad %5 %242
%311 = OpCompositeConstruct %6 %310 %32
OpStore %259 %311
%312 = OpLoad %5 %242
%313 = OpCompositeConstruct %6 %31 %312
OpStore %261 %313
%314 = OpLoad %7 %248
%315 = OpCompositeConstruct %8 %314 %42
%316 = OpCompositeConstruct %9 %315 %46
OpStore %263 %316
%317 = OpLoad %7 %248
%318 = OpCompositeConstruct %8 %41 %317
%319 = OpCompositeConstruct %9 %318 %46
OpStore %265 %319
%320 = OpLoad %7 %248
%321 = OpCompositeConstruct %8 %320 %45
%322 = OpCompositeConstruct %9 %43 %321
OpStore %267 %322
%323 = OpLoad %7 %248
%324 = OpCompositeConstruct %8 %44 %323
%325 = OpCompositeConstruct %9 %43 %324
OpStore %269 %325
%326 = OpLoad %7 %248
%327 = OpCompositeConstruct %10 %326 %42
OpStore %271 %327
%328 = OpLoad %7 %248
%329 = OpCompositeConstruct %10 %41 %328
OpStore %273 %329
%330 = OpLoad %7 %248
%331 = OpCompositeConstruct %10 %330 %42
OpStore %275 %331
%332 = OpLoad %7 %248
%333 = OpCompositeConstruct %10 %41 %332
OpStore %277 %333
%334 = OpLoad %3 %245
%335 = OpCompositeConstruct %12 %334 %53
OpStore %279 %335
%336 = OpLoad %3 %245
%337 = OpCompositeConstruct %12 %48 %336
OpStore %281 %337
%338 = OpLoad %7 %248
%339 = OpCompositeConstruct %10 %338 %42
OpStore %283 %339
%340 = OpLoad %7 %248
%341 = OpCompositeConstruct %10 %41 %340
OpStore %285 %341
%342 = OpLoad %7 %248
%343 = OpCompositeConstruct %10 %342 %42
OpStore %287 %343
%344 = OpLoad %7 %248
%345 = OpCompositeConstruct %10 %41 %344
OpStore %289 %345
%346 = OpLoad %3 %245
%347 = OpCompositeConstruct %12 %346 %53
OpStore %291 %347
%348 = OpLoad %3 %245
%349 = OpCompositeConstruct %12 %48 %348
OpStore %293 %349
%350 = OpLoad %3 %245
%351 = OpCompositeConstruct %4 %350 %350
OpStore %295 %351
%352 = OpLoad %5 %242
%353 = OpCompositeConstruct %6 %352 %352
OpStore %297 %353
%354 = OpLoad %7 %248
%355 = OpCompositeConstruct %8 %354 %354
OpStore %299 %355
%356 = OpLoad %5 %242
%357 = OpCompositeConstruct %6 %356 %32
OpStore %251 %357
%358 = OpLoad %5 %242
%359 = OpCompositeConstruct %6 %31 %358
OpStore %253 %359
%360 = OpLoad %5 %242
%361 = OpCompositeConstruct %6 %360 %32
OpStore %259 %361
%362 = OpLoad %5 %242
%363 = OpCompositeConstruct %6 %31 %362
OpStore %261 %363
%364 = OpLoad %7 %248
%365 = OpCompositeConstruct %8 %364 %42
%366 = OpCompositeConstruct %9 %365 %46
OpStore %263 %366
%367 = OpLoad %7 %248
%368 = OpCompositeConstruct %8 %41 %367
%369 = OpCompositeConstruct %9 %368 %46
OpStore %265 %369
%370 = OpLoad %7 %248
%371 = OpCompositeConstruct %8 %370 %45
%372 = OpCompositeConstruct %9 %43 %371
OpStore %267 %372
%373 = OpLoad %7 %248
%374 = OpCompositeConstruct %8 %44 %373
%375 = OpCompositeConstruct %9 %43 %374
OpStore %269 %375
%376 = OpLoad %7 %248
%377 = OpCompositeConstruct %10 %376 %42
OpStore %271 %377
%378 = OpLoad %7 %248
%379 = OpCompositeConstruct %10 %41 %378
OpStore %273 %379
%380 = OpLoad %7 %248
%381 = OpCompositeConstruct %10 %380 %42
OpStore %275 %381
%382 = OpLoad %7 %248
%383 = OpCompositeConstruct %10 %41 %382
OpStore %277 %383
%384 = OpLoad %3 %245
%385 = OpCompositeConstruct %12 %384 %53
OpStore %279 %385
%386 = OpLoad %3 %245
%387 = OpCompositeConstruct %12 %48 %386
OpStore %281 %387
%388 = OpLoad %7 %248
%389 = OpCompositeConstruct %10 %388 %42
OpStore %283 %389
%390 = OpLoad %7 %248
%391 = OpCompositeConstruct %10 %41 %390
OpStore %285 %391
%392 = OpLoad %7 %248
%393 = OpCompositeConstruct %10 %392 %42
OpStore %287 %393
%394 = OpLoad %7 %248
%395 = OpCompositeConstruct %10 %41 %394
OpStore %289 %395
%396 = OpLoad %3 %245
%397 = OpCompositeConstruct %12 %396 %53
OpStore %291 %397
%398 = OpLoad %3 %245
%399 = OpCompositeConstruct %12 %48 %398
OpStore %293 %399
%400 = OpLoad %3 %245
%401 = OpCompositeConstruct %4 %400 %400
OpStore %295 %401
%402 = OpLoad %5 %242
%403 = OpCompositeConstruct %6 %402 %402
OpStore %297 %403
%404 = OpLoad %7 %248
%405 = OpCompositeConstruct %8 %404 %404
OpStore %299 %405
OpReturn
OpFunctionEnd
%407 = OpFunction %2 None %124
%406 = OpLabel
OpBranch %408
%408 = OpLabel
%409 = OpFunctionCall %2 %123
%410 = OpFunctionCall %2 %180
%411 = OpFunctionCall %2 %241
OpReturn
OpFunctionEnd

413
naga/tests/out/spv/wgsl-access.spvasm
View File

@ -1,16 +1,16 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 415
; Bound: 416
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %333 "foo_vert" %328 %331
OpEntryPoint Fragment %388 "foo_frag" %387
OpEntryPoint GLCompute %406 "foo_compute"
OpExecutionMode %388 OriginUpperLeft
OpExecutionMode %406 LocalSize 1 1 1
OpEntryPoint Fragment %389 "foo_frag" %388
OpEntryPoint GLCompute %407 "foo_compute"
OpExecutionMode %389 OriginUpperLeft
OpExecutionMode %407 LocalSize 1 1 1
%3 = OpString "access.wgsl"
OpSource Unknown 0 %3 "// This snapshot tests accessing various containers, dereferencing pointers.
@ -20,26 +20,26 @@ struct GlobalConst {
c: i32,
}
// tests msl padding insertion for global constants
var<private> global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0);
var<private> msl_padding_global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0);
struct AlignedWrapper {
\t@align(8) value: i32
@align(8) value: i32
}
struct Bar {
\t_matrix: mat4x3<f32>,
\tmatrix_array: array<mat2x2<f32>, 2>,
\tatom: atomic<i32>,
\tatom_arr: array<atomic<i32>, 10>,
\tarr: array<vec2<u32>, 2>,
\tdata: array<AlignedWrapper>,
_matrix: mat4x3<f32>,
matrix_array: array<mat2x2<f32>, 2>,
atom: atomic<i32>,
atom_arr: array<atomic<i32>, 10>,
arr: array<vec2<u32>, 2>,
data: array<AlignedWrapper>,
}
@group(0) @binding(0)
var<storage,read_write> bar: Bar;
struct Baz {
\tm: mat3x2<f32>,
m: mat3x2<f32>,
}
@group(0) @binding(1)
@ -49,11 +49,11 @@ var<uniform> baz: Baz;
var<storage,read_write> qux: vec2<i32>;
fn test_matrix_within_struct_accesses() {
\tvar idx = 1;
var idx = 1;
idx--;
\t// loads
// loads
let l0 = baz.m;
let l1 = baz.m[0];
let l2 = baz.m[idx];
@ -66,7 +66,7 @@ fn test_matrix_within_struct_accesses() {
idx++;
\t// stores
// stores
t.m = mat3x2<f32>(vec2<f32>(6.0), vec2<f32>(5.0), vec2<f32>(4.0));
t.m[0] = vec2<f32>(9.0);
t.m[idx] = vec2<f32>(90.0);
@ -77,18 +77,18 @@ fn test_matrix_within_struct_accesses() {
}
struct MatCx2InArray {
\tam: array<mat4x2<f32>, 2>,
am: array<mat4x2<f32>, 2>,
}
@group(0) @binding(3)
var<uniform> nested_mat_cx2: MatCx2InArray;
fn test_matrix_within_array_within_struct_accesses() {
\tvar idx = 1;
var idx = 1;
idx--;
\t// loads
// loads
let l0 = nested_mat_cx2.am;
let l1 = nested_mat_cx2.am[0];
let l2 = nested_mat_cx2.am[0][0];
@ -102,7 +102,7 @@ fn test_matrix_within_array_within_struct_accesses() {
idx++;
\t// stores
// stores
t.am = array<mat4x2<f32>, 2>();
t.am[0] = mat4x2<f32>(vec2<f32>(8.0), vec2<f32>(7.0), vec2<f32>(6.0), vec2<f32>(5.0));
t.am[0][0] = vec2<f32>(9.0);
@ -128,41 +128,42 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let baz: f32 = foo;
foo = 1.0;
\ttest_matrix_within_struct_accesses();
\ttest_matrix_within_array_within_struct_accesses();
_ = msl_padding_global_const;
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
// test storage loads
\tlet _matrix = bar._matrix;
\tlet arr = bar.arr;
\tlet index = 3u;
\tlet b = bar._matrix[index].x;
\tlet a = bar.data[arrayLength(&bar.data) - 2u].value;
\tlet c = qux;
let _matrix = bar._matrix;
let arr = bar.arr;
let index = 3u;
let b = bar._matrix[index].x;
let a = bar.data[arrayLength(&bar.data) - 2u].value;
let c = qux;
\t// test pointer types
\tlet data_pointer: ptr<storage, i32, read_write> = &bar.data[0].value;
\tlet foo_value = read_from_private(&foo);
// test pointer types
let data_pointer: ptr<storage, i32, read_write> = &bar.data[0].value;
let foo_value = read_from_private(&foo);
\t// test array indexing
\tvar c2 = array<i32, 5>(a, i32(b), 3, 4, 5);
\tc2[vi + 1u] = 42;
\tlet value = c2[vi];
// test array indexing
var c2 = array<i32, 5>(a, i32(b), 3, 4, 5);
c2[vi + 1u] = 42;
let value = c2[vi];
\ttest_arr_as_arg(array<array<f32, 10>, 5>());
test_arr_as_arg(array<array<f32, 10>, 5>());
\treturn vec4<f32>(_matrix * vec4<f32>(vec4<i32>(value)), 2.0);
return vec4<f32>(_matrix * vec4<f32>(vec4<i32>(value)), 2.0);
}
@fragment
fn foo_frag() -> @location(0) vec4<f32> {
\t// test storage stores
\tbar._matrix[1].z = 1.0;
\tbar._matrix = mat4x3<f32>(vec3<f32>(0.0), vec3<f32>(1.0), vec3<f32>(2.0), vec3<f32>(3.0));
\tbar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
\tbar.data[1].value = 1;
\tqux = vec2<i32>();
// test storage stores
bar._matrix[1].z = 1.0;
bar._matrix = mat4x3<f32>(vec3<f32>(0.0), vec3<f32>(1.0), vec3<f32>(2.0), vec3<f32>(3.0));
bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
bar.data[1].value = 1;
qux = vec2<i32>();
\treturn vec4<f32>(0.0);
return vec4<f32>(0.0);
}
fn assign_through_ptr_fn(p: ptr<function, u32>) {
@ -177,7 +178,7 @@ fn assign_through_ptr() {
var val = 33u;
assign_through_ptr_fn(&val);
\tvar arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
assign_array_through_ptr_fn(&arr);
}
@ -291,7 +292,7 @@ OpName %46 "Inner"
OpMemberName %47 0 "om_nom_nom"
OpMemberName %47 1 "thing"
OpName %47 "Outer"
OpName %52 "global_const"
OpName %52 "msl_padding_global_const"
OpName %54 "bar"
OpName %56 "baz"
OpName %59 "qux"
@ -338,8 +339,8 @@ OpName %328 "vi"
OpName %333 "foo_vert"
OpName %344 "foo"
OpName %345 "c2"
OpName %388 "foo_frag"
OpName %406 "foo_compute"
OpName %389 "foo_frag"
OpName %407 "foo_compute"
OpMemberDecorate %7 0 Offset 0
OpMemberDecorate %7 1 Offset 16
OpMemberDecorate %7 2 Offset 28
@ -393,7 +394,7 @@ OpDecorate %63 Block
OpMemberDecorate %63 0 Offset 0
OpDecorate %328 BuiltIn VertexIndex
OpDecorate %331 BuiltIn Position
OpDecorate %387 Location 0
OpDecorate %388 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 0
%5 = OpTypeVector %4 3
@ -546,28 +547,28 @@ OpDecorate %387 Location 0
%343 = OpConstantNull %30
%346 = OpTypePointer Function %33
%347 = OpConstantNull %33
%352 = OpTypePointer StorageBuffer %10
%355 = OpTypePointer StorageBuffer %19
%358 = OpTypePointer StorageBuffer %11
%359 = OpTypePointer StorageBuffer %9
%362 = OpTypePointer StorageBuffer %20
%365 = OpTypePointer StorageBuffer %8
%366 = OpTypePointer StorageBuffer %6
%371 = OpConstant %9 -2147483600.0
%372 = OpConstant %9 2147483500.0
%381 = OpTypeVector %6 4
%387 = OpVariable %332 Output
%390 = OpConstantComposite %11 %338 %338 %338
%391 = OpConstantComposite %11 %71 %71 %71
%392 = OpConstantComposite %11 %73 %73 %73
%393 = OpConstantComposite %11 %75 %75 %75
%394 = OpConstantComposite %10 %390 %391 %392 %393
%395 = OpConstantComposite %18 %48 %48
%396 = OpConstantComposite %18 %44 %44
%397 = OpConstantComposite %19 %395 %396
%398 = OpConstantNull %24
%399 = OpConstantComposite %32 %338 %338 %338 %338
%407 = OpConstantTrue %42
%353 = OpTypePointer StorageBuffer %10
%356 = OpTypePointer StorageBuffer %19
%359 = OpTypePointer StorageBuffer %11
%360 = OpTypePointer StorageBuffer %9
%363 = OpTypePointer StorageBuffer %20
%366 = OpTypePointer StorageBuffer %8
%367 = OpTypePointer StorageBuffer %6
%372 = OpConstant %9 -2147483600.0
%373 = OpConstant %9 2147483500.0
%382 = OpTypeVector %6 4
%388 = OpVariable %332 Output
%391 = OpConstantComposite %11 %338 %338 %338
%392 = OpConstantComposite %11 %71 %71 %71
%393 = OpConstantComposite %11 %73 %73 %73
%394 = OpConstantComposite %11 %75 %75 %75
%395 = OpConstantComposite %10 %391 %392 %393 %394
%396 = OpConstantComposite %18 %48 %48
%397 = OpConstantComposite %18 %44 %44
%398 = OpConstantComposite %19 %396 %397
%399 = OpConstantNull %24
%400 = OpConstantComposite %32 %338 %338 %338 %338
%408 = OpConstantTrue %42
%66 = OpFunction %2 None %67
%65 = OpLabel
%94 = OpVariable %95 Function %70
@ -805,7 +806,7 @@ OpFunctionEnd
%209 = OpLabel
OpBranch %214
%214 = OpLabel
OpLine %3 155 5
OpLine %3 156 5
OpStore %210 %213
OpReturn
OpFunctionEnd
@ -814,11 +815,11 @@ OpFunctionEnd
%215 = OpLabel
OpBranch %222
%222 = OpLabel
OpLine %3 159 32
OpLine %3 159 43
OpLine %3 159 32
OpLine %3 159 12
OpLine %3 159 5
OpLine %3 160 32
OpLine %3 160 43
OpLine %3 160 32
OpLine %3 160 12
OpLine %3 160 5
OpStore %216 %221
OpReturn
OpFunctionEnd
@ -828,13 +829,13 @@ OpFunctionEnd
%230 = OpVariable %36 Function %228
OpBranch %231
%231 = OpLabel
OpLine %3 164 5
OpLine %3 165 5
%232 = OpFunctionCall %2 %211 %229
OpLine %3 166 32
OpLine %3 166 43
OpLine %3 166 32
OpLine %3 166 12
OpLine %3 167 5
OpLine %3 167 35
OpLine %3 167 46
OpLine %3 167 35
OpLine %3 167 15
OpLine %3 168 5
%233 = OpFunctionCall %2 %217 %230
OpReturn
OpFunctionEnd
@ -843,7 +844,7 @@ OpFunctionEnd
%234 = OpLabel
OpBranch %238
%238 = OpLabel
OpLine %3 175 10
OpLine %3 176 10
%239 = OpAccessChain %34 %235 %48
%240 = OpLoad %4 %239
OpReturnValue %240
@ -853,8 +854,8 @@ OpFunctionEnd
%241 = OpLabel
OpBranch %245
%245 = OpLabel
OpLine %3 179 3
OpLine %3 179 3
OpLine %3 180 3
OpLine %3 180 3
%246 = OpAccessChain %34 %242 %48
OpStore %246 %17
OpReturn
@ -864,7 +865,7 @@ OpFunctionEnd
%247 = OpLabel
OpBranch %251
%251 = OpLabel
OpLine %3 183 10
OpLine %3 184 10
%252 = OpAccessChain %34 %248 %44
%253 = OpLoad %4 %252
OpReturnValue %253
@ -874,8 +875,8 @@ OpFunctionEnd
%254 = OpLabel
OpBranch %258
%258 = OpLabel
OpLine %3 187 3
OpLine %3 187 3
OpLine %3 188 3
OpLine %3 188 3
%259 = OpAccessChain %34 %255 %44
OpStore %259 %17
OpReturn
@ -886,13 +887,13 @@ OpFunctionEnd
%264 = OpVariable %41 Function %265
OpBranch %266
%266 = OpLabel
OpLine %3 192 4
%267 = OpFunctionCall %2 %243 %262
OpLine %3 193 4
%267 = OpFunctionCall %2 %243 %262
OpLine %3 194 4
%268 = OpFunctionCall %4 %236 %262
OpLine %3 196 4
%269 = OpFunctionCall %2 %256 %264
OpLine %3 197 4
%269 = OpFunctionCall %2 %256 %264
OpLine %3 198 4
%270 = OpFunctionCall %4 %249 %264
OpReturn
OpFunctionEnd
@ -902,11 +903,11 @@ OpFunctionEnd
%275 = OpVariable %276 Function %277
OpBranch %278
%278 = OpLabel
OpLine %3 201 13
OpLine %3 202 13
%279 = OpCompositeConstruct %43 %272
OpLine %3 201 5
OpLine %3 202 5
OpStore %275 %279
OpLine %3 203 12
OpLine %3 204 12
%281 = OpAccessChain %280 %275 %48
%282 = OpLoad %42 %281
OpReturnValue %282
@ -916,8 +917,8 @@ OpFunctionEnd
%288 = OpVariable %289 Function %287
OpBranch %290
%290 = OpLabel
OpLine %3 209 16
OpLine %3 211 12
OpLine %3 210 16
OpLine %3 212 12
%291 = OpAccessChain %95 %288 %48
%292 = OpLoad %6 %291
OpReturnValue %292
@ -926,19 +927,19 @@ OpFunctionEnd
%293 = OpLabel
OpBranch %296
%296 = OpLabel
OpLine %3 221 17
OpLine %3 222 17
%297 = OpCompositeExtract %46 %295 0
OpLine %3 222 20
OpLine %3 223 20
%298 = OpCompositeExtract %6 %297 0
OpLine %3 224 9
OpLine %3 225 9
%299 = OpCompositeExtract %4 %295 1
%300 = OpBitcast %4 %298
%301 = OpINotEqual %42 %299 %300
OpLine %3 224 5
OpLine %3 225 5
OpSelectionMerge %302 None
OpBranchConditional %301 %302 %302
%302 = OpLabel
OpLine %3 228 12
OpLine %3 229 12
%303 = OpCompositeExtract %46 %295 0
%304 = OpCompositeExtract %6 %303 0
OpReturnValue %304
@ -950,27 +951,27 @@ OpFunctionEnd
%312 = OpVariable %95 Function %313
OpBranch %314
%314 = OpLabel
OpLine %3 234 17
OpLine %3 235 17
%315 = OpAccessChain %310 %307 %48
%316 = OpLoad %46 %315
OpLine %3 234 5
OpLine %3 235 5
OpStore %309 %316
OpLine %3 235 20
OpLine %3 236 20
%317 = OpAccessChain %95 %309 %48
%318 = OpLoad %6 %317
OpLine %3 235 5
OpLine %3 236 5
OpStore %312 %318
OpLine %3 237 9
OpLine %3 238 9
%319 = OpAccessChain %34 %307 %44
%320 = OpLoad %4 %319
%321 = OpLoad %6 %312
%322 = OpBitcast %4 %321
%323 = OpINotEqual %42 %320 %322
OpLine %3 237 5
OpLine %3 238 5
OpSelectionMerge %324 None
OpBranchConditional %323 %324 %324
%324 = OpLabel
OpLine %3 241 12
OpLine %3 242 12
%325 = OpAccessChain %95 %307 %48 %48
%326 = OpLoad %6 %325
OpReturnValue %326
@ -989,108 +990,110 @@ OpLine %3 1 1
%349 = OpLoad %9 %344
OpLine %3 115 5
OpStore %344 %71
OpLine %3 117 2
%350 = OpFunctionCall %2 %66
OpLine %3 118 2
%351 = OpFunctionCall %2 %140
OpLine %3 121 16
%353 = OpAccessChain %352 %54 %48
%354 = OpLoad %10 %353
OpLine %3 122 12
%356 = OpAccessChain %355 %54 %40
%357 = OpLoad %19 %356
OpLine %3 124 10
%360 = OpAccessChain %359 %54 %48 %339 %48
%361 = OpLoad %9 %360
OpLine %3 125 10
OpLine %3 125 19
%363 = OpArrayLength %4 %54 5
OpLine %3 125 10
%364 = OpISub %4 %363 %15
%367 = OpAccessChain %366 %54 %31 %364 %48
%368 = OpLoad %6 %367
OpLine %3 126 10
%369 = OpLoad %24 %336
OpLine %3 129 53
OpLine %3 129 53
OpLine %3 130 18
%370 = OpFunctionCall %9 %198 %344
OpLine %3 133 28
%373 = OpExtInst %9 %1 FClamp %361 %371 %372
%374 = OpConvertFToS %6 %373
OpLine %3 133 11
%375 = OpCompositeConstruct %33 %368 %374 %340 %341 %342
OpLine %3 133 2
OpStore %345 %375
OpLine %3 134 2
%376 = OpIAdd %4 %330 %44
OpLine %3 134 2
%377 = OpAccessChain %95 %345 %376
OpStore %377 %286
OpLine %3 135 14
%378 = OpAccessChain %95 %345 %330
%379 = OpLoad %6 %378
OpLine %3 137 2
%380 = OpFunctionCall %9 %204 %343
OpLine %3 139 19
%382 = OpCompositeConstruct %381 %379 %379 %379 %379
%383 = OpConvertSToF %32 %382
%384 = OpMatrixTimesVector %11 %354 %383
OpLine %3 139 9
%385 = OpCompositeConstruct %32 %384 %73
OpStore %331 %385
OpLine %3 117 9
%350 = OpLoad %7 %52
OpLine %3 118 5
%351 = OpFunctionCall %2 %66
OpLine %3 119 5
%352 = OpFunctionCall %2 %140
OpLine %3 122 19
%354 = OpAccessChain %353 %54 %48
%355 = OpLoad %10 %354
OpLine %3 123 15
%357 = OpAccessChain %356 %54 %40
%358 = OpLoad %19 %357
OpLine %3 125 13
%361 = OpAccessChain %360 %54 %48 %339 %48
%362 = OpLoad %9 %361
OpLine %3 126 13
OpLine %3 126 22
%364 = OpArrayLength %4 %54 5
OpLine %3 126 13
%365 = OpISub %4 %364 %15
%368 = OpAccessChain %367 %54 %31 %365 %48
%369 = OpLoad %6 %368
OpLine %3 127 13
%370 = OpLoad %24 %336
OpLine %3 130 56
OpLine %3 130 56
OpLine %3 131 21
%371 = OpFunctionCall %9 %198 %344
OpLine %3 134 31
%374 = OpExtInst %9 %1 FClamp %362 %372 %373
%375 = OpConvertFToS %6 %374
OpLine %3 134 14
%376 = OpCompositeConstruct %33 %369 %375 %340 %341 %342
OpLine %3 134 5
OpStore %345 %376
OpLine %3 135 5
%377 = OpIAdd %4 %330 %44
OpLine %3 135 5
%378 = OpAccessChain %95 %345 %377
OpStore %378 %286
OpLine %3 136 17
%379 = OpAccessChain %95 %345 %330
%380 = OpLoad %6 %379
OpLine %3 138 5
%381 = OpFunctionCall %9 %204 %343
OpLine %3 140 22
%383 = OpCompositeConstruct %382 %380 %380 %380 %380
%384 = OpConvertSToF %32 %383
%385 = OpMatrixTimesVector %11 %355 %384
OpLine %3 140 12
%386 = OpCompositeConstruct %32 %385 %73
OpStore %331 %386
OpReturn
OpFunctionEnd
%388 = OpFunction %2 None %67
%386 = OpLabel
%389 = OpAccessChain %335 %59 %48
OpBranch %400
%400 = OpLabel
OpLine %3 145 2
OpLine %3 145 2
OpLine %3 145 2
%401 = OpAccessChain %359 %54 %48 %44 %15
OpStore %401 %71
OpLine %3 146 2
OpLine %3 146 28
OpLine %3 146 44
OpLine %3 146 60
OpLine %3 146 16
OpLine %3 146 2
%402 = OpAccessChain %352 %54 %48
OpStore %402 %394
OpLine %3 147 2
OpLine %3 147 32
OpLine %3 147 12
OpLine %3 147 2
%403 = OpAccessChain %355 %54 %40
OpStore %403 %397
OpLine %3 148 2
OpLine %3 148 2
OpLine %3 148 2
%404 = OpAccessChain %366 %54 %31 %44 %48
OpStore %404 %70
OpLine %3 149 2
OpStore %389 %398
OpLine %3 151 9
OpStore %387 %399
%389 = OpFunction %2 None %67
%387 = OpLabel
%390 = OpAccessChain %335 %59 %48
OpBranch %401
%401 = OpLabel
OpLine %3 146 5
OpLine %3 146 5
OpLine %3 146 5
%402 = OpAccessChain %360 %54 %48 %44 %15
OpStore %402 %71
OpLine %3 147 5
OpLine %3 147 31
OpLine %3 147 47
OpLine %3 147 63
OpLine %3 147 19
OpLine %3 147 5
%403 = OpAccessChain %353 %54 %48
OpStore %403 %395
OpLine %3 148 5
OpLine %3 148 35
OpLine %3 148 15
OpLine %3 148 5
%404 = OpAccessChain %356 %54 %40
OpStore %404 %398
OpLine %3 149 5
OpLine %3 149 5
OpLine %3 149 5
%405 = OpAccessChain %367 %54 %31 %44 %48
OpStore %405 %70
OpLine %3 150 5
OpStore %390 %399
OpLine %3 152 12
OpStore %388 %400
OpReturn
OpFunctionEnd
%406 = OpFunction %2 None %67
%405 = OpLabel
OpBranch %408
%408 = OpLabel
OpLine %3 246 5
%409 = OpFunctionCall %2 %224
%407 = OpFunction %2 None %67
%406 = OpLabel
OpBranch %409
%409 = OpLabel
OpLine %3 247 5
%410 = OpFunctionCall %2 %261
%410 = OpFunctionCall %2 %224
OpLine %3 248 5
%411 = OpFunctionCall %42 %273 %407
%411 = OpFunctionCall %2 %261
OpLine %3 249 5
%412 = OpFunctionCall %6 %284
%412 = OpFunctionCall %42 %273 %408
OpLine %3 250 5
%413 = OpFunctionCall %6 %294
%413 = OpFunctionCall %6 %284
OpLine %3 251 5
%414 = OpFunctionCall %6 %306
%414 = OpFunctionCall %6 %294
OpLine %3 252 5
%415 = OpFunctionCall %6 %306
OpReturn
OpFunctionEnd

183
naga/tests/out/spv/wgsl-debug-symbol-large-source.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 678
; Bound: 682
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
@ -7457,6 +7457,11 @@ fn color23(p: vec2<f32>) -> vec3<f32> {
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
_ = t_diffuse;
_ = s_diffuse;
_ = t_normal;
_ = s_normal;
color23(vec2(1, 2));
var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));
@ -7582,7 +7587,7 @@ OpName %608 "clip_position"
OpName %610 "normal"
OpName %612 "world_pos"
OpName %615 "fs_main"
OpName %625 "color"
OpName %629 "color"
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 8
OpMemberDecorate %13 2 Offset 16
@ -7826,15 +7831,15 @@ OpDecorate %614 Location 0
%612 = OpVariable %582 Input
%614 = OpVariable %441 Output
%617 = OpTypePointer Uniform %25
%619 = OpConstantComposite %6 %56 %81
%620 = OpConstantComposite %5 %286 %286 %286
%621 = OpConstant %4 0.7
%622 = OpConstantComposite %5 %78 %286 %621
%623 = OpConstant %4 0.2
%624 = OpConstantComposite %5 %623 %623 %623
%626 = OpConstantNull %5
%642 = OpTypePointer Uniform %5
%651 = OpTypePointer Uniform %7
%623 = OpConstantComposite %6 %56 %81
%624 = OpConstantComposite %5 %286 %286 %286
%625 = OpConstant %4 0.7
%626 = OpConstantComposite %5 %78 %286 %625
%627 = OpConstant %4 0.2
%628 = OpConstantComposite %5 %627 %627 %627
%630 = OpConstantNull %5
%646 = OpTypePointer Uniform %5
%655 = OpTypePointer Uniform %7
%53 = OpFunction %5 None %54
%52 = OpFunctionParameter %5
%51 = OpLabel
@ -8594,87 +8599,91 @@ OpReturn
OpFunctionEnd
%615 = OpFunction %2 None %368
%606 = OpLabel
%625 = OpVariable %95 Function %626
%629 = OpVariable %95 Function %630
%609 = OpLoad %7 %608
%611 = OpLoad %5 %610
%613 = OpLoad %5 %612
%607 = OpCompositeConstruct %26 %609 %611 %613
%616 = OpAccessChain %591 %39 %135
%618 = OpAccessChain %617 %42 %135
OpBranch %627
%627 = OpLabel
OpLine %3 7421 13
OpLine %3 7421 13
OpLine %3 7421 5
%628 = OpFunctionCall %5 %342 %619
OpLine %3 7423 28
OpLine %3 7423 17
%629 = OpCompositeExtract %5 %607 2
%630 = OpExtInst %5 %1 Fract %629
%631 = OpExtInst %5 %1 SmoothStep %80 %620 %630
OpLine %3 7423 5
OpStore %625 %631
OpLine %3 7424 17
OpLine %3 7424 13
%632 = OpAccessChain %125 %625 %135
%633 = OpLoad %4 %632
%634 = OpAccessChain %125 %625 %126
%635 = OpLoad %4 %634
%636 = OpFMul %4 %633 %635
%637 = OpAccessChain %125 %625 %372
%638 = OpLoad %4 %637
%639 = OpFMul %4 %636 %638
%640 = OpCompositeConstruct %5 %639 %639 %639
%641 = OpExtInst %5 %1 FMix %622 %624 %640
OpLine %3 7424 5
OpStore %625 %641
OpLine %3 7427 25
%643 = OpAccessChain %642 %618 %126
%644 = OpLoad %5 %643
%645 = OpVectorTimesScalar %5 %644 %286
OpLine %3 7429 21
%646 = OpAccessChain %642 %618 %135
%647 = OpLoad %5 %646
%648 = OpCompositeExtract %5 %607 2
%649 = OpFSub %5 %647 %648
%650 = OpExtInst %5 %1 Normalize %649
OpLine %3 7430 20
%652 = OpAccessChain %651 %616 %135
%653 = OpLoad %7 %652
%654 = OpVectorShuffle %5 %653 %653 0 1 2
%655 = OpCompositeExtract %5 %607 2
%656 = OpFSub %5 %654 %655
%657 = OpExtInst %5 %1 Normalize %656
OpLine %3 7431 20
%658 = OpFAdd %5 %657 %650
%659 = OpExtInst %5 %1 Normalize %658
OpLine %3 7433 32
%660 = OpCompositeExtract %5 %607 1
%661 = OpDot %4 %660 %650
OpLine %3 7433 28
%662 = OpExtInst %4 %1 FMax %661 %74
OpLine %3 7434 25
%663 = OpAccessChain %642 %618 %126
%664 = OpLoad %5 %663
%665 = OpVectorTimesScalar %5 %664 %662
OpLine %3 7436 37
%666 = OpCompositeExtract %5 %607 1
%667 = OpDot %4 %666 %659
OpLine %3 7436 33
%668 = OpExtInst %4 %1 FMax %667 %74
OpLine %3 7436 29
%669 = OpExtInst %4 %1 Pow %668 %345
OpLine %3 7437 26
%670 = OpAccessChain %642 %618 %126
%671 = OpLoad %5 %670
%672 = OpVectorTimesScalar %5 %671 %669
OpLine %3 7439 18
%673 = OpFAdd %5 %645 %665
%674 = OpFAdd %5 %673 %672
%675 = OpLoad %5 %625
%676 = OpFMul %5 %674 %675
OpLine %3 7441 12
%677 = OpCompositeConstruct %7 %676 %56
OpStore %614 %677
%619 = OpLoad %27 %45
%620 = OpLoad %28 %47
%621 = OpLoad %27 %49
%622 = OpLoad %28 %50
OpBranch %631
%631 = OpLabel
OpLine %3 7426 13
OpLine %3 7426 13
OpLine %3 7426 5
%632 = OpFunctionCall %5 %342 %623
OpLine %3 7428 28
OpLine %3 7428 17
%633 = OpCompositeExtract %5 %607 2
%634 = OpExtInst %5 %1 Fract %633
%635 = OpExtInst %5 %1 SmoothStep %80 %624 %634
OpLine %3 7428 5
OpStore %629 %635
OpLine %3 7429 17
OpLine %3 7429 13
%636 = OpAccessChain %125 %629 %135
%637 = OpLoad %4 %636
%638 = OpAccessChain %125 %629 %126
%639 = OpLoad %4 %638
%640 = OpFMul %4 %637 %639
%641 = OpAccessChain %125 %629 %372
%642 = OpLoad %4 %641
%643 = OpFMul %4 %640 %642
%644 = OpCompositeConstruct %5 %643 %643 %643
%645 = OpExtInst %5 %1 FMix %626 %628 %644
OpLine %3 7429 5
OpStore %629 %645
OpLine %3 7432 25
%647 = OpAccessChain %646 %618 %126
%648 = OpLoad %5 %647
%649 = OpVectorTimesScalar %5 %648 %286
OpLine %3 7434 21
%650 = OpAccessChain %646 %618 %135
%651 = OpLoad %5 %650
%652 = OpCompositeExtract %5 %607 2
%653 = OpFSub %5 %651 %652
%654 = OpExtInst %5 %1 Normalize %653
OpLine %3 7435 20
%656 = OpAccessChain %655 %616 %135
%657 = OpLoad %7 %656
%658 = OpVectorShuffle %5 %657 %657 0 1 2
%659 = OpCompositeExtract %5 %607 2
%660 = OpFSub %5 %658 %659
%661 = OpExtInst %5 %1 Normalize %660
OpLine %3 7436 20
%662 = OpFAdd %5 %661 %654
%663 = OpExtInst %5 %1 Normalize %662
OpLine %3 7438 32
%664 = OpCompositeExtract %5 %607 1
%665 = OpDot %4 %664 %654
OpLine %3 7438 28
%666 = OpExtInst %4 %1 FMax %665 %74
OpLine %3 7439 25
%667 = OpAccessChain %646 %618 %126
%668 = OpLoad %5 %667
%669 = OpVectorTimesScalar %5 %668 %666
OpLine %3 7441 37
%670 = OpCompositeExtract %5 %607 1
%671 = OpDot %4 %670 %663
OpLine %3 7441 33
%672 = OpExtInst %4 %1 FMax %671 %74
OpLine %3 7441 29
%673 = OpExtInst %4 %1 Pow %672 %345
OpLine %3 7442 26
%674 = OpAccessChain %646 %618 %126
%675 = OpLoad %5 %674
%676 = OpVectorTimesScalar %5 %675 %673
OpLine %3 7444 18
%677 = OpFAdd %5 %649 %669
%678 = OpFAdd %5 %677 %676
%679 = OpLoad %5 %629
%680 = OpFMul %5 %678 %679
OpLine %3 7446 12
%681 = OpCompositeConstruct %7 %680 %56
OpStore %614 %681
OpReturn
OpFunctionEnd

185
naga/tests/out/spv/wgsl-debug-symbol-terrain.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 678
; Bound: 682
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
@ -203,7 +203,7 @@ fn gen_terrain_fragment(in: GenVertexOutput) -> GenFragmentOutput {
let v = terrain_vertex(p, gen_data.min_max_height);
var vert_component: f32 = 0.;
switch comp_index {
case 0u: { vert_component = v.position.x; }
case 1u: { vert_component = v.position.y; }
@ -292,6 +292,11 @@ fn color23(p: vec2<f32>) -> vec3<f32> {
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
_ = t_diffuse;
_ = s_diffuse;
_ = t_normal;
_ = s_normal;
color23(vec2(1, 2));
var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));
@ -417,7 +422,7 @@ OpName %608 "clip_position"
OpName %610 "normal"
OpName %612 "world_pos"
OpName %615 "fs_main"
OpName %625 "color"
OpName %629 "color"
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 8
OpMemberDecorate %13 2 Offset 16
@ -661,15 +666,15 @@ OpDecorate %614 Location 0
%612 = OpVariable %582 Input
%614 = OpVariable %441 Output
%617 = OpTypePointer Uniform %25
%619 = OpConstantComposite %6 %56 %81
%620 = OpConstantComposite %5 %286 %286 %286
%621 = OpConstant %4 0.7
%622 = OpConstantComposite %5 %78 %286 %621
%623 = OpConstant %4 0.2
%624 = OpConstantComposite %5 %623 %623 %623
%626 = OpConstantNull %5
%642 = OpTypePointer Uniform %5
%651 = OpTypePointer Uniform %7
%623 = OpConstantComposite %6 %56 %81
%624 = OpConstantComposite %5 %286 %286 %286
%625 = OpConstant %4 0.7
%626 = OpConstantComposite %5 %78 %286 %625
%627 = OpConstant %4 0.2
%628 = OpConstantComposite %5 %627 %627 %627
%630 = OpConstantNull %5
%646 = OpTypePointer Uniform %5
%655 = OpTypePointer Uniform %7
%53 = OpFunction %5 None %54
%52 = OpFunctionParameter %5
%51 = OpLabel
@ -1429,87 +1434,91 @@ OpReturn
OpFunctionEnd
%615 = OpFunction %2 None %368
%606 = OpLabel
%625 = OpVariable %95 Function %626
%629 = OpVariable %95 Function %630
%609 = OpLoad %7 %608
%611 = OpLoad %5 %610
%613 = OpLoad %5 %612
%607 = OpCompositeConstruct %26 %609 %611 %613
%616 = OpAccessChain %591 %39 %135
%618 = OpAccessChain %617 %42 %135
OpBranch %627
%627 = OpLabel
OpLine %3 278 13
OpLine %3 278 13
OpLine %3 278 5
%628 = OpFunctionCall %5 %342 %619
OpLine %3 280 28
OpLine %3 280 17
%629 = OpCompositeExtract %5 %607 2
%630 = OpExtInst %5 %1 Fract %629
%631 = OpExtInst %5 %1 SmoothStep %80 %620 %630
OpLine %3 280 5
OpStore %625 %631
OpLine %3 281 17
OpLine %3 281 13
%632 = OpAccessChain %125 %625 %135
%633 = OpLoad %4 %632
%634 = OpAccessChain %125 %625 %126
%635 = OpLoad %4 %634
%636 = OpFMul %4 %633 %635
%637 = OpAccessChain %125 %625 %372
%638 = OpLoad %4 %637
%639 = OpFMul %4 %636 %638
%640 = OpCompositeConstruct %5 %639 %639 %639
%641 = OpExtInst %5 %1 FMix %622 %624 %640
OpLine %3 281 5
OpStore %625 %641
OpLine %3 284 25
%643 = OpAccessChain %642 %618 %126
%644 = OpLoad %5 %643
%645 = OpVectorTimesScalar %5 %644 %286
OpLine %3 286 21
%646 = OpAccessChain %642 %618 %135
%647 = OpLoad %5 %646
%648 = OpCompositeExtract %5 %607 2
%649 = OpFSub %5 %647 %648
%650 = OpExtInst %5 %1 Normalize %649
OpLine %3 287 20
%652 = OpAccessChain %651 %616 %135
%653 = OpLoad %7 %652
%654 = OpVectorShuffle %5 %653 %653 0 1 2
%655 = OpCompositeExtract %5 %607 2
%656 = OpFSub %5 %654 %655
%657 = OpExtInst %5 %1 Normalize %656
OpLine %3 288 20
%658 = OpFAdd %5 %657 %650
%659 = OpExtInst %5 %1 Normalize %658
OpLine %3 290 32
%660 = OpCompositeExtract %5 %607 1
%661 = OpDot %4 %660 %650
OpLine %3 290 28
%662 = OpExtInst %4 %1 FMax %661 %74
OpLine %3 291 25
%663 = OpAccessChain %642 %618 %126
%664 = OpLoad %5 %663
%665 = OpVectorTimesScalar %5 %664 %662
OpLine %3 293 37
%666 = OpCompositeExtract %5 %607 1
%667 = OpDot %4 %666 %659
OpLine %3 293 33
%668 = OpExtInst %4 %1 FMax %667 %74
OpLine %3 293 29
%669 = OpExtInst %4 %1 Pow %668 %345
OpLine %3 294 26
%670 = OpAccessChain %642 %618 %126
%671 = OpLoad %5 %670
%672 = OpVectorTimesScalar %5 %671 %669
OpLine %3 296 18
%673 = OpFAdd %5 %645 %665
%674 = OpFAdd %5 %673 %672
%675 = OpLoad %5 %625
%676 = OpFMul %5 %674 %675
OpLine %3 298 12
%677 = OpCompositeConstruct %7 %676 %56
OpStore %614 %677
%619 = OpLoad %27 %45
%620 = OpLoad %28 %47
%621 = OpLoad %27 %49
%622 = OpLoad %28 %50
OpBranch %631
%631 = OpLabel
OpLine %3 283 13
OpLine %3 283 13
OpLine %3 283 5
%632 = OpFunctionCall %5 %342 %623
OpLine %3 285 28
OpLine %3 285 17
%633 = OpCompositeExtract %5 %607 2
%634 = OpExtInst %5 %1 Fract %633
%635 = OpExtInst %5 %1 SmoothStep %80 %624 %634
OpLine %3 285 5
OpStore %629 %635
OpLine %3 286 17
OpLine %3 286 13
%636 = OpAccessChain %125 %629 %135
%637 = OpLoad %4 %636
%638 = OpAccessChain %125 %629 %126
%639 = OpLoad %4 %638
%640 = OpFMul %4 %637 %639
%641 = OpAccessChain %125 %629 %372
%642 = OpLoad %4 %641
%643 = OpFMul %4 %640 %642
%644 = OpCompositeConstruct %5 %643 %643 %643
%645 = OpExtInst %5 %1 FMix %626 %628 %644
OpLine %3 286 5
OpStore %629 %645
OpLine %3 289 25
%647 = OpAccessChain %646 %618 %126
%648 = OpLoad %5 %647
%649 = OpVectorTimesScalar %5 %648 %286
OpLine %3 291 21
%650 = OpAccessChain %646 %618 %135
%651 = OpLoad %5 %650
%652 = OpCompositeExtract %5 %607 2
%653 = OpFSub %5 %651 %652
%654 = OpExtInst %5 %1 Normalize %653
OpLine %3 292 20
%656 = OpAccessChain %655 %616 %135
%657 = OpLoad %7 %656
%658 = OpVectorShuffle %5 %657 %657 0 1 2
%659 = OpCompositeExtract %5 %607 2
%660 = OpFSub %5 %658 %659
%661 = OpExtInst %5 %1 Normalize %660
OpLine %3 293 20
%662 = OpFAdd %5 %661 %654
%663 = OpExtInst %5 %1 Normalize %662
OpLine %3 295 32
%664 = OpCompositeExtract %5 %607 1
%665 = OpDot %4 %664 %654
OpLine %3 295 28
%666 = OpExtInst %4 %1 FMax %665 %74
OpLine %3 296 25
%667 = OpAccessChain %646 %618 %126
%668 = OpLoad %5 %667
%669 = OpVectorTimesScalar %5 %668 %666
OpLine %3 298 37
%670 = OpCompositeExtract %5 %607 1
%671 = OpDot %4 %670 %663
OpLine %3 298 33
%672 = OpExtInst %4 %1 FMax %671 %74
OpLine %3 298 29
%673 = OpExtInst %4 %1 Pow %672 %345
OpLine %3 299 26
%674 = OpAccessChain %646 %618 %126
%675 = OpLoad %5 %674
%676 = OpVectorTimesScalar %5 %675 %673
OpLine %3 301 18
%677 = OpFAdd %5 %649 %669
%678 = OpFAdd %5 %677 %676
%679 = OpLoad %5 %629
%680 = OpFMul %5 %678 %679
OpLine %3 303 12
%681 = OpCompositeConstruct %7 %680 %56
OpStore %614 %681
OpReturn
OpFunctionEnd

1010
naga/tests/out/spv/wgsl-f16.spvasm
View File

File diff suppressed because it is too large Load Diff

35
naga/tests/out/spv/wgsl-f64.spvasm
View File

@ -1,13 +1,13 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 32
; Bound: 33
OpCapability Shader
OpCapability Float64
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %27 "main"
OpExecutionMode %27 LocalSize 1 1 1
OpEntryPoint GLCompute %28 "main"
OpExecutionMode %28 LocalSize 1 1 1
%2 = OpTypeVoid
%3 = OpTypeFloat 64
%4 = OpConstant %3 1.0
@ -21,8 +21,8 @@ OpExecutionMode %27 LocalSize 1 1 1
%15 = OpConstant %3 -1.0
%17 = OpTypePointer Function %3
%18 = OpConstantNull %3
%28 = OpTypeFunction %2
%29 = OpConstant %3 6.0
%29 = OpTypeFunction %2
%30 = OpConstant %3 6.0
%10 = OpFunction %3 None %11
%9 = OpFunctionParameter %3
%8 = OpLabel
@ -30,18 +30,19 @@ OpExecutionMode %27 LocalSize 1 1 1
%19 = OpVariable %17 Function %15
OpBranch %20
%20 = OpLabel
%21 = OpFAdd %3 %12 %13
%22 = OpFAdd %3 %21 %14
OpStore %16 %22
%23 = OpFAdd %3 %9 %21
%24 = OpFAdd %3 %23 %5
%25 = OpFAdd %3 %24 %14
OpReturnValue %25
%21 = OpLoad %3 %6
%22 = OpFAdd %3 %12 %13
%23 = OpFAdd %3 %22 %14
OpStore %16 %23
%24 = OpFAdd %3 %9 %22
%25 = OpFAdd %3 %24 %5
%26 = OpFAdd %3 %25 %14
OpReturnValue %26
OpFunctionEnd
%27 = OpFunction %2 None %28
%26 = OpLabel
OpBranch %30
%30 = OpLabel
%31 = OpFunctionCall %3 %10 %29
%28 = OpFunction %2 None %29
%27 = OpLabel
OpBranch %31
%31 = OpLabel
%32 = OpFunctionCall %3 %10 %30
OpReturn
OpFunctionEnd

1132
naga/tests/out/spv/wgsl-image.spvasm
View File

File diff suppressed because it is too large Load Diff

741
naga/tests/out/spv/wgsl-int64.spvasm
View File

@ -1,14 +1,14 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 384
; Bound: 385
OpCapability Shader
OpCapability Int64
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %368 "main"
OpExecutionMode %368 LocalSize 1 1 1
OpEntryPoint GLCompute %369 "main"
OpExecutionMode %369 LocalSize 1 1 1
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 4
OpMemberDecorate %14 2 Offset 8
@ -95,51 +95,51 @@ OpMemberDecorate %36 0 Offset 0
%55 = OpConstant %3 5
%56 = OpConstant %3 -9223372036854775808
%58 = OpTypePointer Function %3
%68 = OpTypePointer Uniform %5
%77 = OpTypePointer Uniform %6
%78 = OpConstant %5 1
%87 = OpTypePointer Uniform %7
%93 = OpConstant %7 -9.223372e18
%94 = OpConstant %7 9.2233715e18
%99 = OpTypePointer Uniform %3
%100 = OpConstant %5 7
%107 = OpTypePointer Uniform %4
%108 = OpConstant %5 3
%114 = OpTypePointer Uniform %8
%115 = OpConstant %5 4
%122 = OpTypePointer Uniform %9
%123 = OpConstant %5 5
%130 = OpTypePointer Uniform %10
%131 = OpConstant %5 6
%140 = OpTypePointer StorageBuffer %3
%147 = OpTypePointer StorageBuffer %11
%148 = OpTypePointer Uniform %11
%149 = OpConstant %5 8
%156 = OpTypePointer StorageBuffer %12
%157 = OpTypePointer Uniform %12
%158 = OpConstant %5 9
%165 = OpTypePointer StorageBuffer %13
%166 = OpTypePointer Uniform %13
%167 = OpConstant %5 10
%174 = OpTypePointer StorageBuffer %17
%194 = OpConstantNull %3
%222 = OpTypeFunction %4 %4
%228 = OpConstant %4 31
%229 = OpConstant %4 18446744073709551615
%230 = OpConstant %4 5
%232 = OpTypePointer Function %4
%263 = OpConstant %7 0.0
%264 = OpConstant %7 1.8446743e19
%298 = OpTypePointer StorageBuffer %4
%305 = OpTypePointer StorageBuffer %8
%312 = OpTypePointer StorageBuffer %9
%319 = OpTypePointer StorageBuffer %10
%326 = OpTypePointer StorageBuffer %15
%346 = OpConstantNull %4
%369 = OpTypeFunction %2
%375 = OpConstant %4 67
%376 = OpConstant %3 60
%382 = OpConstant %5 11
%69 = OpTypePointer Uniform %5
%78 = OpTypePointer Uniform %6
%79 = OpConstant %5 1
%88 = OpTypePointer Uniform %7
%94 = OpConstant %7 -9.223372e18
%95 = OpConstant %7 9.2233715e18
%100 = OpTypePointer Uniform %3
%101 = OpConstant %5 7
%108 = OpTypePointer Uniform %4
%109 = OpConstant %5 3
%115 = OpTypePointer Uniform %8
%116 = OpConstant %5 4
%123 = OpTypePointer Uniform %9
%124 = OpConstant %5 5
%131 = OpTypePointer Uniform %10
%132 = OpConstant %5 6
%141 = OpTypePointer StorageBuffer %3
%148 = OpTypePointer StorageBuffer %11
%149 = OpTypePointer Uniform %11
%150 = OpConstant %5 8
%157 = OpTypePointer StorageBuffer %12
%158 = OpTypePointer Uniform %12
%159 = OpConstant %5 9
%166 = OpTypePointer StorageBuffer %13
%167 = OpTypePointer Uniform %13
%168 = OpConstant %5 10
%175 = OpTypePointer StorageBuffer %17
%195 = OpConstantNull %3
%223 = OpTypeFunction %4 %4
%229 = OpConstant %4 31
%230 = OpConstant %4 18446744073709551615
%231 = OpConstant %4 5
%233 = OpTypePointer Function %4
%264 = OpConstant %7 0.0
%265 = OpConstant %7 1.8446743e19
%299 = OpTypePointer StorageBuffer %4
%306 = OpTypePointer StorageBuffer %8
%313 = OpTypePointer StorageBuffer %9
%320 = OpTypePointer StorageBuffer %10
%327 = OpTypePointer StorageBuffer %15
%347 = OpConstantNull %4
%370 = OpTypeFunction %2
%376 = OpConstant %4 67
%377 = OpConstant %3 60
%383 = OpConstant %5 11
%40 = OpFunction %3 None %41
%39 = OpFunctionParameter %3
%38 = OpLabel
@ -151,333 +151,334 @@ OpMemberDecorate %36 0 Offset 0
%50 = OpAccessChain %47 %35 %43
OpBranch %59
%59 = OpLabel
%60 = OpISub %3 %52 %53
%61 = OpIAdd %3 %60 %54
%62 = OpLoad %3 %57
%63 = OpIAdd %3 %62 %61
OpStore %57 %63
%64 = OpLoad %3 %57
%65 = OpIAdd %3 %64 %55
%66 = OpLoad %3 %57
%67 = OpIAdd %3 %66 %65
OpStore %57 %67
%69 = OpAccessChain %68 %44 %43
%70 = OpLoad %5 %69
%71 = OpLoad %3 %57
%72 = OpUConvert %5 %71
%73 = OpIAdd %5 %70 %72
%74 = OpSConvert %3 %73
%75 = OpLoad %3 %57
%76 = OpIAdd %3 %75 %74
OpStore %57 %76
%79 = OpAccessChain %77 %44 %78
%80 = OpLoad %6 %79
%81 = OpLoad %3 %57
%82 = OpSConvert %6 %81
%83 = OpIAdd %6 %80 %82
%84 = OpSConvert %3 %83
%85 = OpLoad %3 %57
%86 = OpIAdd %3 %85 %84
OpStore %57 %86
%88 = OpAccessChain %87 %44 %16
%89 = OpLoad %7 %88
%90 = OpLoad %3 %57
%91 = OpConvertSToF %7 %90
%92 = OpFAdd %7 %89 %91
%95 = OpExtInst %7 %1 FClamp %92 %93 %94
%96 = OpConvertFToS %3 %95
%97 = OpLoad %3 %57
%98 = OpIAdd %3 %97 %96
OpStore %57 %98
%101 = OpAccessChain %99 %44 %100
%102 = OpLoad %3 %101
%103 = OpCompositeConstruct %12 %102 %102 %102
%104 = OpCompositeExtract %3 %103 2
%105 = OpLoad %3 %57
%106 = OpIAdd %3 %105 %104
OpStore %57 %106
%109 = OpAccessChain %107 %44 %108
%110 = OpLoad %4 %109
%111 = OpBitcast %3 %110
%112 = OpLoad %3 %57
%113 = OpIAdd %3 %112 %111
OpStore %57 %113
%116 = OpAccessChain %114 %44 %115
%117 = OpLoad %8 %116
%118 = OpBitcast %11 %117
%119 = OpCompositeExtract %3 %118 1
%120 = OpLoad %3 %57
%121 = OpIAdd %3 %120 %119
OpStore %57 %121
%124 = OpAccessChain %122 %44 %123
%125 = OpLoad %9 %124
%126 = OpBitcast %12 %125
%127 = OpCompositeExtract %3 %126 2
%128 = OpLoad %3 %57
%129 = OpIAdd %3 %128 %127
OpStore %57 %129
%132 = OpAccessChain %130 %44 %131
%133 = OpLoad %10 %132
%134 = OpBitcast %13 %133
%135 = OpCompositeExtract %3 %134 3
%136 = OpLoad %3 %57
%137 = OpIAdd %3 %136 %135
OpStore %57 %137
%138 = OpLoad %3 %57
%139 = OpIAdd %3 %138 %56
OpStore %57 %139
%141 = OpAccessChain %99 %44 %100
%142 = OpLoad %3 %141
%143 = OpAccessChain %140 %46 %100
%144 = OpLoad %3 %143
%145 = OpIAdd %3 %142 %144
%146 = OpAccessChain %140 %49 %100
OpStore %146 %145
%150 = OpAccessChain %148 %44 %149
%151 = OpLoad %11 %150
%152 = OpAccessChain %147 %46 %149
%153 = OpLoad %11 %152
%154 = OpIAdd %11 %151 %153
%155 = OpAccessChain %147 %49 %149
OpStore %155 %154
%159 = OpAccessChain %157 %44 %158
%160 = OpLoad %12 %159
%161 = OpAccessChain %156 %46 %158
%162 = OpLoad %12 %161
%163 = OpIAdd %12 %160 %162
%164 = OpAccessChain %156 %49 %158
OpStore %164 %163
%168 = OpAccessChain %166 %44 %167
%169 = OpLoad %13 %168
%170 = OpAccessChain %165 %46 %167
%171 = OpLoad %13 %170
%172 = OpIAdd %13 %169 %171
%173 = OpAccessChain %165 %49 %167
OpStore %173 %172
%175 = OpAccessChain %174 %48 %78
%176 = OpLoad %17 %175
%177 = OpAccessChain %174 %50 %78
OpStore %177 %176
%178 = OpLoad %3 %57
%179 = OpExtInst %3 %1 SAbs %178
%180 = OpLoad %3 %57
%181 = OpIAdd %3 %180 %179
OpStore %57 %181
%182 = OpLoad %3 %57
%60 = OpLoad %3 %21
%61 = OpISub %3 %52 %53
%62 = OpIAdd %3 %61 %54
%63 = OpLoad %3 %57
%64 = OpIAdd %3 %63 %62
OpStore %57 %64
%65 = OpLoad %3 %57
%66 = OpIAdd %3 %65 %55
%67 = OpLoad %3 %57
%68 = OpIAdd %3 %67 %66
OpStore %57 %68
%70 = OpAccessChain %69 %44 %43
%71 = OpLoad %5 %70
%72 = OpLoad %3 %57
%73 = OpUConvert %5 %72
%74 = OpIAdd %5 %71 %73
%75 = OpSConvert %3 %74
%76 = OpLoad %3 %57
%77 = OpIAdd %3 %76 %75
OpStore %57 %77
%80 = OpAccessChain %78 %44 %79
%81 = OpLoad %6 %80
%82 = OpLoad %3 %57
%83 = OpSConvert %6 %82
%84 = OpIAdd %6 %81 %83
%85 = OpSConvert %3 %84
%86 = OpLoad %3 %57
%87 = OpIAdd %3 %86 %85
OpStore %57 %87
%89 = OpAccessChain %88 %44 %16
%90 = OpLoad %7 %89
%91 = OpLoad %3 %57
%92 = OpConvertSToF %7 %91
%93 = OpFAdd %7 %90 %92
%96 = OpExtInst %7 %1 FClamp %93 %94 %95
%97 = OpConvertFToS %3 %96
%98 = OpLoad %3 %57
%99 = OpIAdd %3 %98 %97
OpStore %57 %99
%102 = OpAccessChain %100 %44 %101
%103 = OpLoad %3 %102
%104 = OpCompositeConstruct %12 %103 %103 %103
%105 = OpCompositeExtract %3 %104 2
%106 = OpLoad %3 %57
%107 = OpIAdd %3 %106 %105
OpStore %57 %107
%110 = OpAccessChain %108 %44 %109
%111 = OpLoad %4 %110
%112 = OpBitcast %3 %111
%113 = OpLoad %3 %57
%114 = OpIAdd %3 %113 %112
OpStore %57 %114
%117 = OpAccessChain %115 %44 %116
%118 = OpLoad %8 %117
%119 = OpBitcast %11 %118
%120 = OpCompositeExtract %3 %119 1
%121 = OpLoad %3 %57
%122 = OpIAdd %3 %121 %120
OpStore %57 %122
%125 = OpAccessChain %123 %44 %124
%126 = OpLoad %9 %125
%127 = OpBitcast %12 %126
%128 = OpCompositeExtract %3 %127 2
%129 = OpLoad %3 %57
%130 = OpIAdd %3 %129 %128
OpStore %57 %130
%133 = OpAccessChain %131 %44 %132
%134 = OpLoad %10 %133
%135 = OpBitcast %13 %134
%136 = OpCompositeExtract %3 %135 3
%137 = OpLoad %3 %57
%138 = OpIAdd %3 %137 %136
OpStore %57 %138
%139 = OpLoad %3 %57
%140 = OpIAdd %3 %139 %56
OpStore %57 %140
%142 = OpAccessChain %100 %44 %101
%143 = OpLoad %3 %142
%144 = OpAccessChain %141 %46 %101
%145 = OpLoad %3 %144
%146 = OpIAdd %3 %143 %145
%147 = OpAccessChain %141 %49 %101
OpStore %147 %146
%151 = OpAccessChain %149 %44 %150
%152 = OpLoad %11 %151
%153 = OpAccessChain %148 %46 %150
%154 = OpLoad %11 %153
%155 = OpIAdd %11 %152 %154
%156 = OpAccessChain %148 %49 %150
OpStore %156 %155
%160 = OpAccessChain %158 %44 %159
%161 = OpLoad %12 %160
%162 = OpAccessChain %157 %46 %159
%163 = OpLoad %12 %162
%164 = OpIAdd %12 %161 %163
%165 = OpAccessChain %157 %49 %159
OpStore %165 %164
%169 = OpAccessChain %167 %44 %168
%170 = OpLoad %13 %169
%171 = OpAccessChain %166 %46 %168
%172 = OpLoad %13 %171
%173 = OpIAdd %13 %170 %172
%174 = OpAccessChain %166 %49 %168
OpStore %174 %173
%176 = OpAccessChain %175 %48 %79
%177 = OpLoad %17 %176
%178 = OpAccessChain %175 %50 %79
OpStore %178 %177
%179 = OpLoad %3 %57
%180 = OpExtInst %3 %1 SAbs %179
%181 = OpLoad %3 %57
%182 = OpIAdd %3 %181 %180
OpStore %57 %182
%183 = OpLoad %3 %57
%184 = OpLoad %3 %57
%186 = OpExtInst %3 %1 SMax %182 %183
%185 = OpExtInst %3 %1 SMin %186 %184
%187 = OpLoad %3 %57
%188 = OpIAdd %3 %187 %185
OpStore %57 %188
%189 = OpLoad %3 %57
%190 = OpCompositeConstruct %11 %189 %189
%191 = OpLoad %3 %57
%192 = OpCompositeConstruct %11 %191 %191
%195 = OpCompositeExtract %3 %190 0
%196 = OpCompositeExtract %3 %192 0
%197 = OpIMul %3 %195 %196
%198 = OpIAdd %3 %194 %197
%199 = OpCompositeExtract %3 %190 1
%200 = OpCompositeExtract %3 %192 1
%201 = OpIMul %3 %199 %200
%193 = OpIAdd %3 %198 %201
%202 = OpLoad %3 %57
%203 = OpIAdd %3 %202 %193
OpStore %57 %203
%204 = OpLoad %3 %57
%185 = OpLoad %3 %57
%187 = OpExtInst %3 %1 SMax %183 %184
%186 = OpExtInst %3 %1 SMin %187 %185
%188 = OpLoad %3 %57
%189 = OpIAdd %3 %188 %186
OpStore %57 %189
%190 = OpLoad %3 %57
%191 = OpCompositeConstruct %11 %190 %190
%192 = OpLoad %3 %57
%193 = OpCompositeConstruct %11 %192 %192
%196 = OpCompositeExtract %3 %191 0
%197 = OpCompositeExtract %3 %193 0
%198 = OpIMul %3 %196 %197
%199 = OpIAdd %3 %195 %198
%200 = OpCompositeExtract %3 %191 1
%201 = OpCompositeExtract %3 %193 1
%202 = OpIMul %3 %200 %201
%194 = OpIAdd %3 %199 %202
%203 = OpLoad %3 %57
%204 = OpIAdd %3 %203 %194
OpStore %57 %204
%205 = OpLoad %3 %57
%206 = OpExtInst %3 %1 SMax %204 %205
%207 = OpLoad %3 %57
%208 = OpIAdd %3 %207 %206
OpStore %57 %208
%209 = OpLoad %3 %57
%206 = OpLoad %3 %57
%207 = OpExtInst %3 %1 SMax %205 %206
%208 = OpLoad %3 %57
%209 = OpIAdd %3 %208 %207
OpStore %57 %209
%210 = OpLoad %3 %57
%211 = OpExtInst %3 %1 SMin %209 %210
%212 = OpLoad %3 %57
%213 = OpIAdd %3 %212 %211
OpStore %57 %213
%214 = OpLoad %3 %57
%215 = OpExtInst %3 %1 SSign %214
%216 = OpLoad %3 %57
%217 = OpIAdd %3 %216 %215
OpStore %57 %217
%218 = OpLoad %3 %57
OpReturnValue %218
%211 = OpLoad %3 %57
%212 = OpExtInst %3 %1 SMin %210 %211
%213 = OpLoad %3 %57
%214 = OpIAdd %3 %213 %212
OpStore %57 %214
%215 = OpLoad %3 %57
%216 = OpExtInst %3 %1 SSign %215
%217 = OpLoad %3 %57
%218 = OpIAdd %3 %217 %216
OpStore %57 %218
%219 = OpLoad %3 %57
OpReturnValue %219
OpFunctionEnd
%221 = OpFunction %4 None %222
%220 = OpFunctionParameter %4
%219 = OpLabel
%231 = OpVariable %232 Function %20
%223 = OpAccessChain %42 %23 %43
%224 = OpAccessChain %45 %26 %43
%225 = OpAccessChain %47 %29 %43
%226 = OpAccessChain %45 %32 %43
%227 = OpAccessChain %47 %35 %43
OpBranch %233
%233 = OpLabel
%234 = OpIAdd %4 %228 %229
%235 = OpISub %4 %234 %229
%236 = OpLoad %4 %231
%237 = OpIAdd %4 %236 %235
OpStore %231 %237
%238 = OpLoad %4 %231
%239 = OpIAdd %4 %238 %230
%240 = OpLoad %4 %231
%241 = OpIAdd %4 %240 %239
OpStore %231 %241
%242 = OpAccessChain %68 %223 %43
%243 = OpLoad %5 %242
%244 = OpLoad %4 %231
%245 = OpUConvert %5 %244
%246 = OpIAdd %5 %243 %245
%247 = OpUConvert %4 %246
%248 = OpLoad %4 %231
%249 = OpIAdd %4 %248 %247
OpStore %231 %249
%250 = OpAccessChain %77 %223 %78
%251 = OpLoad %6 %250
%252 = OpLoad %4 %231
%253 = OpSConvert %6 %252
%254 = OpIAdd %6 %251 %253
%255 = OpUConvert %4 %254
%256 = OpLoad %4 %231
%257 = OpIAdd %4 %256 %255
OpStore %231 %257
%258 = OpAccessChain %87 %223 %16
%259 = OpLoad %7 %258
%260 = OpLoad %4 %231
%261 = OpConvertUToF %7 %260
%262 = OpFAdd %7 %259 %261
%265 = OpExtInst %7 %1 FClamp %262 %263 %264
%266 = OpConvertFToU %4 %265
%267 = OpLoad %4 %231
%268 = OpIAdd %4 %267 %266
OpStore %231 %268
%269 = OpAccessChain %107 %223 %108
%270 = OpLoad %4 %269
%271 = OpCompositeConstruct %9 %270 %270 %270
%272 = OpCompositeExtract %4 %271 2
%273 = OpLoad %4 %231
%274 = OpIAdd %4 %273 %272
OpStore %231 %274
%275 = OpAccessChain %99 %223 %100
%276 = OpLoad %3 %275
%277 = OpBitcast %4 %276
%278 = OpLoad %4 %231
%279 = OpIAdd %4 %278 %277
OpStore %231 %279
%280 = OpAccessChain %148 %223 %149
%281 = OpLoad %11 %280
%282 = OpBitcast %8 %281
%283 = OpCompositeExtract %4 %282 1
%284 = OpLoad %4 %231
%285 = OpIAdd %4 %284 %283
OpStore %231 %285
%286 = OpAccessChain %157 %223 %158
%287 = OpLoad %12 %286
%288 = OpBitcast %9 %287
%289 = OpCompositeExtract %4 %288 2
%290 = OpLoad %4 %231
%291 = OpIAdd %4 %290 %289
OpStore %231 %291
%292 = OpAccessChain %166 %223 %167
%293 = OpLoad %13 %292
%294 = OpBitcast %10 %293
%295 = OpCompositeExtract %4 %294 3
%296 = OpLoad %4 %231
%297 = OpIAdd %4 %296 %295
OpStore %231 %297
%299 = OpAccessChain %107 %223 %108
%300 = OpLoad %4 %299
%301 = OpAccessChain %298 %224 %108
%302 = OpLoad %4 %301
%303 = OpIAdd %4 %300 %302
%304 = OpAccessChain %298 %226 %108
OpStore %304 %303
%306 = OpAccessChain %114 %223 %115
%307 = OpLoad %8 %306
%308 = OpAccessChain %305 %224 %115
%309 = OpLoad %8 %308
%310 = OpIAdd %8 %307 %309
%311 = OpAccessChain %305 %226 %115
OpStore %311 %310
%313 = OpAccessChain %122 %223 %123
%314 = OpLoad %9 %313
%315 = OpAccessChain %312 %224 %123
%316 = OpLoad %9 %315
%317 = OpIAdd %9 %314 %316
%318 = OpAccessChain %312 %226 %123
OpStore %318 %317
%320 = OpAccessChain %130 %223 %131
%321 = OpLoad %10 %320
%322 = OpAccessChain %319 %224 %131
%323 = OpLoad %10 %322
%324 = OpIAdd %10 %321 %323
%325 = OpAccessChain %319 %226 %131
OpStore %325 %324
%327 = OpAccessChain %326 %225 %43
%328 = OpLoad %15 %327
%329 = OpAccessChain %326 %227 %43
OpStore %329 %328
%330 = OpLoad %4 %231
%331 = OpCopyObject %4 %330
%332 = OpLoad %4 %231
%333 = OpIAdd %4 %332 %331
OpStore %231 %333
%334 = OpLoad %4 %231
%335 = OpLoad %4 %231
%336 = OpLoad %4 %231
%338 = OpExtInst %4 %1 UMax %334 %335
%337 = OpExtInst %4 %1 UMin %338 %336
%339 = OpLoad %4 %231
%340 = OpIAdd %4 %339 %337
OpStore %231 %340
%341 = OpLoad %4 %231
%342 = OpCompositeConstruct %8 %341 %341
%343 = OpLoad %4 %231
%344 = OpCompositeConstruct %8 %343 %343
%347 = OpCompositeExtract %4 %342 0
%348 = OpCompositeExtract %4 %344 0
%349 = OpIMul %4 %347 %348
%350 = OpIAdd %4 %346 %349
%351 = OpCompositeExtract %4 %342 1
%352 = OpCompositeExtract %4 %344 1
%353 = OpIMul %4 %351 %352
%345 = OpIAdd %4 %350 %353
%354 = OpLoad %4 %231
%355 = OpIAdd %4 %354 %345
OpStore %231 %355
%356 = OpLoad %4 %231
%357 = OpLoad %4 %231
%358 = OpExtInst %4 %1 UMax %356 %357
%359 = OpLoad %4 %231
%360 = OpIAdd %4 %359 %358
OpStore %231 %360
%361 = OpLoad %4 %231
%362 = OpLoad %4 %231
%363 = OpExtInst %4 %1 UMin %361 %362
%364 = OpLoad %4 %231
%365 = OpIAdd %4 %364 %363
OpStore %231 %365
%366 = OpLoad %4 %231
OpReturnValue %366
%222 = OpFunction %4 None %223
%221 = OpFunctionParameter %4
%220 = OpLabel
%232 = OpVariable %233 Function %20
%224 = OpAccessChain %42 %23 %43
%225 = OpAccessChain %45 %26 %43
%226 = OpAccessChain %47 %29 %43
%227 = OpAccessChain %45 %32 %43
%228 = OpAccessChain %47 %35 %43
OpBranch %234
%234 = OpLabel
%235 = OpIAdd %4 %229 %230
%236 = OpISub %4 %235 %230
%237 = OpLoad %4 %232
%238 = OpIAdd %4 %237 %236
OpStore %232 %238
%239 = OpLoad %4 %232
%240 = OpIAdd %4 %239 %231
%241 = OpLoad %4 %232
%242 = OpIAdd %4 %241 %240
OpStore %232 %242
%243 = OpAccessChain %69 %224 %43
%244 = OpLoad %5 %243
%245 = OpLoad %4 %232
%246 = OpUConvert %5 %245
%247 = OpIAdd %5 %244 %246
%248 = OpUConvert %4 %247
%249 = OpLoad %4 %232
%250 = OpIAdd %4 %249 %248
OpStore %232 %250
%251 = OpAccessChain %78 %224 %79
%252 = OpLoad %6 %251
%253 = OpLoad %4 %232
%254 = OpSConvert %6 %253
%255 = OpIAdd %6 %252 %254
%256 = OpUConvert %4 %255
%257 = OpLoad %4 %232
%258 = OpIAdd %4 %257 %256
OpStore %232 %258
%259 = OpAccessChain %88 %224 %16
%260 = OpLoad %7 %259
%261 = OpLoad %4 %232
%262 = OpConvertUToF %7 %261
%263 = OpFAdd %7 %260 %262
%266 = OpExtInst %7 %1 FClamp %263 %264 %265
%267 = OpConvertFToU %4 %266
%268 = OpLoad %4 %232
%269 = OpIAdd %4 %268 %267
OpStore %232 %269
%270 = OpAccessChain %108 %224 %109
%271 = OpLoad %4 %270
%272 = OpCompositeConstruct %9 %271 %271 %271
%273 = OpCompositeExtract %4 %272 2
%274 = OpLoad %4 %232
%275 = OpIAdd %4 %274 %273
OpStore %232 %275
%276 = OpAccessChain %100 %224 %101
%277 = OpLoad %3 %276
%278 = OpBitcast %4 %277
%279 = OpLoad %4 %232
%280 = OpIAdd %4 %279 %278
OpStore %232 %280
%281 = OpAccessChain %149 %224 %150
%282 = OpLoad %11 %281
%283 = OpBitcast %8 %282
%284 = OpCompositeExtract %4 %283 1
%285 = OpLoad %4 %232
%286 = OpIAdd %4 %285 %284
OpStore %232 %286
%287 = OpAccessChain %158 %224 %159
%288 = OpLoad %12 %287
%289 = OpBitcast %9 %288
%290 = OpCompositeExtract %4 %289 2
%291 = OpLoad %4 %232
%292 = OpIAdd %4 %291 %290
OpStore %232 %292
%293 = OpAccessChain %167 %224 %168
%294 = OpLoad %13 %293
%295 = OpBitcast %10 %294
%296 = OpCompositeExtract %4 %295 3
%297 = OpLoad %4 %232
%298 = OpIAdd %4 %297 %296
OpStore %232 %298
%300 = OpAccessChain %108 %224 %109
%301 = OpLoad %4 %300
%302 = OpAccessChain %299 %225 %109
%303 = OpLoad %4 %302
%304 = OpIAdd %4 %301 %303
%305 = OpAccessChain %299 %227 %109
OpStore %305 %304
%307 = OpAccessChain %115 %224 %116
%308 = OpLoad %8 %307
%309 = OpAccessChain %306 %225 %116
%310 = OpLoad %8 %309
%311 = OpIAdd %8 %308 %310
%312 = OpAccessChain %306 %227 %116
OpStore %312 %311
%314 = OpAccessChain %123 %224 %124
%315 = OpLoad %9 %314
%316 = OpAccessChain %313 %225 %124
%317 = OpLoad %9 %316
%318 = OpIAdd %9 %315 %317
%319 = OpAccessChain %313 %227 %124
OpStore %319 %318
%321 = OpAccessChain %131 %224 %132
%322 = OpLoad %10 %321
%323 = OpAccessChain %320 %225 %132
%324 = OpLoad %10 %323
%325 = OpIAdd %10 %322 %324
%326 = OpAccessChain %320 %227 %132
OpStore %326 %325
%328 = OpAccessChain %327 %226 %43
%329 = OpLoad %15 %328
%330 = OpAccessChain %327 %228 %43
OpStore %330 %329
%331 = OpLoad %4 %232
%332 = OpCopyObject %4 %331
%333 = OpLoad %4 %232
%334 = OpIAdd %4 %333 %332
OpStore %232 %334
%335 = OpLoad %4 %232
%336 = OpLoad %4 %232
%337 = OpLoad %4 %232
%339 = OpExtInst %4 %1 UMax %335 %336
%338 = OpExtInst %4 %1 UMin %339 %337
%340 = OpLoad %4 %232
%341 = OpIAdd %4 %340 %338
OpStore %232 %341
%342 = OpLoad %4 %232
%343 = OpCompositeConstruct %8 %342 %342
%344 = OpLoad %4 %232
%345 = OpCompositeConstruct %8 %344 %344
%348 = OpCompositeExtract %4 %343 0
%349 = OpCompositeExtract %4 %345 0
%350 = OpIMul %4 %348 %349
%351 = OpIAdd %4 %347 %350
%352 = OpCompositeExtract %4 %343 1
%353 = OpCompositeExtract %4 %345 1
%354 = OpIMul %4 %352 %353
%346 = OpIAdd %4 %351 %354
%355 = OpLoad %4 %232
%356 = OpIAdd %4 %355 %346
OpStore %232 %356
%357 = OpLoad %4 %232
%358 = OpLoad %4 %232
%359 = OpExtInst %4 %1 UMax %357 %358
%360 = OpLoad %4 %232
%361 = OpIAdd %4 %360 %359
OpStore %232 %361
%362 = OpLoad %4 %232
%363 = OpLoad %4 %232
%364 = OpExtInst %4 %1 UMin %362 %363
%365 = OpLoad %4 %232
%366 = OpIAdd %4 %365 %364
OpStore %232 %366
%367 = OpLoad %4 %232
OpReturnValue %367
OpFunctionEnd
%368 = OpFunction %2 None %369
%367 = OpLabel
%370 = OpAccessChain %42 %23 %43
%371 = OpAccessChain %45 %26 %43
%372 = OpAccessChain %47 %29 %43
%373 = OpAccessChain %45 %32 %43
%374 = OpAccessChain %47 %35 %43
OpBranch %377
%377 = OpLabel
%378 = OpFunctionCall %4 %221 %375
%379 = OpFunctionCall %3 %40 %376
%380 = OpBitcast %4 %379
%381 = OpIAdd %4 %378 %380
%383 = OpAccessChain %298 %373 %382
OpStore %383 %381
%369 = OpFunction %2 None %370
%368 = OpLabel
%371 = OpAccessChain %42 %23 %43
%372 = OpAccessChain %45 %26 %43
%373 = OpAccessChain %47 %29 %43
%374 = OpAccessChain %45 %32 %43
%375 = OpAccessChain %47 %35 %43
OpBranch %378
%378 = OpLabel
%379 = OpFunctionCall %4 %222 %376
%380 = OpFunctionCall %3 %40 %377
%381 = OpBitcast %4 %380
%382 = OpIAdd %4 %379 %381
%384 = OpAccessChain %299 %374 %383
OpStore %384 %382
OpReturn
OpFunctionEnd

58
naga/tests/out/wgsl/wgsl-abstract-types-var.wgsl
View File

@ -52,6 +52,63 @@ var<private> iavipai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
var<private> iavfpai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
var<private> iavfpaf: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
fn globals() {
let phony = xvipaiai_1;
let phony_1 = xvupaiai_1;
let phony_2 = xvfpaiai_1;
let phony_3 = xvfpafaf_1;
let phony_4 = xvfpaiaf_1;
let phony_5 = xvupuai_2;
let phony_6 = xvupaiu_2;
let phony_7 = xvuuai_2;
let phony_8 = xvuaiu_2;
let phony_9 = xvip_1;
let phony_10 = xvup_1;
let phony_11 = xvfp_1;
let phony_12 = xmfp_1;
let phony_13 = xmfpaiaiaiai_1;
let phony_14 = xmfpafaiaiai_1;
let phony_15 = xmfpaiafaiai_1;
let phony_16 = xmfpaiaiafai_1;
let phony_17 = xmfpaiaiaiaf_1;
let phony_18 = xvispai_1;
let phony_19 = xvfspaf_1;
let phony_20 = xvis_ai_1;
let phony_21 = xvus_ai_1;
let phony_22 = xvfs_ai_1;
let phony_23 = xvfs_af_1;
let phony_24 = xafafaf_1;
let phony_25 = xafaiai_1;
let phony_26 = xaipaiai_1;
let phony_27 = xaupaiai;
let phony_28 = xafpaiaf_1;
let phony_29 = xafpafai_1;
let phony_30 = xafpafaf_1;
let phony_31 = xavipai_1;
let phony_32 = xavfpai_1;
let phony_33 = xavfpaf_1;
let phony_34 = xvisai_1;
let phony_35 = xvusai_1;
let phony_36 = xvfsai_1;
let phony_37 = xvfsaf_1;
let phony_38 = ivispai;
let phony_39 = ivfspaf;
let phony_40 = ivis_ai;
let phony_41 = ivus_ai;
let phony_42 = ivfs_ai;
let phony_43 = ivfs_af;
let phony_44 = iafafaf;
let phony_45 = iafaiai;
let phony_46 = iaipaiai_1;
let phony_47 = iafpafaf_1;
let phony_48 = iafpaiaf_1;
let phony_49 = iafpafai_1;
let phony_50 = iavipai;
let phony_51 = iavfpai;
let phony_52 = iavfpaf;
return;
}
fn all_constant_arguments() {
var xvipaiai: vec2<i32> = vec2<i32>(42i, 43i);
var xvupaiai: vec2<u32> = vec2<u32>(44u, 45u);
@ -288,6 +345,7 @@ fn mixed_constant_and_runtime_arguments() {
@compute @workgroup_size(1, 1, 1)
fn main() {
globals();
all_constant_arguments();
mixed_constant_and_runtime_arguments();
return;

7
naga/tests/out/wgsl/wgsl-access.wgsl
View File

@ -42,7 +42,7 @@ struct Outer {
thing: u32,
}
var<private> global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0i);
var<private> msl_padding_global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0i);
@group(0) @binding(0)
var<storage, read_write> bar: Bar;
@group(0) @binding(1)
@ -230,6 +230,7 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let baz_1 = foo;
foo = 1f;
let phony = msl_padding_global_const;
test_matrix_within_struct_accesses();
test_matrix_within_array_within_struct_accesses();
let _matrix = bar._matrix;
@ -238,11 +239,11 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let a_2 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
let c = qux;
let data_pointer = (&bar.data[0].value);
let _e33 = read_from_private((&foo));
let _e35 = read_from_private((&foo));
c2_ = array<i32, 5>(a_2, i32(b), 3i, 4i, 5i);
c2_[(vi + 1u)] = 42i;
let value_1 = c2_[vi];
let _e47 = test_arr_as_arg(array<array<f32, 10>, 5>());
let _e49 = test_arr_as_arg(array<array<f32, 10>, 5>());
return vec4<f32>((_matrix * vec4<f32>(vec4(value_1))), 2f);
}

184
naga/tests/out/wgsl/wgsl-f16.wgsl
View File

@ -48,121 +48,123 @@ var<storage, read_write> output: UniformCompatible;
var<storage, read_write> output_arrays: StorageCompatible;
fn f16_function(x: f16) -> f16 {
var l: LayoutTest;
var val: f16 = 15.203125h;
let _e4 = val;
val = (_e4 + -33344h);
let phony = private_variable;
let _e6 = val;
let _e9 = val;
val = (_e9 + (_e6 + 5h));
let _e13 = input_uniform.val_f32_;
let _e14 = val;
let _e18 = val;
val = (_e18 + f16((_e13 + f32(_e14))));
let _e22 = input_uniform.val_f16_;
let _e25 = val;
val = (_e25 + vec3(_e22).z);
val = (_e6 + -33344h);
let _e8 = val;
let _e11 = val;
val = (_e11 + (_e8 + 5h));
let _e15 = input_uniform.val_f32_;
let _e16 = val;
let _e20 = val;
val = (_e20 + f16((_e15 + f32(_e16))));
let _e24 = input_uniform.val_f16_;
let _e27 = val;
val = (_e27 + vec3(_e24).z);
output.val_i32_ = 65504i;
output.val_i32_ = -65504i;
output.val_u32_ = 65504u;
output.val_u32_ = 0u;
output.val_f32_ = 65504f;
output.val_f32_ = -65504f;
let _e49 = input_uniform.val_f16_;
let _e52 = input_storage.val_f16_;
output.val_f16_ = (_e49 + _e52);
let _e58 = input_uniform.val_f16_2_;
let _e61 = input_storage.val_f16_2_;
output.val_f16_2_ = (_e58 + _e61);
let _e67 = input_uniform.val_f16_3_;
let _e70 = input_storage.val_f16_3_;
output.val_f16_3_ = (_e67 + _e70);
let _e76 = input_uniform.val_f16_4_;
let _e79 = input_storage.val_f16_4_;
output.val_f16_4_ = (_e76 + _e79);
let _e85 = input_uniform.val_mat2x2_;
let _e88 = input_storage.val_mat2x2_;
output.val_mat2x2_ = (_e85 + _e88);
let _e94 = input_uniform.val_mat2x3_;
let _e97 = input_storage.val_mat2x3_;
output.val_mat2x3_ = (_e94 + _e97);
let _e103 = input_uniform.val_mat2x4_;
let _e106 = input_storage.val_mat2x4_;
output.val_mat2x4_ = (_e103 + _e106);
let _e112 = input_uniform.val_mat3x2_;
let _e115 = input_storage.val_mat3x2_;
output.val_mat3x2_ = (_e112 + _e115);
let _e121 = input_uniform.val_mat3x3_;
let _e124 = input_storage.val_mat3x3_;
output.val_mat3x3_ = (_e121 + _e124);
let _e130 = input_uniform.val_mat3x4_;
let _e133 = input_storage.val_mat3x4_;
output.val_mat3x4_ = (_e130 + _e133);
let _e139 = input_uniform.val_mat4x2_;
let _e142 = input_storage.val_mat4x2_;
output.val_mat4x2_ = (_e139 + _e142);
let _e148 = input_uniform.val_mat4x3_;
let _e151 = input_storage.val_mat4x3_;
output.val_mat4x3_ = (_e148 + _e151);
let _e157 = input_uniform.val_mat4x4_;
let _e160 = input_storage.val_mat4x4_;
output.val_mat4x4_ = (_e157 + _e160);
let _e166 = input_arrays.val_f16_array_2_;
output_arrays.val_f16_array_2_ = _e166;
let _e167 = val;
let _e51 = input_uniform.val_f16_;
let _e54 = input_storage.val_f16_;
output.val_f16_ = (_e51 + _e54);
let _e60 = input_uniform.val_f16_2_;
let _e63 = input_storage.val_f16_2_;
output.val_f16_2_ = (_e60 + _e63);
let _e69 = input_uniform.val_f16_3_;
let _e72 = input_storage.val_f16_3_;
output.val_f16_3_ = (_e69 + _e72);
let _e78 = input_uniform.val_f16_4_;
let _e81 = input_storage.val_f16_4_;
output.val_f16_4_ = (_e78 + _e81);
let _e87 = input_uniform.val_mat2x2_;
let _e90 = input_storage.val_mat2x2_;
output.val_mat2x2_ = (_e87 + _e90);
let _e96 = input_uniform.val_mat2x3_;
let _e99 = input_storage.val_mat2x3_;
output.val_mat2x3_ = (_e96 + _e99);
let _e105 = input_uniform.val_mat2x4_;
let _e108 = input_storage.val_mat2x4_;
output.val_mat2x4_ = (_e105 + _e108);
let _e114 = input_uniform.val_mat3x2_;
let _e117 = input_storage.val_mat3x2_;
output.val_mat3x2_ = (_e114 + _e117);
let _e123 = input_uniform.val_mat3x3_;
let _e126 = input_storage.val_mat3x3_;
output.val_mat3x3_ = (_e123 + _e126);
let _e132 = input_uniform.val_mat3x4_;
let _e135 = input_storage.val_mat3x4_;
output.val_mat3x4_ = (_e132 + _e135);
let _e141 = input_uniform.val_mat4x2_;
let _e144 = input_storage.val_mat4x2_;
output.val_mat4x2_ = (_e141 + _e144);
let _e150 = input_uniform.val_mat4x3_;
let _e153 = input_storage.val_mat4x3_;
output.val_mat4x3_ = (_e150 + _e153);
let _e159 = input_uniform.val_mat4x4_;
let _e162 = input_storage.val_mat4x4_;
output.val_mat4x4_ = (_e159 + _e162);
let _e168 = input_arrays.val_f16_array_2_;
output_arrays.val_f16_array_2_ = _e168;
let _e169 = val;
val = (_e169 + abs(_e167));
let _e171 = val;
let _e172 = val;
val = (_e171 + abs(_e169));
let _e173 = val;
let _e174 = val;
let _e175 = val;
val = (_e175 + clamp(_e171, _e172, _e173));
let _e177 = val;
val = (_e177 + clamp(_e173, _e174, _e175));
let _e179 = val;
let _e182 = val;
val = (_e182 + dot(vec2(_e177), vec2(_e179)));
let _e181 = val;
let _e184 = val;
let _e185 = val;
val = (_e184 + dot(vec2(_e179), vec2(_e181)));
let _e186 = val;
let _e187 = val;
val = (_e187 + max(_e184, _e185));
let _e189 = val;
let _e190 = val;
val = (_e189 + max(_e186, _e187));
let _e191 = val;
let _e192 = val;
val = (_e192 + min(_e189, _e190));
let _e194 = val;
val = (_e194 + min(_e191, _e192));
let _e196 = val;
val = (_e196 + sign(_e194));
let _e199 = val;
val = (_e199 + 1h);
let _e203 = input_uniform.val_f16_2_;
let float_vec2_ = vec2<f32>(_e203);
let _e198 = val;
val = (_e198 + sign(_e196));
let _e201 = val;
val = (_e201 + 1h);
let _e205 = input_uniform.val_f16_2_;
let float_vec2_ = vec2<f32>(_e205);
output.val_f16_2_ = vec2<f16>(float_vec2_);
let _e210 = input_uniform.val_f16_3_;
let float_vec3_ = vec3<f32>(_e210);
let _e212 = input_uniform.val_f16_3_;
let float_vec3_ = vec3<f32>(_e212);
output.val_f16_3_ = vec3<f16>(float_vec3_);
let _e217 = input_uniform.val_f16_4_;
let float_vec4_ = vec4<f32>(_e217);
let _e219 = input_uniform.val_f16_4_;
let float_vec4_ = vec4<f32>(_e219);
output.val_f16_4_ = vec4<f16>(float_vec4_);
let _e226 = input_uniform.val_mat2x2_;
output.val_mat2x2_ = mat2x2<f16>(mat2x2<f32>(_e226));
let _e233 = input_uniform.val_mat2x3_;
output.val_mat2x3_ = mat2x3<f16>(mat2x3<f32>(_e233));
let _e240 = input_uniform.val_mat2x4_;
output.val_mat2x4_ = mat2x4<f16>(mat2x4<f32>(_e240));
let _e247 = input_uniform.val_mat3x2_;
output.val_mat3x2_ = mat3x2<f16>(mat3x2<f32>(_e247));
let _e254 = input_uniform.val_mat3x3_;
output.val_mat3x3_ = mat3x3<f16>(mat3x3<f32>(_e254));
let _e261 = input_uniform.val_mat3x4_;
output.val_mat3x4_ = mat3x4<f16>(mat3x4<f32>(_e261));
let _e268 = input_uniform.val_mat4x2_;
output.val_mat4x2_ = mat4x2<f16>(mat4x2<f32>(_e268));
let _e275 = input_uniform.val_mat4x3_;
output.val_mat4x3_ = mat4x3<f16>(mat4x3<f32>(_e275));
let _e282 = input_uniform.val_mat4x4_;
output.val_mat4x4_ = mat4x4<f16>(mat4x4<f32>(_e282));
let _e285 = val;
return _e285;
let _e228 = input_uniform.val_mat2x2_;
output.val_mat2x2_ = mat2x2<f16>(mat2x2<f32>(_e228));
let _e235 = input_uniform.val_mat2x3_;
output.val_mat2x3_ = mat2x3<f16>(mat2x3<f32>(_e235));
let _e242 = input_uniform.val_mat2x4_;
output.val_mat2x4_ = mat2x4<f16>(mat2x4<f32>(_e242));
let _e249 = input_uniform.val_mat3x2_;
output.val_mat3x2_ = mat3x2<f16>(mat3x2<f32>(_e249));
let _e256 = input_uniform.val_mat3x3_;
output.val_mat3x3_ = mat3x3<f16>(mat3x3<f32>(_e256));
let _e263 = input_uniform.val_mat3x4_;
output.val_mat3x4_ = mat3x4<f16>(mat3x4<f32>(_e263));
let _e270 = input_uniform.val_mat4x2_;
output.val_mat4x2_ = mat4x2<f16>(mat4x2<f32>(_e270));
let _e277 = input_uniform.val_mat4x3_;
output.val_mat4x3_ = mat4x3<f16>(mat4x3<f32>(_e277));
let _e284 = input_uniform.val_mat4x4_;
output.val_mat4x4_ = mat4x4<f16>(mat4x4<f32>(_e284));
let _e287 = val;
return _e287;
}
@compute @workgroup_size(1, 1, 1)

1
naga/tests/out/wgsl/wgsl-f64.wgsl
View File

@ -6,6 +6,7 @@ fn f(x: f64) -> f64 {
var z: f64;
var w: f64 = -1.0lf;
let phony = v;
let y = (30.0lf + 400.0lf);
z = (y + 5.0lf);
return (((x + y) + k) + 5.0lf);

1
naga/tests/out/wgsl/wgsl-image.wgsl
View File

@ -55,6 +55,7 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
let value5_ = textureLoad(image_array_src, itc, local_id.z, (i32(local_id.z) + 1i));
let value6_ = textureLoad(image_array_src, itc, i32(local_id.z), (i32(local_id.z) + 1i));
let value7_ = textureLoad(image_1d_src, i32(local_id.x), i32(local_id.z));
let value8_ = textureLoad(image_dup_src, i32(local_id.x));
let value1u = textureLoad(image_mipmapped_src, vec2<u32>(itc), i32(local_id.z));
let value2u = textureLoad(image_multisampled_src, vec2<u32>(itc), i32(local_id.z));
let value3u = textureLoad(image_multisampled_src, vec2<u32>(itc), u32(local_id.z));

119
naga/tests/out/wgsl/wgsl-int64.wgsl
View File

@ -35,79 +35,80 @@ var<storage, read_write> output_arrays: StorageCompatible;
fn int64_function(x: i64) -> i64 {
var val: i64 = 20li;
let _e8 = val;
val = (_e8 + ((31li - 1002003004005006li) + -9223372036854775807li));
let phony = private_variable;
let _e10 = val;
let _e13 = val;
val = (_e13 + (_e10 + 5li));
let _e17 = input_uniform.val_u32_;
let _e18 = val;
let _e22 = val;
val = (_e22 + i64((_e17 + u32(_e18))));
let _e26 = input_uniform.val_i32_;
let _e27 = val;
let _e31 = val;
val = (_e31 + i64((_e26 + i32(_e27))));
let _e35 = input_uniform.val_f32_;
let _e36 = val;
let _e40 = val;
val = (_e40 + i64((_e35 + f32(_e36))));
let _e44 = input_uniform.val_i64_;
let _e47 = val;
val = (_e47 + vec3(_e44).z);
let _e51 = input_uniform.val_u64_;
let _e53 = val;
val = (_e53 + bitcast<i64>(_e51));
let _e57 = input_uniform.val_u64_2_;
let _e60 = val;
val = (_e60 + bitcast<vec2<i64>>(_e57).y);
let _e64 = input_uniform.val_u64_3_;
let _e67 = val;
val = (_e67 + bitcast<vec3<i64>>(_e64).z);
let _e71 = input_uniform.val_u64_4_;
let _e74 = val;
val = (_e74 + bitcast<vec4<i64>>(_e71).w);
let _e77 = val;
val = (_e77 + i64(-9223372036854775807 - 1));
let _e83 = input_uniform.val_i64_;
let _e86 = input_storage.val_i64_;
output.val_i64_ = (_e83 + _e86);
let _e92 = input_uniform.val_i64_2_;
let _e95 = input_storage.val_i64_2_;
output.val_i64_2_ = (_e92 + _e95);
let _e101 = input_uniform.val_i64_3_;
let _e104 = input_storage.val_i64_3_;
output.val_i64_3_ = (_e101 + _e104);
let _e110 = input_uniform.val_i64_4_;
let _e113 = input_storage.val_i64_4_;
output.val_i64_4_ = (_e110 + _e113);
let _e119 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e119;
let _e120 = val;
val = (_e10 + ((31li - 1002003004005006li) + -9223372036854775807li));
let _e12 = val;
let _e15 = val;
val = (_e15 + (_e12 + 5li));
let _e19 = input_uniform.val_u32_;
let _e20 = val;
let _e24 = val;
val = (_e24 + i64((_e19 + u32(_e20))));
let _e28 = input_uniform.val_i32_;
let _e29 = val;
let _e33 = val;
val = (_e33 + i64((_e28 + i32(_e29))));
let _e37 = input_uniform.val_f32_;
let _e38 = val;
let _e42 = val;
val = (_e42 + i64((_e37 + f32(_e38))));
let _e46 = input_uniform.val_i64_;
let _e49 = val;
val = (_e49 + vec3(_e46).z);
let _e53 = input_uniform.val_u64_;
let _e55 = val;
val = (_e55 + bitcast<i64>(_e53));
let _e59 = input_uniform.val_u64_2_;
let _e62 = val;
val = (_e62 + bitcast<vec2<i64>>(_e59).y);
let _e66 = input_uniform.val_u64_3_;
let _e69 = val;
val = (_e69 + bitcast<vec3<i64>>(_e66).z);
let _e73 = input_uniform.val_u64_4_;
let _e76 = val;
val = (_e76 + bitcast<vec4<i64>>(_e73).w);
let _e79 = val;
val = (_e79 + i64(-9223372036854775807 - 1));
let _e85 = input_uniform.val_i64_;
let _e88 = input_storage.val_i64_;
output.val_i64_ = (_e85 + _e88);
let _e94 = input_uniform.val_i64_2_;
let _e97 = input_storage.val_i64_2_;
output.val_i64_2_ = (_e94 + _e97);
let _e103 = input_uniform.val_i64_3_;
let _e106 = input_storage.val_i64_3_;
output.val_i64_3_ = (_e103 + _e106);
let _e112 = input_uniform.val_i64_4_;
let _e115 = input_storage.val_i64_4_;
output.val_i64_4_ = (_e112 + _e115);
let _e121 = input_arrays.val_i64_array_2_;
output_arrays.val_i64_array_2_ = _e121;
let _e122 = val;
val = (_e122 + abs(_e120));
let _e124 = val;
let _e125 = val;
val = (_e124 + abs(_e122));
let _e126 = val;
let _e127 = val;
let _e128 = val;
val = (_e128 + clamp(_e124, _e125, _e126));
let _e130 = val;
val = (_e130 + clamp(_e126, _e127, _e128));
let _e132 = val;
let _e135 = val;
val = (_e135 + dot(vec2(_e130), vec2(_e132)));
let _e134 = val;
let _e137 = val;
let _e138 = val;
val = (_e137 + dot(vec2(_e132), vec2(_e134)));
let _e139 = val;
let _e140 = val;
val = (_e140 + max(_e137, _e138));
let _e142 = val;
let _e143 = val;
val = (_e142 + max(_e139, _e140));
let _e144 = val;
let _e145 = val;
val = (_e145 + min(_e142, _e143));
let _e147 = val;
val = (_e147 + min(_e144, _e145));
let _e149 = val;
val = (_e149 + sign(_e147));
let _e151 = val;
return _e151;
val = (_e151 + sign(_e149));
let _e153 = val;
return _e153;
}
fn uint64_function(x_1: u64) -> u64 {