nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-22 14:55:05 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix texture built-ins where u32 was expected (#2245)

Browse Source 
- The Typifier was updated to expect `uint`
- Both `glsl` and `spv` frontends where updated to cast the result to `sint`.
- Both `glsl` and `spv` backends where updated to cast the result to `uint`.
- Remove cast in `msl` backend.
This commit is contained in:
Evan Mark Hopkins 2023-02-13 08:13:58 -05:00 committed by GitHub
parent 40b8f66146
commit 1ad47f732d
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
23 changed files with 1836 additions and 1704 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
src/back/glsl/mod.rs
View File

@ -2534,6 +2534,16 @@ impl<'a, W: Write> Writer<'a, W> {
crate::ImageDimension::D3 => 3,
crate::ImageDimension::Cube => 2,
};
if let crate::ImageQuery::Size { .. } = query {
match components {
1 => write!(self.out, "uint(")?,
_ => write!(self.out, "uvec{components}(")?,
}
} else {
write!(self.out, "uint(")?;
}
match query {
crate::ImageQuery::Size { level } => {
match class {
@ -2593,6 +2603,8 @@ impl<'a, W: Write> Writer<'a, W> {
write!(self.out, ")",)?;
}
}
write!(self.out, ")")?;
}
// `Unary` is pretty straightforward
// "-" - for `Negate`

11
src/back/msl/writer.rs
View File

@ -1453,24 +1453,21 @@ impl<W: Write> Writer<W> {
self.put_image_size_query(
image,
level.map(LevelOfDetail::Direct),
crate::ScalarKind::Sint,
crate::ScalarKind::Uint,
context,
)?;
}
crate::ImageQuery::NumLevels => {
write!(self.out, "int(")?;
self.put_expression(image, context, false)?;
write!(self.out, ".get_num_mip_levels())")?;
write!(self.out, ".get_num_mip_levels()")?;
}
crate::ImageQuery::NumLayers => {
write!(self.out, "int(")?;
self.put_expression(image, context, false)?;
write!(self.out, ".get_array_size())")?;
write!(self.out, ".get_array_size()")?;
}
crate::ImageQuery::NumSamples => {
write!(self.out, "int(")?;
self.put_expression(image, context, false)?;
write!(self.out, ".get_num_samples())")?;
write!(self.out, ".get_num_samples()")?;
}
},
crate::Expression::Unary { op, expr } => {

122
src/back/spv/image.rs
View File

@ -1032,21 +1032,19 @@ impl<'w> BlockContext<'w> {
Id::D2 | Id::Cube => 2,
Id::D3 => 3,
};
let extended_size_type_id = {
let array_coords = usize::from(arrayed);
let vector_size = match dim_coords + array_coords {
2 => Some(crate::VectorSize::Bi),
3 => Some(crate::VectorSize::Tri),
4 => Some(crate::VectorSize::Quad),
_ => None,
};
self.get_type_id(LookupType::Local(LocalType::Value {
vector_size,
kind: crate::ScalarKind::Sint,
width: 4,
pointer_space: None,
}))
let array_coords = usize::from(arrayed);
let vector_size = match dim_coords + array_coords {
2 => Some(crate::VectorSize::Bi),
3 => Some(crate::VectorSize::Tri),
4 => Some(crate::VectorSize::Quad),
_ => None,
};
let extended_size_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
vector_size,
kind: crate::ScalarKind::Sint,
width: 4,
pointer_space: None,
}));
let (query_op, level_id) = match class {
Ic::Sampled { multi: true, .. }
@ -1075,7 +1073,24 @@ impl<'w> BlockContext<'w> {
}
block.body.push(inst);
if result_type_id != extended_size_type_id {
let bitcast_type_id = self.get_type_id(
LocalType::Value {
vector_size,
kind: crate::ScalarKind::Uint,
width: 4,
pointer_space: None,
}
.into(),
);
let bitcast_id = self.gen_id();
block.body.push(Instruction::unary(
spirv::Op::Bitcast,
bitcast_type_id,
bitcast_id,
id_extended,
));
if result_type_id != bitcast_type_id {
let id = self.gen_id();
let components = match dim {
// always pick the first component, and duplicate it for all 3 dimensions
@ -1085,23 +1100,41 @@ impl<'w> BlockContext<'w> {
block.body.push(Instruction::vector_shuffle(
result_type_id,
id,
id_extended,
id_extended,
bitcast_id,
bitcast_id,
components,
));
id
} else {
id_extended
bitcast_id
}
}
Iq::NumLevels => {
let id = self.gen_id();
let query_id = self.gen_id();
block.body.push(Instruction::image_query(
spirv::Op::ImageQueryLevels,
result_type_id,
id,
self.get_type_id(
LocalType::Value {
vector_size: None,
kind: crate::ScalarKind::Sint,
width: 4,
pointer_space: None,
}
.into(),
),
query_id,
image_id,
));
let id = self.gen_id();
block.body.push(Instruction::unary(
spirv::Op::Bitcast,
result_type_id,
id,
query_id,
));
id
}
Iq::NumLayers => {
@ -1125,23 +1158,58 @@ impl<'w> BlockContext<'w> {
);
inst.add_operand(self.get_index_constant(0));
block.body.push(inst);
let id = self.gen_id();
let extract_id = self.gen_id();
block.body.push(Instruction::composite_extract(
result_type_id,
id,
self.get_type_id(
LocalType::Value {
vector_size: None,
kind: crate::ScalarKind::Sint,
width: 4,
pointer_space: None,
}
.into(),
),
extract_id,
id_extended,
&[vec_size as u32 - 1],
));
let id = self.gen_id();
block.body.push(Instruction::unary(
spirv::Op::Bitcast,
result_type_id,
id,
extract_id,
));
id
}
Iq::NumSamples => {
let id = self.gen_id();
let query_id = self.gen_id();
block.body.push(Instruction::image_query(
spirv::Op::ImageQuerySamples,
result_type_id,
id,
self.get_type_id(
LocalType::Value {
vector_size: None,
kind: crate::ScalarKind::Sint,
width: 4,
pointer_space: None,
}
.into(),
),
query_id,
image_id,
));
let id = self.gen_id();
block.body.push(Instruction::unary(
spirv::Op::Bitcast,
result_type_id,
id,
query_id,
));
id
}
};

12
src/front/glsl/builtins.rs
View File

@ -1873,7 +1873,7 @@ impl MacroCall {
name: None,
inner: TypeInner::Vector {
size,
kind: crate::ScalarKind::Sint,
kind: crate::ScalarKind::Uint,
width: 4,
},
},
@ -1883,7 +1883,15 @@ impl MacroCall {
expr = ctx.add_expression(Expression::Compose { components, ty }, meta, body)
}
expr
ctx.add_expression(
Expression::As {
expr,
kind: Sk::Sint,
convert: Some(4),
},
Span::default(),
body,
)
}
MacroCall::ImageLoad { multi } => {
let comps =

10
src/front/spv/image.rs
View File

@ -684,6 +684,11 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
image: image_lexp.handle,
query: crate::ImageQuery::Size { level },
};
let expr = crate::Expression::As {
expr: ctx.expressions.append(expr, self.span_from_with_op(start)),
kind: crate::ScalarKind::Sint,
convert: Some(4),
};
self.lookup_expression.insert(
result_id,
LookupExpression {
@ -714,6 +719,11 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
image: image_lexp.handle,
query,
};
let expr = crate::Expression::As {
expr: expressions.append(expr, self.span_from_with_op(start)),
kind: crate::ScalarKind::Sint,
convert: Some(4),
};
self.lookup_expression.insert(
result_id,
LookupExpression {

6
src/front/wgsl/tests.rs
View File

@ -357,10 +357,10 @@ fn parse_texture_query() {
"
var t: texture_multisampled_2d_array<f32>;
fn foo() {
var dim: vec2<i32> = textureDimensions(t);
var dim: vec2<u32> = textureDimensions(t);
dim = textureDimensions(t, 0);
let layers: i32 = textureNumLayers(t);
let samples: i32 = textureNumSamples(t);
let layers: u32 = textureNumLayers(t);
let samples: u32 = textureNumSamples(t);
}
",
)

8
src/proc/typifier.rs
View File

@ -525,17 +525,17 @@ impl<'a> ResolveContext<'a> {
crate::ImageQuery::Size { level: _ } => match *past(image)?.inner_with(types) {
Ti::Image { dim, .. } => match dim {
crate::ImageDimension::D1 => Ti::Scalar {
kind: crate::ScalarKind::Sint,
kind: crate::ScalarKind::Uint,
width: 4,
},
crate::ImageDimension::D2 | crate::ImageDimension::Cube => Ti::Vector {
size: crate::VectorSize::Bi,
kind: crate::ScalarKind::Sint,
kind: crate::ScalarKind::Uint,
width: 4,
},
crate::ImageDimension::D3 => Ti::Vector {
size: crate::VectorSize::Tri,
kind: crate::ScalarKind::Sint,
kind: crate::ScalarKind::Uint,
width: 4,
},
},
@ -547,7 +547,7 @@ impl<'a> ResolveContext<'a> {
crate::ImageQuery::NumLevels
| crate::ImageQuery::NumLayers
| crate::ImageQuery::NumSamples => Ti::Scalar {
kind: crate::ScalarKind::Sint,
kind: crate::ScalarKind::Uint,
width: 4,
},
}),

30
tests/in/binding-arrays.wgsl
View File

@ -30,8 +30,8 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
let uniform_index = uni.index;
let non_uniform_index = fragment_in.index;
var i1 = 0;
var i2 = vec2<i32>(0);
var u1 = 0u;
var u2 = vec2<u32>(0u);
var v1 = 0.0;
var v4 = vec4<f32>(0.0);
@ -46,9 +46,9 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
let uv = vec2<f32>(0.0);
let pix = vec2<i32>(0);
i2 += textureDimensions(texture_array_unbounded[0]);
i2 += textureDimensions(texture_array_unbounded[uniform_index]);
i2 += textureDimensions(texture_array_unbounded[non_uniform_index]);
u2 += textureDimensions(texture_array_unbounded[0]);
u2 += textureDimensions(texture_array_unbounded[uniform_index]);
u2 += textureDimensions(texture_array_unbounded[non_uniform_index]);
v4 += textureGather(0, texture_array_bounded[0], samp[0], uv);
v4 += textureGather(0, texture_array_bounded[uniform_index], samp[uniform_index], uv);
@ -62,17 +62,17 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
v4 += textureLoad(texture_array_unbounded[uniform_index], pix, 0);
v4 += textureLoad(texture_array_unbounded[non_uniform_index], pix, 0);
i1 += textureNumLayers(texture_array_2darray[0]);
i1 += textureNumLayers(texture_array_2darray[uniform_index]);
i1 += textureNumLayers(texture_array_2darray[non_uniform_index]);
u1 += textureNumLayers(texture_array_2darray[0]);
u1 += textureNumLayers(texture_array_2darray[uniform_index]);
u1 += textureNumLayers(texture_array_2darray[non_uniform_index]);
i1 += textureNumLevels(texture_array_bounded[0]);
i1 += textureNumLevels(texture_array_bounded[uniform_index]);
i1 += textureNumLevels(texture_array_bounded[non_uniform_index]);
u1 += textureNumLevels(texture_array_bounded[0]);
u1 += textureNumLevels(texture_array_bounded[uniform_index]);
u1 += textureNumLevels(texture_array_bounded[non_uniform_index]);
i1 += textureNumSamples(texture_array_multisampled[0]);
i1 += textureNumSamples(texture_array_multisampled[uniform_index]);
i1 += textureNumSamples(texture_array_multisampled[non_uniform_index]);
u1 += textureNumSamples(texture_array_multisampled[0]);
u1 += textureNumSamples(texture_array_multisampled[uniform_index]);
u1 += textureNumSamples(texture_array_multisampled[non_uniform_index]);
v4 += textureSample(texture_array_bounded[0], samp[0], uv);
v4 += textureSample(texture_array_bounded[uniform_index], samp[uniform_index], uv);
@ -102,7 +102,7 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
textureStore(texture_array_storage[uniform_index], pix, v4);
textureStore(texture_array_storage[non_uniform_index], pix, v4);
let v2 = vec2<f32>(i2 + vec2<i32>(i1));
let v2 = vec2<f32>(u2 + vec2<u32>(u1));
return v4 + vec4<f32>(v2.x, v2.y, v2.x, v2.y) + v1;
}

6
tests/in/image.wgsl
View File

@ -18,7 +18,7 @@ var image_dst: texture_storage_1d<r32uint,write>;
@compute @workgroup_size(16)
fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
let dim = textureDimensions(image_storage_src);
let itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20);
let itc = vec2<i32>(dim * local_id.xy) % vec2<i32>(10, 20);
// loads with ivec2 coords.
let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z));
@ -39,8 +39,8 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
@compute @workgroup_size(16, 1, 1)
fn depth_load(@builtin(local_invocation_id) local_id: vec3<u32>) {
let dim: vec2<i32> = textureDimensions(image_storage_src);
let itc: vec2<i32> = ((dim * vec2<i32>(local_id.xy)) % vec2<i32>(10, 20));
let dim: vec2<u32> = textureDimensions(image_storage_src);
let itc: vec2<i32> = (vec2<i32>(dim * local_id.xy) % vec2<i32>(10, 20));
let val: f32 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z));
textureStore(image_dst, itc.x, vec4<u32>(u32(val)));
return;

4
tests/out/glsl/image.main.Compute.glsl
View File

@ -21,8 +21,8 @@ layout(r32ui) writeonly uniform highp uimage2D _group_0_binding_2_cs;
void main() {
uvec3 local_id = gl_LocalInvocationID;
ivec2 dim = imageSize(_group_0_binding_1_cs).xy;
ivec2 itc = ((dim * ivec2(local_id.xy)) % ivec2(10, 20));
uvec2 dim = uvec2(imageSize(_group_0_binding_1_cs).xy);
ivec2 itc = (ivec2((dim * local_id.xy)) % ivec2(10, 20));
uvec4 value1_ = texelFetch(_group_0_binding_0_cs, itc, int(local_id.z));
uvec4 value2_ = texelFetch(_group_0_binding_3_cs, itc, int(local_id.z));
uvec4 value4_ = imageLoad(_group_0_binding_1_cs, itc);

28
tests/out/glsl/image.queries.Vertex.glsl
View File

@ -20,20 +20,20 @@ uniform highp sampler2DMS _group_0_binding_8_vs;
void main() {
int dim_1d = textureSize(_group_0_binding_0_vs, 0).x;
int dim_1d_lod = textureSize(_group_0_binding_0_vs, int(dim_1d)).x;
ivec2 dim_2d = textureSize(_group_0_binding_1_vs, 0).xy;
ivec2 dim_2d_lod = textureSize(_group_0_binding_1_vs, 1).xy;
ivec2 dim_2d_array = textureSize(_group_0_binding_4_vs, 0).xy;
ivec2 dim_2d_array_lod = textureSize(_group_0_binding_4_vs, 1).xy;
ivec2 dim_cube = textureSize(_group_0_binding_5_vs, 0).xy;
ivec2 dim_cube_lod = textureSize(_group_0_binding_5_vs, 1).xy;
ivec2 dim_cube_array = textureSize(_group_0_binding_6_vs, 0).xy;
ivec2 dim_cube_array_lod = textureSize(_group_0_binding_6_vs, 1).xy;
ivec3 dim_3d = textureSize(_group_0_binding_7_vs, 0).xyz;
ivec3 dim_3d_lod = textureSize(_group_0_binding_7_vs, 1).xyz;
ivec2 dim_2s_ms = textureSize(_group_0_binding_8_vs).xy;
int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
uint dim_1d = uint(textureSize(_group_0_binding_0_vs, 0).x);
uint dim_1d_lod = uint(textureSize(_group_0_binding_0_vs, int(dim_1d)).x);
uvec2 dim_2d = uvec2(textureSize(_group_0_binding_1_vs, 0).xy);
uvec2 dim_2d_lod = uvec2(textureSize(_group_0_binding_1_vs, 1).xy);
uvec2 dim_2d_array = uvec2(textureSize(_group_0_binding_4_vs, 0).xy);
uvec2 dim_2d_array_lod = uvec2(textureSize(_group_0_binding_4_vs, 1).xy);
uvec2 dim_cube = uvec2(textureSize(_group_0_binding_5_vs, 0).xy);
uvec2 dim_cube_lod = uvec2(textureSize(_group_0_binding_5_vs, 1).xy);
uvec2 dim_cube_array = uvec2(textureSize(_group_0_binding_6_vs, 0).xy);
uvec2 dim_cube_array_lod = uvec2(textureSize(_group_0_binding_6_vs, 1).xy);
uvec3 dim_3d = uvec3(textureSize(_group_0_binding_7_vs, 0).xyz);
uvec3 dim_3d_lod = uvec3(textureSize(_group_0_binding_7_vs, 1).xyz);
uvec2 dim_2s_ms = uvec2(textureSize(_group_0_binding_8_vs).xy);
uint sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
gl_Position = vec4(float(sum));
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;

2
tests/out/glsl/variations-glsl.main.Fragment.glsl
View File

@ -9,7 +9,7 @@ uniform highp samplerCube _group_0_binding_0_fs;
void main_1() {
ivec2 sizeCube = ivec2(0);
float a = 0.0;
sizeCube = textureSize(_group_0_binding_0_fs, 0).xy;
sizeCube = ivec2(uvec2(textureSize(_group_0_binding_0_fs, 0).xy));
a = ceil(1.0);
return;
}

68
tests/out/hlsl/binding-arrays.hlsl
View File

@ -21,28 +21,28 @@ struct FragmentInput_main {
nointerpolation uint index : LOC0;
};
int2 NagaDimensions2D(Texture2D<float4> tex)
uint2 NagaDimensions2D(Texture2D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy;
}
int NagaNumLayers2DArray(Texture2DArray<float4> tex)
uint NagaNumLayers2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLevels2D(Texture2D<float4> tex)
uint NagaNumLevels2D(Texture2D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int NagaMSNumSamples2D(Texture2DMS<float4> tex)
uint NagaMSNumSamples2D(Texture2DMS<float4> tex)
{
uint4 ret;
tex.GetDimensions(ret.x, ret.y, ret.z);
@ -52,25 +52,25 @@ int NagaMSNumSamples2D(Texture2DMS<float4> tex)
float4 main(FragmentInput_main fragmentinput_main) : SV_Target0
{
FragmentIn fragment_in = { fragmentinput_main.index };
int i1_ = (int)0;
int2 i2_ = (int2)0;
uint u1_ = (uint)0;
uint2 u2_ = (uint2)0;
float v1_ = (float)0;
float4 v4_ = (float4)0;
uint uniform_index = uni.index;
uint non_uniform_index = fragment_in.index;
i1_ = 0;
i2_ = (0).xx;
u1_ = 0u;
u2_ = (0u).xx;
v1_ = 0.0;
v4_ = (0.0).xxxx;
float2 uv = (0.0).xx;
int2 pix = (0).xx;
int2 _expr23 = i2_;
i2_ = (_expr23 + NagaDimensions2D(texture_array_unbounded[0]));
int2 _expr28 = i2_;
i2_ = (_expr28 + NagaDimensions2D(texture_array_unbounded[uniform_index]));
int2 _expr33 = i2_;
i2_ = (_expr33 + NagaDimensions2D(texture_array_unbounded[NonUniformResourceIndex(non_uniform_index)]));
uint2 _expr23 = u2_;
u2_ = (_expr23 + NagaDimensions2D(texture_array_unbounded[0]));
uint2 _expr28 = u2_;
u2_ = (_expr28 + NagaDimensions2D(texture_array_unbounded[uniform_index]));
uint2 _expr33 = u2_;
u2_ = (_expr33 + NagaDimensions2D(texture_array_unbounded[NonUniformResourceIndex(non_uniform_index)]));
float4 _expr41 = texture_array_bounded[0].Gather(samp[0], uv);
float4 _expr42 = v4_;
v4_ = (_expr42 + _expr41);
@ -98,24 +98,24 @@ float4 main(FragmentInput_main fragmentinput_main) : SV_Target0
float4 _expr100 = texture_array_unbounded[NonUniformResourceIndex(non_uniform_index)].Load(int3(pix, 0));
float4 _expr101 = v4_;
v4_ = (_expr101 + _expr100);
int _expr107 = i1_;
i1_ = (_expr107 + NagaNumLayers2DArray(texture_array_2darray[0]));
int _expr112 = i1_;
i1_ = (_expr112 + NagaNumLayers2DArray(texture_array_2darray[uniform_index]));
int _expr117 = i1_;
i1_ = (_expr117 + NagaNumLayers2DArray(texture_array_2darray[NonUniformResourceIndex(non_uniform_index)]));
int _expr123 = i1_;
i1_ = (_expr123 + NagaNumLevels2D(texture_array_bounded[0]));
int _expr128 = i1_;
i1_ = (_expr128 + NagaNumLevels2D(texture_array_bounded[uniform_index]));
int _expr133 = i1_;
i1_ = (_expr133 + NagaNumLevels2D(texture_array_bounded[NonUniformResourceIndex(non_uniform_index)]));
int _expr139 = i1_;
i1_ = (_expr139 + NagaMSNumSamples2D(texture_array_multisampled[0]));
int _expr144 = i1_;
i1_ = (_expr144 + NagaMSNumSamples2D(texture_array_multisampled[uniform_index]));
int _expr149 = i1_;
i1_ = (_expr149 + NagaMSNumSamples2D(texture_array_multisampled[NonUniformResourceIndex(non_uniform_index)]));
uint _expr107 = u1_;
u1_ = (_expr107 + NagaNumLayers2DArray(texture_array_2darray[0]));
uint _expr112 = u1_;
u1_ = (_expr112 + NagaNumLayers2DArray(texture_array_2darray[uniform_index]));
uint _expr117 = u1_;
u1_ = (_expr117 + NagaNumLayers2DArray(texture_array_2darray[NonUniformResourceIndex(non_uniform_index)]));
uint _expr123 = u1_;
u1_ = (_expr123 + NagaNumLevels2D(texture_array_bounded[0]));
uint _expr128 = u1_;
u1_ = (_expr128 + NagaNumLevels2D(texture_array_bounded[uniform_index]));
uint _expr133 = u1_;
u1_ = (_expr133 + NagaNumLevels2D(texture_array_bounded[NonUniformResourceIndex(non_uniform_index)]));
uint _expr139 = u1_;
u1_ = (_expr139 + NagaMSNumSamples2D(texture_array_multisampled[0]));
uint _expr144 = u1_;
u1_ = (_expr144 + NagaMSNumSamples2D(texture_array_multisampled[uniform_index]));
uint _expr149 = u1_;
u1_ = (_expr149 + NagaMSNumSamples2D(texture_array_multisampled[NonUniformResourceIndex(non_uniform_index)]));
float4 _expr157 = texture_array_bounded[0].Sample(samp[0], uv);
float4 _expr158 = v4_;
v4_ = (_expr158 + _expr157);
@ -176,8 +176,8 @@ float4 main(FragmentInput_main fragmentinput_main) : SV_Target0
texture_array_storage[uniform_index][pix] = _expr307;
float4 _expr310 = v4_;
texture_array_storage[NonUniformResourceIndex(non_uniform_index)][pix] = _expr310;
int2 _expr311 = i2_;
int _expr312 = i1_;
uint2 _expr311 = u2_;
uint _expr312 = u1_;
float2 v2_ = float2((_expr311 + (_expr312).xx));
float4 _expr316 = v4_;
float _expr323 = v1_;

98
tests/out/hlsl/image.hlsl
View File

@ -21,7 +21,7 @@ SamplerComparisonState sampler_cmp : register(s1, space1);
Texture2D<float> image_2d_depth : register(t2, space1);
TextureCube<float> image_cube_depth : register(t3, space1);
int2 NagaRWDimensions2D(RWTexture2D<uint4> tex)
uint2 NagaRWDimensions2D(RWTexture2D<uint4> tex)
{
uint4 ret;
tex.GetDimensions(ret.x, ret.y);
@ -31,8 +31,8 @@ int2 NagaRWDimensions2D(RWTexture2D<uint4> tex)
[numthreads(16, 1, 1)]
void main(uint3 local_id : SV_GroupThreadID)
{
int2 dim = NagaRWDimensions2D(image_storage_src);
int2 itc = ((dim * int2(local_id.xy)) % int2(10, 20));
uint2 dim = NagaRWDimensions2D(image_storage_src);
int2 itc = (int2((dim * local_id.xy)) % int2(10, 20));
uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(local_id.z)));
uint4 value2_ = image_multisampled_src.Load(itc, int(local_id.z));
uint4 value4_ = image_storage_src.Load(itc);
@ -51,98 +51,98 @@ void main(uint3 local_id : SV_GroupThreadID)
[numthreads(16, 1, 1)]
void depth_load(uint3 local_id_1 : SV_GroupThreadID)
{
int2 dim_1 = NagaRWDimensions2D(image_storage_src);
int2 itc_1 = ((dim_1 * int2(local_id_1.xy)) % int2(10, 20));
uint2 dim_1 = NagaRWDimensions2D(image_storage_src);
int2 itc_1 = (int2((dim_1 * local_id_1.xy)) % int2(10, 20));
float val = image_depth_multisampled_src.Load(itc_1, int(local_id_1.z)).x;
image_dst[itc_1.x] = (uint(val)).xxxx;
return;
}
int NagaDimensions1D(Texture1D<float4> tex)
uint NagaDimensions1D(Texture1D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y);
return ret.x;
}
int NagaMipDimensions1D(Texture1D<float4> tex, uint mip_level)
uint NagaMipDimensions1D(Texture1D<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y);
return ret.x;
}
int2 NagaDimensions2D(Texture2D<float4> tex)
uint2 NagaDimensions2D(Texture2D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy;
}
int2 NagaMipDimensions2D(Texture2D<float4> tex, uint mip_level)
uint2 NagaMipDimensions2D(Texture2D<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z);
return ret.xy;
}
int2 NagaDimensions2DArray(Texture2DArray<float4> tex)
uint2 NagaDimensions2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int2 NagaMipDimensions2DArray(Texture2DArray<float4> tex, uint mip_level)
uint2 NagaMipDimensions2DArray(Texture2DArray<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int2 NagaDimensionsCube(TextureCube<float4> tex)
uint2 NagaDimensionsCube(TextureCube<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy;
}
int2 NagaMipDimensionsCube(TextureCube<float4> tex, uint mip_level)
uint2 NagaMipDimensionsCube(TextureCube<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z);
return ret.xy;
}
int2 NagaDimensionsCubeArray(TextureCubeArray<float4> tex)
uint2 NagaDimensionsCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int2 NagaMipDimensionsCubeArray(TextureCubeArray<float4> tex, uint mip_level)
uint2 NagaMipDimensionsCubeArray(TextureCubeArray<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int3 NagaDimensions3D(Texture3D<float4> tex)
uint3 NagaDimensions3D(Texture3D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.xyz;
}
int3 NagaMipDimensions3D(Texture3D<float4> tex, uint mip_level)
uint3 NagaMipDimensions3D(Texture3D<float4> tex, uint mip_level)
{
uint4 ret;
tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xyz;
}
int2 NagaMSDimensions2D(Texture2DMS<float4> tex)
uint2 NagaMSDimensions2D(Texture2DMS<float4> tex)
{
uint4 ret;
tex.GetDimensions(ret.x, ret.y, ret.z);
@ -151,73 +151,73 @@ int2 NagaMSDimensions2D(Texture2DMS<float4> tex)
float4 queries() : SV_Position
{
int dim_1d = NagaDimensions1D(image_1d);
int dim_1d_lod = NagaMipDimensions1D(image_1d, int(dim_1d));
int2 dim_2d = NagaDimensions2D(image_2d);
int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1);
int2 dim_2d_array = NagaDimensions2DArray(image_2d_array);
int2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1);
int2 dim_cube = NagaDimensionsCube(image_cube);
int2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1);
int2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array);
int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1);
int3 dim_3d = NagaDimensions3D(image_3d);
int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1);
int2 dim_2s_ms = NagaMSDimensions2D(image_aa);
int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
uint dim_1d = NagaDimensions1D(image_1d);
uint dim_1d_lod = NagaMipDimensions1D(image_1d, int(dim_1d));
uint2 dim_2d = NagaDimensions2D(image_2d);
uint2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1);
uint2 dim_2d_array = NagaDimensions2DArray(image_2d_array);
uint2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1);
uint2 dim_cube = NagaDimensionsCube(image_cube);
uint2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1);
uint2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array);
uint2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1);
uint3 dim_3d = NagaDimensions3D(image_3d);
uint3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1);
uint2 dim_2s_ms = NagaMSDimensions2D(image_aa);
uint sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
return (float(sum)).xxxx;
}
int NagaNumLevels2D(Texture2D<float4> tex)
uint NagaNumLevels2D(Texture2D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int NagaNumLevels2DArray(Texture2DArray<float4> tex)
uint NagaNumLevels2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLayers2DArray(Texture2DArray<float4> tex)
uint NagaNumLayers2DArray(Texture2DArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLevelsCube(TextureCube<float4> tex)
uint NagaNumLevelsCube(TextureCube<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int NagaNumLevelsCubeArray(TextureCubeArray<float4> tex)
uint NagaNumLevelsCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLayersCubeArray(TextureCubeArray<float4> tex)
uint NagaNumLayersCubeArray(TextureCubeArray<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaNumLevels3D(Texture3D<float4> tex)
uint NagaNumLevels3D(Texture3D<float4> tex)
{
uint4 ret;
tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int NagaMSNumSamples2D(Texture2DMS<float4> tex)
uint NagaMSNumSamples2D(Texture2DMS<float4> tex)
{
uint4 ret;
tex.GetDimensions(ret.x, ret.y, ret.z);
@ -226,15 +226,15 @@ int NagaMSNumSamples2D(Texture2DMS<float4> tex)
float4 levels_queries() : SV_Position
{
int num_levels_2d = NagaNumLevels2D(image_2d);
int num_levels_2d_array = NagaNumLevels2DArray(image_2d_array);
int num_layers_2d = NagaNumLayers2DArray(image_2d_array);
int num_levels_cube = NagaNumLevelsCube(image_cube);
int num_levels_cube_array = NagaNumLevelsCubeArray(image_cube_array);
int num_layers_cube = NagaNumLayersCubeArray(image_cube_array);
int num_levels_3d = NagaNumLevels3D(image_3d);
int num_samples_aa = NagaMSNumSamples2D(image_aa);
int sum_1 = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
uint num_levels_2d = NagaNumLevels2D(image_2d);
uint num_levels_2d_array = NagaNumLevels2DArray(image_2d_array);
uint num_layers_2d = NagaNumLayers2DArray(image_2d_array);
uint num_levels_cube = NagaNumLevelsCube(image_cube);
uint num_levels_cube_array = NagaNumLevelsCubeArray(image_cube_array);
uint num_layers_cube = NagaNumLayersCubeArray(image_cube_array);
uint num_levels_3d = NagaNumLevels3D(image_3d);
uint num_samples_aa = NagaMSNumSamples2D(image_aa);
uint sum_1 = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array);
return (float(sum_1)).xxxx;
}

62
tests/out/msl/binding-arrays.msl
View File

@ -36,24 +36,24 @@ fragment main_Output main_(
, constant UniformIndex& uni [[user(fake0)]]
) {
const FragmentIn fragment_in = { varyings.index };
int i1_ = {};
metal::int2 i2_ = {};
uint u1_ = {};
metal::uint2 u2_ = {};
float v1_ = {};
metal::float4 v4_ = {};
uint uniform_index = uni.index;
uint non_uniform_index = fragment_in.index;
i1_ = 0;
i2_ = metal::int2(0);
u1_ = 0u;
u2_ = metal::uint2(0u);
v1_ = 0.0;
v4_ = metal::float4(0.0);
metal::float2 uv = metal::float2(0.0);
metal::int2 pix = metal::int2(0);
metal::int2 _e23 = i2_;
i2_ = _e23 + metal::int2(texture_array_unbounded[0].get_width(), texture_array_unbounded[0].get_height());
metal::int2 _e28 = i2_;
i2_ = _e28 + metal::int2(texture_array_unbounded[uniform_index].get_width(), texture_array_unbounded[uniform_index].get_height());
metal::int2 _e33 = i2_;
i2_ = _e33 + metal::int2(texture_array_unbounded[non_uniform_index].get_width(), texture_array_unbounded[non_uniform_index].get_height());
metal::uint2 _e23 = u2_;
u2_ = _e23 + metal::uint2(texture_array_unbounded[0].get_width(), texture_array_unbounded[0].get_height());
metal::uint2 _e28 = u2_;
u2_ = _e28 + metal::uint2(texture_array_unbounded[uniform_index].get_width(), texture_array_unbounded[uniform_index].get_height());
metal::uint2 _e33 = u2_;
u2_ = _e33 + metal::uint2(texture_array_unbounded[non_uniform_index].get_width(), texture_array_unbounded[non_uniform_index].get_height());
metal::float4 _e41 = texture_array_bounded[0].gather(samp[0], uv);
metal::float4 _e42 = v4_;
v4_ = _e42 + _e41;
@ -81,24 +81,24 @@ fragment main_Output main_(
metal::float4 _e100 = (uint(0) < texture_array_unbounded[non_uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[non_uniform_index].get_width(0), texture_array_unbounded[non_uniform_index].get_height(0))) ? texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e101 = v4_;
v4_ = _e101 + _e100;
int _e107 = i1_;
i1_ = _e107 + int(texture_array_2darray[0].get_array_size());
int _e112 = i1_;
i1_ = _e112 + int(texture_array_2darray[uniform_index].get_array_size());
int _e117 = i1_;
i1_ = _e117 + int(texture_array_2darray[non_uniform_index].get_array_size());
int _e123 = i1_;
i1_ = _e123 + int(texture_array_bounded[0].get_num_mip_levels());
int _e128 = i1_;
i1_ = _e128 + int(texture_array_bounded[uniform_index].get_num_mip_levels());
int _e133 = i1_;
i1_ = _e133 + int(texture_array_bounded[non_uniform_index].get_num_mip_levels());
int _e139 = i1_;
i1_ = _e139 + int(texture_array_multisampled[0].get_num_samples());
int _e144 = i1_;
i1_ = _e144 + int(texture_array_multisampled[uniform_index].get_num_samples());
int _e149 = i1_;
i1_ = _e149 + int(texture_array_multisampled[non_uniform_index].get_num_samples());
uint _e107 = u1_;
u1_ = _e107 + texture_array_2darray[0].get_array_size();
uint _e112 = u1_;
u1_ = _e112 + texture_array_2darray[uniform_index].get_array_size();
uint _e117 = u1_;
u1_ = _e117 + texture_array_2darray[non_uniform_index].get_array_size();
uint _e123 = u1_;
u1_ = _e123 + texture_array_bounded[0].get_num_mip_levels();
uint _e128 = u1_;
u1_ = _e128 + texture_array_bounded[uniform_index].get_num_mip_levels();
uint _e133 = u1_;
u1_ = _e133 + texture_array_bounded[non_uniform_index].get_num_mip_levels();
uint _e139 = u1_;
u1_ = _e139 + texture_array_multisampled[0].get_num_samples();
uint _e144 = u1_;
u1_ = _e144 + texture_array_multisampled[uniform_index].get_num_samples();
uint _e149 = u1_;
u1_ = _e149 + texture_array_multisampled[non_uniform_index].get_num_samples();
metal::float4 _e157 = texture_array_bounded[0].sample(samp[0], uv);
metal::float4 _e158 = v4_;
v4_ = _e158 + _e157;
@ -165,9 +165,9 @@ fragment main_Output main_(
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[non_uniform_index].get_width(), texture_array_storage[non_uniform_index].get_height()))) {
texture_array_storage[non_uniform_index].write(_e310, metal::uint2(pix));
}
metal::int2 _e311 = i2_;
int _e312 = i1_;
metal::float2 v2_ = static_cast<metal::float2>(_e311 + metal::int2(_e312));
metal::uint2 _e311 = u2_;
uint _e312 = u1_;
metal::float2 v2_ = static_cast<metal::float2>(_e311 + metal::uint2(_e312));
metal::float4 _e316 = v4_;
float _e323 = v1_;
return main_Output { (_e316 + metal::float4(v2_.x, v2_.y, v2_.x, v2_.y)) + metal::float4(_e323) };

54
tests/out/msl/image.msl
View File

@ -17,8 +17,8 @@ kernel void main_(
, metal::texture1d<uint, metal::access::sample> image_1d_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
) {
metal::int2 dim = metal::int2(image_storage_src.get_width(), image_storage_src.get_height());
metal::int2 itc = (dim * static_cast<metal::int2>(local_id.xy)) % metal::int2(10, 20);
metal::uint2 dim = metal::uint2(image_storage_src.get_width(), image_storage_src.get_height());
metal::int2 itc = static_cast<metal::int2>(dim * local_id.xy) % metal::int2(10, 20);
metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc));
@ -43,8 +43,8 @@ kernel void depth_load(
, metal::texture2d<uint, metal::access::read> image_storage_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
) {
metal::int2 dim_1 = metal::int2(image_storage_src.get_width(), image_storage_src.get_height());
metal::int2 itc_1 = (dim_1 * static_cast<metal::int2>(local_id_1.xy)) % metal::int2(10, 20);
metal::uint2 dim_1 = metal::uint2(image_storage_src.get_width(), image_storage_src.get_height());
metal::int2 itc_1 = static_cast<metal::int2>(dim_1 * local_id_1.xy) % metal::int2(10, 20);
float val = image_depth_multisampled_src.read(metal::uint2(itc_1), static_cast<int>(local_id_1.z));
image_dst.write(metal::uint4(static_cast<uint>(val)), uint(itc_1.x));
return;
@ -63,20 +63,20 @@ vertex queriesOutput queries(
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
, metal::texture2d_ms<float, metal::access::read> image_aa [[user(fake0)]]
) {
int dim_1d = int(image_1d.get_width());
int dim_1d_lod = int(image_1d.get_width());
metal::int2 dim_2d = metal::int2(image_2d.get_width(), image_2d.get_height());
metal::int2 dim_2d_lod = metal::int2(image_2d.get_width(1), image_2d.get_height(1));
metal::int2 dim_2d_array = metal::int2(image_2d_array.get_width(), image_2d_array.get_height());
metal::int2 dim_2d_array_lod = metal::int2(image_2d_array.get_width(1), image_2d_array.get_height(1));
metal::int2 dim_cube = metal::int2(image_cube.get_width());
metal::int2 dim_cube_lod = metal::int2(image_cube.get_width(1));
metal::int2 dim_cube_array = metal::int2(image_cube_array.get_width());
metal::int2 dim_cube_array_lod = metal::int2(image_cube_array.get_width(1));
metal::int3 dim_3d = metal::int3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth());
metal::int3 dim_3d_lod = metal::int3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1));
metal::int2 dim_2s_ms = metal::int2(image_aa.get_width(), image_aa.get_height());
int sum = (((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z;
uint dim_1d = image_1d.get_width();
uint dim_1d_lod = image_1d.get_width();
metal::uint2 dim_2d = metal::uint2(image_2d.get_width(), image_2d.get_height());
metal::uint2 dim_2d_lod = metal::uint2(image_2d.get_width(1), image_2d.get_height(1));
metal::uint2 dim_2d_array = metal::uint2(image_2d_array.get_width(), image_2d_array.get_height());
metal::uint2 dim_2d_array_lod = metal::uint2(image_2d_array.get_width(1), image_2d_array.get_height(1));
metal::uint2 dim_cube = metal::uint2(image_cube.get_width());
metal::uint2 dim_cube_lod = metal::uint2(image_cube.get_width(1));
metal::uint2 dim_cube_array = metal::uint2(image_cube_array.get_width());
metal::uint2 dim_cube_array_lod = metal::uint2(image_cube_array.get_width(1));
metal::uint3 dim_3d = metal::uint3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth());
metal::uint3 dim_3d_lod = metal::uint3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1));
metal::uint2 dim_2s_ms = metal::uint2(image_aa.get_width(), image_aa.get_height());
uint sum = (((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z;
return queriesOutput { metal::float4(static_cast<float>(sum)) };
}
@ -92,15 +92,15 @@ vertex levels_queriesOutput levels_queries(
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
, metal::texture2d_ms<float, metal::access::read> image_aa [[user(fake0)]]
) {
int num_levels_2d = int(image_2d.get_num_mip_levels());
int num_levels_2d_array = int(image_2d_array.get_num_mip_levels());
int num_layers_2d = int(image_2d_array.get_array_size());
int num_levels_cube = int(image_cube.get_num_mip_levels());
int num_levels_cube_array = int(image_cube_array.get_num_mip_levels());
int num_layers_cube = int(image_cube_array.get_array_size());
int num_levels_3d = int(image_3d.get_num_mip_levels());
int num_samples_aa = int(image_aa.get_num_samples());
int sum_1 = ((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array;
uint num_levels_2d = image_2d.get_num_mip_levels();
uint num_levels_2d_array = image_2d_array.get_num_mip_levels();
uint num_layers_2d = image_2d_array.get_array_size();
uint num_levels_cube = image_cube.get_num_mip_levels();
uint num_levels_cube_array = image_cube_array.get_num_mip_levels();
uint num_layers_cube = image_cube_array.get_array_size();
uint num_levels_3d = image_3d.get_num_mip_levels();
uint num_samples_aa = image_aa.get_num_samples();
uint sum_1 = ((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array;
return levels_queriesOutput { metal::float4(static_cast<float>(sum_1)) };
}

1073
tests/out/spv/binding-arrays.spvasm
View File

File diff suppressed because it is too large Load Diff

626
tests/out/spv/image.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 353
; Bound: 376
OpCapability SampledCubeArray
OpCapability ImageQuery
OpCapability Image1D
@ -10,19 +10,19 @@ OpCapability Sampled1D
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %84 "main" %81
OpEntryPoint GLCompute %157 "depth_load" %155
OpEntryPoint Vertex %178 "queries" %176
OpEntryPoint Vertex %230 "levels_queries" %229
OpEntryPoint Fragment %259 "texture_sample" %258
OpEntryPoint Fragment %287 "texture_sample_comparison" %285
OpEntryPoint Fragment %306 "gather" %305
OpEntryPoint Fragment %341 "depth_no_comparison" %340
OpEntryPoint GLCompute %158 "depth_load" %156
OpEntryPoint Vertex %180 "queries" %178
OpEntryPoint Vertex %245 "levels_queries" %244
OpEntryPoint Fragment %282 "texture_sample" %281
OpEntryPoint Fragment %310 "texture_sample_comparison" %308
OpEntryPoint Fragment %329 "gather" %328
OpEntryPoint Fragment %364 "depth_no_comparison" %363
OpExecutionMode %84 LocalSize 16 1 1
OpExecutionMode %157 LocalSize 16 1 1
OpExecutionMode %259 OriginUpperLeft
OpExecutionMode %287 OriginUpperLeft
OpExecutionMode %306 OriginUpperLeft
OpExecutionMode %341 OriginUpperLeft
OpExecutionMode %158 LocalSize 16 1 1
OpExecutionMode %282 OriginUpperLeft
OpExecutionMode %310 OriginUpperLeft
OpExecutionMode %329 OriginUpperLeft
OpExecutionMode %364 OriginUpperLeft
OpSource GLSL 450
OpName %39 "image_mipmapped_src"
OpName %41 "image_multisampled_src"
@ -47,14 +47,14 @@ OpName %76 "image_2d_depth"
OpName %78 "image_cube_depth"
OpName %81 "local_id"
OpName %84 "main"
OpName %155 "local_id"
OpName %157 "depth_load"
OpName %178 "queries"
OpName %230 "levels_queries"
OpName %259 "texture_sample"
OpName %287 "texture_sample_comparison"
OpName %306 "gather"
OpName %341 "depth_no_comparison"
OpName %156 "local_id"
OpName %158 "depth_load"
OpName %180 "queries"
OpName %245 "levels_queries"
OpName %282 "texture_sample"
OpName %310 "texture_sample_comparison"
OpName %329 "gather"
OpName %364 "depth_no_comparison"
OpDecorate %39 DescriptorSet 0
OpDecorate %39 Binding 0
OpDecorate %41 DescriptorSet 0
@ -101,13 +101,13 @@ OpDecorate %76 Binding 2
OpDecorate %78 DescriptorSet 1
OpDecorate %78 Binding 3
OpDecorate %81 BuiltIn LocalInvocationId
OpDecorate %155 BuiltIn LocalInvocationId
OpDecorate %176 BuiltIn Position
OpDecorate %229 BuiltIn Position
OpDecorate %258 Location 0
OpDecorate %285 Location 0
OpDecorate %305 Location 0
OpDecorate %340 Location 0
OpDecorate %156 BuiltIn LocalInvocationId
OpDecorate %178 BuiltIn Position
OpDecorate %244 BuiltIn Position
OpDecorate %281 Location 0
OpDecorate %308 Location 0
OpDecorate %328 Location 0
OpDecorate %363 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10
@ -189,27 +189,27 @@ OpDecorate %340 Location 0
%82 = OpTypePointer Input %20
%81 = OpVariable %82 Input
%85 = OpTypeFunction %2
%111 = OpTypeVector %4 3
%155 = OpVariable %82 Input
%177 = OpTypePointer Output %32
%176 = OpVariable %177 Output
%187 = OpConstant %13 0
%229 = OpVariable %177 Output
%258 = OpVariable %177 Output
%266 = OpTypeSampledImage %24
%269 = OpTypeSampledImage %25
%286 = OpTypePointer Output %8
%285 = OpVariable %286 Output
%293 = OpTypeSampledImage %35
%298 = OpConstant %8 0.0
%300 = OpTypeSampledImage %36
%305 = OpVariable %177 Output
%317 = OpConstant %13 1
%320 = OpConstant %13 3
%325 = OpTypeSampledImage %12
%328 = OpTypeVector %4 4
%329 = OpTypeSampledImage %26
%340 = OpVariable %177 Output
%112 = OpTypeVector %4 3
%156 = OpVariable %82 Input
%179 = OpTypePointer Output %32
%178 = OpVariable %179 Output
%189 = OpConstant %13 0
%244 = OpVariable %179 Output
%281 = OpVariable %179 Output
%289 = OpTypeSampledImage %24
%292 = OpTypeSampledImage %25
%309 = OpTypePointer Output %8
%308 = OpVariable %309 Output
%316 = OpTypeSampledImage %35
%321 = OpConstant %8 0.0
%323 = OpTypeSampledImage %36
%328 = OpVariable %179 Output
%340 = OpConstant %13 1
%343 = OpConstant %13 3
%348 = OpTypeSampledImage %12
%351 = OpTypeVector %4 4
%352 = OpTypeSampledImage %26
%363 = OpVariable %179 Output
%84 = OpFunction %2 None %85
%80 = OpLabel
%83 = OpLoad %20 %81
@ -222,271 +222,295 @@ OpDecorate %340 Location 0
OpBranch %92
%92 = OpLabel
%93 = OpImageQuerySize %21 %88
%94 = OpVectorShuffle %22 %83 %83 0 1
%95 = OpBitcast %21 %94
%96 = OpIMul %21 %93 %95
%97 = OpCompositeConstruct %21 %3 %5
%98 = OpSRem %21 %96 %97
%99 = OpCompositeExtract %13 %83 2
%100 = OpBitcast %4 %99
%101 = OpImageFetch %23 %86 %98 Lod %100
%102 = OpCompositeExtract %13 %83 2
%103 = OpBitcast %4 %102
%104 = OpImageFetch %23 %87 %98 Sample %103
%105 = OpImageRead %23 %88 %98
%106 = OpCompositeExtract %13 %83 2
%107 = OpBitcast %4 %106
%108 = OpCompositeExtract %13 %83 2
%109 = OpBitcast %4 %108
%110 = OpIAdd %4 %109 %6
%112 = OpCompositeConstruct %111 %98 %107
%113 = OpImageFetch %23 %89 %112 Lod %110
%114 = OpCompositeExtract %13 %83 0
%115 = OpBitcast %4 %114
%116 = OpCompositeExtract %13 %83 2
%117 = OpBitcast %4 %116
%118 = OpImageFetch %23 %90 %115 Lod %117
%119 = OpBitcast %22 %98
%120 = OpCompositeExtract %13 %83 2
%121 = OpBitcast %4 %120
%122 = OpImageFetch %23 %86 %119 Lod %121
%123 = OpBitcast %22 %98
%124 = OpCompositeExtract %13 %83 2
%125 = OpBitcast %4 %124
%126 = OpImageFetch %23 %87 %123 Sample %125
%127 = OpBitcast %22 %98
%128 = OpImageRead %23 %88 %127
%129 = OpBitcast %22 %98
%130 = OpCompositeExtract %13 %83 2
%131 = OpBitcast %4 %130
%132 = OpCompositeExtract %13 %83 2
%133 = OpBitcast %4 %132
%134 = OpIAdd %4 %133 %6
%135 = OpBitcast %13 %131
%136 = OpCompositeConstruct %20 %129 %135
%137 = OpImageFetch %23 %89 %136 Lod %134
%138 = OpCompositeExtract %13 %83 0
%140 = OpCompositeExtract %13 %83 2
%141 = OpBitcast %4 %140
%142 = OpImageFetch %23 %90 %138 Lod %141
%143 = OpCompositeExtract %4 %98 0
%144 = OpIAdd %23 %101 %104
%145 = OpIAdd %23 %144 %105
%146 = OpIAdd %23 %145 %113
%147 = OpIAdd %23 %146 %118
OpImageWrite %91 %143 %147
%148 = OpCompositeExtract %4 %98 0
%149 = OpBitcast %13 %148
%150 = OpIAdd %23 %122 %126
%151 = OpIAdd %23 %150 %128
%152 = OpIAdd %23 %151 %137
%153 = OpIAdd %23 %152 %142
OpImageWrite %91 %149 %153
%94 = OpBitcast %22 %93
%95 = OpVectorShuffle %22 %83 %83 0 1
%96 = OpIMul %22 %94 %95
%97 = OpBitcast %21 %96
%98 = OpCompositeConstruct %21 %3 %5
%99 = OpSRem %21 %97 %98
%100 = OpCompositeExtract %13 %83 2
%101 = OpBitcast %4 %100
%102 = OpImageFetch %23 %86 %99 Lod %101
%103 = OpCompositeExtract %13 %83 2
%104 = OpBitcast %4 %103
%105 = OpImageFetch %23 %87 %99 Sample %104
%106 = OpImageRead %23 %88 %99
%107 = OpCompositeExtract %13 %83 2
%108 = OpBitcast %4 %107
%109 = OpCompositeExtract %13 %83 2
%110 = OpBitcast %4 %109
%111 = OpIAdd %4 %110 %6
%113 = OpCompositeConstruct %112 %99 %108
%114 = OpImageFetch %23 %89 %113 Lod %111
%115 = OpCompositeExtract %13 %83 0
%116 = OpBitcast %4 %115
%117 = OpCompositeExtract %13 %83 2
%118 = OpBitcast %4 %117
%119 = OpImageFetch %23 %90 %116 Lod %118
%120 = OpBitcast %22 %99
%121 = OpCompositeExtract %13 %83 2
%122 = OpBitcast %4 %121
%123 = OpImageFetch %23 %86 %120 Lod %122
%124 = OpBitcast %22 %99
%125 = OpCompositeExtract %13 %83 2
%126 = OpBitcast %4 %125
%127 = OpImageFetch %23 %87 %124 Sample %126
%128 = OpBitcast %22 %99
%129 = OpImageRead %23 %88 %128
%130 = OpBitcast %22 %99
%131 = OpCompositeExtract %13 %83 2
%132 = OpBitcast %4 %131
%133 = OpCompositeExtract %13 %83 2
%134 = OpBitcast %4 %133
%135 = OpIAdd %4 %134 %6
%136 = OpBitcast %13 %132
%137 = OpCompositeConstruct %20 %130 %136
%138 = OpImageFetch %23 %89 %137 Lod %135
%139 = OpCompositeExtract %13 %83 0
%141 = OpCompositeExtract %13 %83 2
%142 = OpBitcast %4 %141
%143 = OpImageFetch %23 %90 %139 Lod %142
%144 = OpCompositeExtract %4 %99 0
%145 = OpIAdd %23 %102 %105
%146 = OpIAdd %23 %145 %106
%147 = OpIAdd %23 %146 %114
%148 = OpIAdd %23 %147 %119
OpImageWrite %91 %144 %148
%149 = OpCompositeExtract %4 %99 0
%150 = OpBitcast %13 %149
%151 = OpIAdd %23 %123 %127
%152 = OpIAdd %23 %151 %129
%153 = OpIAdd %23 %152 %138
%154 = OpIAdd %23 %153 %143
OpImageWrite %91 %150 %154
OpReturn
OpFunctionEnd
%157 = OpFunction %2 None %85
%154 = OpLabel
%156 = OpLoad %20 %155
%158 = OpLoad %15 %43
%159 = OpLoad %16 %45
%160 = OpLoad %18 %53
OpBranch %161
%161 = OpLabel
%162 = OpImageQuerySize %21 %159
%163 = OpVectorShuffle %22 %156 %156 0 1
%164 = OpBitcast %21 %163
%165 = OpIMul %21 %162 %164
%166 = OpCompositeConstruct %21 %3 %5
%167 = OpSRem %21 %165 %166
%168 = OpCompositeExtract %13 %156 2
%169 = OpBitcast %4 %168
%170 = OpImageFetch %32 %158 %167 Sample %169
%171 = OpCompositeExtract %8 %170 0
%172 = OpCompositeExtract %4 %167 0
%173 = OpConvertFToU %13 %171
%174 = OpCompositeConstruct %23 %173 %173 %173 %173
OpImageWrite %160 %172 %174
%158 = OpFunction %2 None %85
%155 = OpLabel
%157 = OpLoad %20 %156
%159 = OpLoad %15 %43
%160 = OpLoad %16 %45
%161 = OpLoad %18 %53
OpBranch %162
%162 = OpLabel
%163 = OpImageQuerySize %21 %160
%164 = OpBitcast %22 %163
%165 = OpVectorShuffle %22 %157 %157 0 1
%166 = OpIMul %22 %164 %165
%167 = OpBitcast %21 %166
%168 = OpCompositeConstruct %21 %3 %5
%169 = OpSRem %21 %167 %168
%170 = OpCompositeExtract %13 %157 2
%171 = OpBitcast %4 %170
%172 = OpImageFetch %32 %159 %169 Sample %171
%173 = OpCompositeExtract %8 %172 0
%174 = OpCompositeExtract %4 %169 0
%175 = OpConvertFToU %13 %173
%176 = OpCompositeConstruct %23 %175 %175 %175 %175
OpImageWrite %161 %174 %176
OpReturn
OpFunctionEnd
%178 = OpFunction %2 None %85
%175 = OpLabel
%179 = OpLoad %24 %55
%180 = OpLoad %25 %57
%181 = OpLoad %27 %62
%182 = OpLoad %28 %64
%183 = OpLoad %29 %66
%184 = OpLoad %30 %68
%185 = OpLoad %31 %70
OpBranch %186
%186 = OpLabel
%188 = OpImageQuerySizeLod %4 %179 %187
%190 = OpImageQuerySizeLod %4 %179 %188
%191 = OpImageQuerySizeLod %21 %180 %187
%192 = OpImageQuerySizeLod %21 %180 %6
%193 = OpImageQuerySizeLod %111 %181 %187
%194 = OpVectorShuffle %21 %193 %193 0 1
%195 = OpImageQuerySizeLod %111 %181 %6
%196 = OpVectorShuffle %21 %195 %195 0 1
%197 = OpImageQuerySizeLod %21 %182 %187
%198 = OpImageQuerySizeLod %21 %182 %6
%199 = OpImageQuerySizeLod %111 %183 %187
%200 = OpVectorShuffle %21 %199 %199 0 0
%201 = OpImageQuerySizeLod %111 %183 %6
%202 = OpVectorShuffle %21 %201 %201 0 0
%203 = OpImageQuerySizeLod %111 %184 %187
%204 = OpImageQuerySizeLod %111 %184 %6
%205 = OpImageQuerySize %21 %185
%206 = OpCompositeExtract %4 %191 1
%207 = OpIAdd %4 %188 %206
%208 = OpCompositeExtract %4 %192 1
%209 = OpIAdd %4 %207 %208
%210 = OpCompositeExtract %4 %194 1
%211 = OpIAdd %4 %209 %210
%212 = OpCompositeExtract %4 %196 1
%213 = OpIAdd %4 %211 %212
%214 = OpCompositeExtract %4 %197 1
%215 = OpIAdd %4 %213 %214
%216 = OpCompositeExtract %4 %198 1
%217 = OpIAdd %4 %215 %216
%218 = OpCompositeExtract %4 %200 1
%219 = OpIAdd %4 %217 %218
%220 = OpCompositeExtract %4 %202 1
%221 = OpIAdd %4 %219 %220
%222 = OpCompositeExtract %4 %203 2
%223 = OpIAdd %4 %221 %222
%224 = OpCompositeExtract %4 %204 2
%225 = OpIAdd %4 %223 %224
%226 = OpConvertSToF %8 %225
%227 = OpCompositeConstruct %32 %226 %226 %226 %226
OpStore %176 %227
%180 = OpFunction %2 None %85
%177 = OpLabel
%181 = OpLoad %24 %55
%182 = OpLoad %25 %57
%183 = OpLoad %27 %62
%184 = OpLoad %28 %64
%185 = OpLoad %29 %66
%186 = OpLoad %30 %68
%187 = OpLoad %31 %70
OpBranch %188
%188 = OpLabel
%190 = OpImageQuerySizeLod %4 %181 %189
%191 = OpBitcast %13 %190
%192 = OpBitcast %4 %191
%193 = OpImageQuerySizeLod %4 %181 %192
%194 = OpBitcast %13 %193
%195 = OpImageQuerySizeLod %21 %182 %189
%196 = OpBitcast %22 %195
%197 = OpImageQuerySizeLod %21 %182 %6
%198 = OpBitcast %22 %197
%199 = OpImageQuerySizeLod %112 %183 %189
%200 = OpBitcast %20 %199
%201 = OpVectorShuffle %22 %200 %200 0 1
%202 = OpImageQuerySizeLod %112 %183 %6
%203 = OpBitcast %20 %202
%204 = OpVectorShuffle %22 %203 %203 0 1
%205 = OpImageQuerySizeLod %21 %184 %189
%206 = OpBitcast %22 %205
%207 = OpImageQuerySizeLod %21 %184 %6
%208 = OpBitcast %22 %207
%209 = OpImageQuerySizeLod %112 %185 %189
%210 = OpBitcast %20 %209
%211 = OpVectorShuffle %22 %210 %210 0 0
%212 = OpImageQuerySizeLod %112 %185 %6
%213 = OpBitcast %20 %212
%214 = OpVectorShuffle %22 %213 %213 0 0
%215 = OpImageQuerySizeLod %112 %186 %189
%216 = OpBitcast %20 %215
%217 = OpImageQuerySizeLod %112 %186 %6
%218 = OpBitcast %20 %217
%219 = OpImageQuerySize %21 %187
%220 = OpBitcast %22 %219
%221 = OpCompositeExtract %13 %196 1
%222 = OpIAdd %13 %191 %221
%223 = OpCompositeExtract %13 %198 1
%224 = OpIAdd %13 %222 %223
%225 = OpCompositeExtract %13 %201 1
%226 = OpIAdd %13 %224 %225
%227 = OpCompositeExtract %13 %204 1
%228 = OpIAdd %13 %226 %227
%229 = OpCompositeExtract %13 %206 1
%230 = OpIAdd %13 %228 %229
%231 = OpCompositeExtract %13 %208 1
%232 = OpIAdd %13 %230 %231
%233 = OpCompositeExtract %13 %211 1
%234 = OpIAdd %13 %232 %233
%235 = OpCompositeExtract %13 %214 1
%236 = OpIAdd %13 %234 %235
%237 = OpCompositeExtract %13 %216 2
%238 = OpIAdd %13 %236 %237
%239 = OpCompositeExtract %13 %218 2
%240 = OpIAdd %13 %238 %239
%241 = OpConvertUToF %8 %240
%242 = OpCompositeConstruct %32 %241 %241 %241 %241
OpStore %178 %242
OpReturn
OpFunctionEnd
%230 = OpFunction %2 None %85
%228 = OpLabel
%231 = OpLoad %25 %57
%232 = OpLoad %27 %62
%233 = OpLoad %28 %64
%234 = OpLoad %29 %66
%235 = OpLoad %30 %68
%236 = OpLoad %31 %70
OpBranch %237
%237 = OpLabel
%238 = OpImageQueryLevels %4 %231
%239 = OpImageQueryLevels %4 %232
%240 = OpImageQuerySizeLod %111 %232 %187
%241 = OpCompositeExtract %4 %240 2
%242 = OpImageQueryLevels %4 %233
%243 = OpImageQueryLevels %4 %234
%244 = OpImageQuerySizeLod %111 %234 %187
%245 = OpCompositeExtract %4 %244 2
%246 = OpImageQueryLevels %4 %235
%247 = OpImageQuerySamples %4 %236
%248 = OpIAdd %4 %241 %245
%249 = OpIAdd %4 %248 %247
%250 = OpIAdd %4 %249 %238
%251 = OpIAdd %4 %250 %239
%252 = OpIAdd %4 %251 %246
%253 = OpIAdd %4 %252 %242
%254 = OpIAdd %4 %253 %243
%255 = OpConvertSToF %8 %254
%256 = OpCompositeConstruct %32 %255 %255 %255 %255
OpStore %229 %256
%245 = OpFunction %2 None %85
%243 = OpLabel
%246 = OpLoad %25 %57
%247 = OpLoad %27 %62
%248 = OpLoad %28 %64
%249 = OpLoad %29 %66
%250 = OpLoad %30 %68
%251 = OpLoad %31 %70
OpBranch %252
%252 = OpLabel
%253 = OpImageQueryLevels %4 %246
%254 = OpBitcast %13 %253
%255 = OpImageQueryLevels %4 %247
%256 = OpBitcast %13 %255
%257 = OpImageQuerySizeLod %112 %247 %189
%258 = OpCompositeExtract %4 %257 2
%259 = OpBitcast %13 %258
%260 = OpImageQueryLevels %4 %248
%261 = OpBitcast %13 %260
%262 = OpImageQueryLevels %4 %249
%263 = OpBitcast %13 %262
%264 = OpImageQuerySizeLod %112 %249 %189
%265 = OpCompositeExtract %4 %264 2
%266 = OpBitcast %13 %265
%267 = OpImageQueryLevels %4 %250
%268 = OpBitcast %13 %267
%269 = OpImageQuerySamples %4 %251
%270 = OpBitcast %13 %269
%271 = OpIAdd %13 %259 %266
%272 = OpIAdd %13 %271 %270
%273 = OpIAdd %13 %272 %254
%274 = OpIAdd %13 %273 %256
%275 = OpIAdd %13 %274 %268
%276 = OpIAdd %13 %275 %261
%277 = OpIAdd %13 %276 %263
%278 = OpConvertUToF %8 %277
%279 = OpCompositeConstruct %32 %278 %278 %278 %278
OpStore %244 %279
OpReturn
OpFunctionEnd
%259 = OpFunction %2 None %85
%257 = OpLabel
%260 = OpLoad %24 %55
%261 = OpLoad %25 %57
%262 = OpLoad %33 %72
OpBranch %263
%263 = OpLabel
%264 = OpCompositeConstruct %34 %7 %7
%265 = OpCompositeExtract %8 %264 0
%267 = OpSampledImage %266 %260 %262
%268 = OpImageSampleImplicitLod %32 %267 %265
%270 = OpSampledImage %269 %261 %262
%271 = OpImageSampleImplicitLod %32 %270 %264
%272 = OpSampledImage %269 %261 %262
%273 = OpImageSampleImplicitLod %32 %272 %264 ConstOffset %38
%274 = OpSampledImage %269 %261 %262
%275 = OpImageSampleExplicitLod %32 %274 %264 Lod %9
%276 = OpSampledImage %269 %261 %262
%277 = OpImageSampleExplicitLod %32 %276 %264 Lod|ConstOffset %9 %38
%278 = OpSampledImage %269 %261 %262
%279 = OpImageSampleImplicitLod %32 %278 %264 Bias|ConstOffset %11 %38
%280 = OpFAdd %32 %268 %271
%281 = OpFAdd %32 %280 %273
%282 = OpFAdd %32 %281 %275
%283 = OpFAdd %32 %282 %277
OpStore %258 %283
%282 = OpFunction %2 None %85
%280 = OpLabel
%283 = OpLoad %24 %55
%284 = OpLoad %25 %57
%285 = OpLoad %33 %72
OpBranch %286
%286 = OpLabel
%287 = OpCompositeConstruct %34 %7 %7
%288 = OpCompositeExtract %8 %287 0
%290 = OpSampledImage %289 %283 %285
%291 = OpImageSampleImplicitLod %32 %290 %288
%293 = OpSampledImage %292 %284 %285
%294 = OpImageSampleImplicitLod %32 %293 %287
%295 = OpSampledImage %292 %284 %285
%296 = OpImageSampleImplicitLod %32 %295 %287 ConstOffset %38
%297 = OpSampledImage %292 %284 %285
%298 = OpImageSampleExplicitLod %32 %297 %287 Lod %9
%299 = OpSampledImage %292 %284 %285
%300 = OpImageSampleExplicitLod %32 %299 %287 Lod|ConstOffset %9 %38
%301 = OpSampledImage %292 %284 %285
%302 = OpImageSampleImplicitLod %32 %301 %287 Bias|ConstOffset %11 %38
%303 = OpFAdd %32 %291 %294
%304 = OpFAdd %32 %303 %296
%305 = OpFAdd %32 %304 %298
%306 = OpFAdd %32 %305 %300
OpStore %281 %306
OpReturn
OpFunctionEnd
%287 = OpFunction %2 None %85
%284 = OpLabel
%288 = OpLoad %33 %74
%289 = OpLoad %35 %76
%290 = OpLoad %36 %78
OpBranch %291
%291 = OpLabel
%292 = OpCompositeConstruct %34 %7 %7
%294 = OpSampledImage %293 %289 %288
%295 = OpImageSampleDrefImplicitLod %8 %294 %292 %7
%296 = OpSampledImage %293 %289 %288
%297 = OpImageSampleDrefExplicitLod %8 %296 %292 %7 Lod %298
%299 = OpCompositeConstruct %37 %7 %7 %7
%301 = OpSampledImage %300 %290 %288
%302 = OpImageSampleDrefExplicitLod %8 %301 %299 %7 Lod %298
%303 = OpFAdd %8 %295 %297
OpStore %285 %303
OpReturn
OpFunctionEnd
%306 = OpFunction %2 None %85
%304 = OpLabel
%307 = OpLoad %25 %57
%308 = OpLoad %12 %59
%309 = OpLoad %26 %60
%310 = OpLoad %33 %72
%310 = OpFunction %2 None %85
%307 = OpLabel
%311 = OpLoad %33 %74
%312 = OpLoad %35 %76
OpBranch %313
%313 = OpLabel
%314 = OpCompositeConstruct %34 %7 %7
%315 = OpSampledImage %269 %307 %310
%316 = OpImageGather %32 %315 %314 %317
%318 = OpSampledImage %269 %307 %310
%319 = OpImageGather %32 %318 %314 %320 ConstOffset %38
%321 = OpSampledImage %293 %312 %311
%322 = OpImageDrefGather %32 %321 %314 %7
%323 = OpSampledImage %293 %312 %311
%324 = OpImageDrefGather %32 %323 %314 %7 ConstOffset %38
%326 = OpSampledImage %325 %308 %310
%327 = OpImageGather %23 %326 %314 %187
%330 = OpSampledImage %329 %309 %310
%331 = OpImageGather %328 %330 %314 %187
%332 = OpConvertUToF %32 %327
%333 = OpConvertSToF %32 %331
%334 = OpFAdd %32 %332 %333
%335 = OpFAdd %32 %316 %319
%336 = OpFAdd %32 %335 %322
%337 = OpFAdd %32 %336 %324
%338 = OpFAdd %32 %337 %334
OpStore %305 %338
%313 = OpLoad %36 %78
OpBranch %314
%314 = OpLabel
%315 = OpCompositeConstruct %34 %7 %7
%317 = OpSampledImage %316 %312 %311
%318 = OpImageSampleDrefImplicitLod %8 %317 %315 %7
%319 = OpSampledImage %316 %312 %311
%320 = OpImageSampleDrefExplicitLod %8 %319 %315 %7 Lod %321
%322 = OpCompositeConstruct %37 %7 %7 %7
%324 = OpSampledImage %323 %313 %311
%325 = OpImageSampleDrefExplicitLod %8 %324 %322 %7 Lod %321
%326 = OpFAdd %8 %318 %320
OpStore %308 %326
OpReturn
OpFunctionEnd
%341 = OpFunction %2 None %85
%339 = OpLabel
%342 = OpLoad %33 %72
%343 = OpLoad %35 %76
OpBranch %344
%344 = OpLabel
%345 = OpCompositeConstruct %34 %7 %7
%346 = OpSampledImage %293 %343 %342
%347 = OpImageSampleImplicitLod %32 %346 %345
%348 = OpCompositeExtract %8 %347 0
%349 = OpSampledImage %293 %343 %342
%350 = OpImageGather %32 %349 %345 %187
%351 = OpCompositeConstruct %32 %348 %348 %348 %348
%352 = OpFAdd %32 %351 %350
OpStore %340 %352
%329 = OpFunction %2 None %85
%327 = OpLabel
%330 = OpLoad %25 %57
%331 = OpLoad %12 %59
%332 = OpLoad %26 %60
%333 = OpLoad %33 %72
%334 = OpLoad %33 %74
%335 = OpLoad %35 %76
OpBranch %336
%336 = OpLabel
%337 = OpCompositeConstruct %34 %7 %7
%338 = OpSampledImage %292 %330 %333
%339 = OpImageGather %32 %338 %337 %340
%341 = OpSampledImage %292 %330 %333
%342 = OpImageGather %32 %341 %337 %343 ConstOffset %38
%344 = OpSampledImage %316 %335 %334
%345 = OpImageDrefGather %32 %344 %337 %7
%346 = OpSampledImage %316 %335 %334
%347 = OpImageDrefGather %32 %346 %337 %7 ConstOffset %38
%349 = OpSampledImage %348 %331 %333
%350 = OpImageGather %23 %349 %337 %189
%353 = OpSampledImage %352 %332 %333
%354 = OpImageGather %351 %353 %337 %189
%355 = OpConvertUToF %32 %350
%356 = OpConvertSToF %32 %354
%357 = OpFAdd %32 %355 %356
%358 = OpFAdd %32 %339 %342
%359 = OpFAdd %32 %358 %345
%360 = OpFAdd %32 %359 %347
%361 = OpFAdd %32 %360 %357
OpStore %328 %361
OpReturn
OpFunctionEnd
%364 = OpFunction %2 None %85
%362 = OpLabel
%365 = OpLoad %33 %72
%366 = OpLoad %35 %76
OpBranch %367
%367 = OpLabel
%368 = OpCompositeConstruct %34 %7 %7
%369 = OpSampledImage %316 %366 %365
%370 = OpImageSampleImplicitLod %32 %369 %368
%371 = OpCompositeExtract %8 %370 0
%372 = OpSampledImage %316 %366 %365
%373 = OpImageGather %32 %372 %368 %189
%374 = OpCompositeConstruct %32 %371 %371 %371 %371
%375 = OpFAdd %32 %374 %373
OpStore %363 %375
OpReturn
OpFunctionEnd

62
tests/out/wgsl/binding-arrays.wgsl
View File

@ -27,28 +27,28 @@ var<uniform> uni: UniformIndex;
@fragment
fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
var i1_: i32;
var i2_: vec2<i32>;
var u1_: u32;
var u2_: vec2<u32>;
var v1_: f32;
var v4_: vec4<f32>;
let uniform_index = uni.index;
let non_uniform_index = fragment_in.index;
i1_ = 0;
i2_ = vec2<i32>(0);
u1_ = 0u;
u2_ = vec2<u32>(0u);
v1_ = 0.0;
v4_ = vec4<f32>(0.0);
let uv = vec2<f32>(0.0);
let pix = vec2<i32>(0);
let _e22 = textureDimensions(texture_array_unbounded[0]);
let _e23 = i2_;
i2_ = (_e23 + _e22);
let _e23 = u2_;
u2_ = (_e23 + _e22);
let _e27 = textureDimensions(texture_array_unbounded[uniform_index]);
let _e28 = i2_;
i2_ = (_e28 + _e27);
let _e28 = u2_;
u2_ = (_e28 + _e27);
let _e32 = textureDimensions(texture_array_unbounded[non_uniform_index]);
let _e33 = i2_;
i2_ = (_e33 + _e32);
let _e33 = u2_;
u2_ = (_e33 + _e32);
let _e41 = textureGather(0, texture_array_bounded[0], samp[0], uv);
let _e42 = v4_;
v4_ = (_e42 + _e41);
@ -77,32 +77,32 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
let _e101 = v4_;
v4_ = (_e101 + _e100);
let _e106 = textureNumLayers(texture_array_2darray[0]);
let _e107 = i1_;
i1_ = (_e107 + _e106);
let _e107 = u1_;
u1_ = (_e107 + _e106);
let _e111 = textureNumLayers(texture_array_2darray[uniform_index]);
let _e112 = i1_;
i1_ = (_e112 + _e111);
let _e112 = u1_;
u1_ = (_e112 + _e111);
let _e116 = textureNumLayers(texture_array_2darray[non_uniform_index]);
let _e117 = i1_;
i1_ = (_e117 + _e116);
let _e117 = u1_;
u1_ = (_e117 + _e116);
let _e122 = textureNumLevels(texture_array_bounded[0]);
let _e123 = i1_;
i1_ = (_e123 + _e122);
let _e123 = u1_;
u1_ = (_e123 + _e122);
let _e127 = textureNumLevels(texture_array_bounded[uniform_index]);
let _e128 = i1_;
i1_ = (_e128 + _e127);
let _e128 = u1_;
u1_ = (_e128 + _e127);
let _e132 = textureNumLevels(texture_array_bounded[non_uniform_index]);
let _e133 = i1_;
i1_ = (_e133 + _e132);
let _e133 = u1_;
u1_ = (_e133 + _e132);
let _e138 = textureNumSamples(texture_array_multisampled[0]);
let _e139 = i1_;
i1_ = (_e139 + _e138);
let _e139 = u1_;
u1_ = (_e139 + _e138);
let _e143 = textureNumSamples(texture_array_multisampled[uniform_index]);
let _e144 = i1_;
i1_ = (_e144 + _e143);
let _e144 = u1_;
u1_ = (_e144 + _e143);
let _e148 = textureNumSamples(texture_array_multisampled[non_uniform_index]);
let _e149 = i1_;
i1_ = (_e149 + _e148);
let _e149 = u1_;
u1_ = (_e149 + _e148);
let _e157 = textureSample(texture_array_bounded[0], samp[0], uv);
let _e158 = v4_;
v4_ = (_e158 + _e157);
@ -163,9 +163,9 @@ fn main(fragment_in: FragmentIn) -> @location(0) vec4<f32> {
textureStore(texture_array_storage[uniform_index], pix, _e307);
let _e310 = v4_;
textureStore(texture_array_storage[non_uniform_index], pix, _e310);
let _e311 = i2_;
let _e312 = i1_;
let v2_ = vec2<f32>((_e311 + vec2<i32>(_e312)));
let _e311 = u2_;
let _e312 = u1_;
let v2_ = vec2<f32>((_e311 + vec2<u32>(_e312)));
let _e316 = v4_;
let _e323 = v1_;
return ((_e316 + vec4<f32>(v2_.x, v2_.y, v2_.x, v2_.y)) + vec4<f32>(_e323));

4
tests/out/wgsl/image.wgsl
View File

@ -44,7 +44,7 @@ var image_cube_depth: texture_depth_cube;
@compute @workgroup_size(16, 1, 1)
fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
let dim = textureDimensions(image_storage_src);
let itc = ((dim * vec2<i32>(local_id.xy)) % vec2<i32>(10, 20));
let itc = (vec2<i32>((dim * local_id.xy)) % vec2<i32>(10, 20));
let value1_ = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
let value2_ = textureLoad(image_multisampled_src, itc, i32(local_id.z));
let value4_ = textureLoad(image_storage_src, itc);
@ -63,7 +63,7 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
@compute @workgroup_size(16, 1, 1)
fn depth_load(@builtin(local_invocation_id) local_id_1: vec3<u32>) {
let dim_1 = textureDimensions(image_storage_src);
let itc_1 = ((dim_1 * vec2<i32>(local_id_1.xy)) % vec2<i32>(10, 20));
let itc_1 = (vec2<i32>((dim_1 * local_id_1.xy)) % vec2<i32>(10, 20));
let val = textureLoad(image_depth_multisampled_src, itc_1, i32(local_id_1.z));
textureStore(image_dst, itc_1.x, vec4<u32>(u32(val)));
return;

76
tests/out/wgsl/images-frag.wgsl
View File

@ -22,15 +22,15 @@ fn testImg1D(coord: i32) {
coord_1 = coord;
let _e10 = textureDimensions(img1D);
size = _e10;
size = i32(_e10);
_ = coord_1;
_ = vec4<f32>(f32(2));
let _e16 = coord_1;
textureStore(img1D, _e16, vec4<f32>(f32(2)));
let _e17 = coord_1;
textureStore(img1D, _e17, vec4<f32>(f32(2)));
_ = coord_1;
let _e21 = coord_1;
let _e22 = textureLoad(img1D, _e21);
c = _e22;
let _e22 = coord_1;
let _e23 = textureLoad(img1D, _e22);
c = _e23;
return;
}
@ -42,15 +42,15 @@ fn testImg1DArray(coord_2: vec2<i32>) {
coord_3 = coord_2;
let _e10 = textureDimensions(img1DArray);
let _e11 = textureNumLayers(img1DArray);
size_1 = vec2<f32>(vec2<i32>(_e10, _e11));
size_1 = vec2<f32>(vec2<i32>(vec2<u32>(_e10, _e11)));
_ = coord_3;
let _e16 = coord_3;
let _e19 = textureLoad(img1DArray, _e16.x, _e16.y);
c_1 = _e19;
let _e17 = coord_3;
let _e20 = textureLoad(img1DArray, _e17.x, _e17.y);
c_1 = _e20;
_ = coord_3;
_ = vec4<f32>(f32(2));
let _e25 = coord_3;
textureStore(img1DArray, _e25.x, _e25.y, vec4<f32>(f32(2)));
let _e26 = coord_3;
textureStore(img1DArray, _e26.x, _e26.y, vec4<f32>(f32(2)));
return;
}
@ -61,15 +61,15 @@ fn testImg2D(coord_4: vec2<i32>) {
coord_5 = coord_4;
let _e10 = textureDimensions(img2D);
size_2 = vec2<f32>(_e10);
size_2 = vec2<f32>(vec2<i32>(_e10));
_ = coord_5;
let _e14 = coord_5;
let _e15 = textureLoad(img2D, _e14);
c_2 = _e15;
let _e15 = coord_5;
let _e16 = textureLoad(img2D, _e15);
c_2 = _e16;
_ = coord_5;
_ = vec4<f32>(f32(2));
let _e21 = coord_5;
textureStore(img2D, _e21, vec4<f32>(f32(2)));
let _e22 = coord_5;
textureStore(img2D, _e22, vec4<f32>(f32(2)));
return;
}
@ -81,15 +81,15 @@ fn testImg2DArray(coord_6: vec3<i32>) {
coord_7 = coord_6;
let _e10 = textureDimensions(img2DArray);
let _e13 = textureNumLayers(img2DArray);
size_3 = vec3<f32>(vec3<i32>(_e10.x, _e10.y, _e13));
size_3 = vec3<f32>(vec3<i32>(vec3<u32>(_e10.x, _e10.y, _e13)));
_ = coord_7;
let _e18 = coord_7;
let _e21 = textureLoad(img2DArray, _e18.xy, _e18.z);
c_3 = _e21;
let _e19 = coord_7;
let _e22 = textureLoad(img2DArray, _e19.xy, _e19.z);
c_3 = _e22;
_ = coord_7;
_ = vec4<f32>(f32(2));
let _e27 = coord_7;
textureStore(img2DArray, _e27.xy, _e27.z, vec4<f32>(f32(2)));
let _e28 = coord_7;
textureStore(img2DArray, _e28.xy, _e28.z, vec4<f32>(f32(2)));
return;
}
@ -100,15 +100,15 @@ fn testImg3D(coord_8: vec3<i32>) {
coord_9 = coord_8;
let _e10 = textureDimensions(img3D);
size_4 = vec3<f32>(_e10);
size_4 = vec3<f32>(vec3<i32>(_e10));
_ = coord_9;
let _e14 = coord_9;
let _e15 = textureLoad(img3D, _e14);
c_4 = _e15;
let _e15 = coord_9;
let _e16 = textureLoad(img3D, _e15);
c_4 = _e16;
_ = coord_9;
_ = vec4<f32>(f32(2));
let _e21 = coord_9;
textureStore(img3D, _e21, vec4<f32>(f32(2)));
let _e22 = coord_9;
textureStore(img3D, _e22, vec4<f32>(f32(2)));
return;
}
@ -119,11 +119,11 @@ fn testImgReadOnly(coord_10: vec2<i32>) {
coord_11 = coord_10;
let _e10 = textureDimensions(img2D);
size_5 = vec2<f32>(_e10);
size_5 = vec2<f32>(vec2<i32>(_e10));
_ = coord_11;
let _e14 = coord_11;
let _e15 = textureLoad(imgReadOnly, _e14);
c_5 = _e15;
let _e15 = coord_11;
let _e16 = textureLoad(imgReadOnly, _e15);
c_5 = _e16;
return;
}
@ -133,11 +133,11 @@ fn testImgWriteOnly(coord_12: vec2<i32>) {
coord_13 = coord_12;
let _e10 = textureDimensions(img2D);
size_6 = vec2<f32>(_e10);
size_6 = vec2<f32>(vec2<i32>(_e10));
_ = coord_13;
_ = vec4<f32>(f32(2));
let _e17 = coord_13;
textureStore(imgWriteOnly, _e17, vec4<f32>(f32(2)));
let _e18 = coord_13;
textureStore(imgWriteOnly, _e18, vec4<f32>(f32(2)));
return;
}
@ -147,7 +147,7 @@ fn testImgWriteReadOnly(coord_14: vec2<i32>) {
coord_15 = coord_14;
let _e10 = textureDimensions(imgWriteReadOnly);
size_7 = vec2<f32>(_e10);
size_7 = vec2<f32>(vec2<i32>(_e10));
return;
}

1158
tests/out/wgsl/samplers-frag.wgsl
View File

File diff suppressed because it is too large Load Diff

8
tests/spirv-capabilities.rs
View File

@ -94,7 +94,7 @@ fn image_queries() {
require(
&[Ca::ImageQuery],
r#"
fn f(i: texture_2d<f32>) -> vec2<i32> {
fn f(i: texture_2d<f32>) -> vec2<u32> {
return textureDimensions(i);
}
"#,
@ -102,7 +102,7 @@ fn image_queries() {
require(
&[Ca::ImageQuery],
r#"
fn f(i: texture_2d_array<f32>) -> i32 {
fn f(i: texture_2d_array<f32>) -> u32 {
return textureNumLayers(i);
}
"#,
@ -110,7 +110,7 @@ fn image_queries() {
require(
&[Ca::ImageQuery],
r#"
fn f(i: texture_2d<f32>) -> i32 {
fn f(i: texture_2d<f32>) -> u32 {
return textureNumLevels(i);
}
"#,
@ -118,7 +118,7 @@ fn image_queries() {
require(
&[Ca::ImageQuery],
r#"
fn f(i: texture_multisampled_2d<f32>) -> i32 {
fn f(i: texture_multisampled_2d<f32>) -> u32 {
return textureNumSamples(i);
}
"#,