diff --git a/naga/src/front/spv/image.rs b/naga/src/front/spv/image.rs index 71ba79e1e..787adf5c3 100644 --- a/naga/src/front/spv/image.rs +++ b/naga/src/front/spv/image.rs @@ -30,6 +30,16 @@ impl<'function> super::BlockContext<'function> { match self.expressions[handle] { crate::Expression::GlobalVariable(handle) => Ok(self.global_arena[handle].ty), crate::Expression::FunctionArgument(i) => Ok(self.arguments[i as usize].ty), + crate::Expression::Access { base, .. } => match self.expressions[base] { + crate::Expression::GlobalVariable(handle) => { + let array_ty = self.global_arena[handle].ty; + match self.type_arena[array_ty].inner { + crate::TypeInner::BindingArray { base, .. } => Ok(base), + _ => Err(Error::InvalidImageBaseType(array_ty)), + } + } + ref other => Err(Error::InvalidImageExpression(other.clone())), + }, ref other => Err(Error::InvalidImageExpression(other.clone())), } } @@ -210,6 +220,22 @@ pub(super) fn patch_comparison_type( arrayed, }, crate::TypeInner::Sampler { .. } => crate::TypeInner::Sampler { comparison: true }, + crate::TypeInner::BindingArray { base, .. } => { + let base_ty = &arena[base]; + match base_ty.inner { + crate::TypeInner::Image { + class: crate::ImageClass::Sampled { multi, .. }, + dim, + arrayed, + } => crate::TypeInner::Image { + class: crate::ImageClass::Depth { multi }, + dim, + arrayed, + }, + crate::TypeInner::Sampler { .. } => crate::TypeInner::Sampler { comparison: true }, + ref other => unreachable!("Unexpected type for comparison mutation: {:?}", other), + } + } ref other => unreachable!("Unexpected type for comparison mutation: {:?}", other), }; diff --git a/naga/tests/in/spv/binding-arrays.stress.param.ron b/naga/tests/in/spv/binding-arrays.stress.param.ron new file mode 100644 index 000000000..c70ada993 --- /dev/null +++ b/naga/tests/in/spv/binding-arrays.stress.param.ron @@ -0,0 +1,3 @@ +( + god_mode: true, +) \ No newline at end of file diff --git a/naga/tests/in/spv/binding-arrays.stress.slang b/naga/tests/in/spv/binding-arrays.stress.slang new file mode 100644 index 000000000..2d16cc671 --- /dev/null +++ b/naga/tests/in/spv/binding-arrays.stress.slang @@ -0,0 +1,44 @@ +// slangc bin + +[[vk::binding(0, 0)]] +Texture2D textures[]; + +[[vk::binding(0, 1)]] +SamplerState samplers[]; + +[[vk::binding(0, 2)]] +Texture2D cmp_texture; + +[[vk::binding(0, 3)]] +SamplerComparisonState cmp_samplers[]; + +[[vk::binding(0, 4)]] +RWTexture2D rw_textures[]; + +[shader("fragment")] +float4 main(uint tex_id, uint samp_id, float2 uv, int3 location, float lod, float clamp, float cmp, float2 ddx, float2 ddy) + : SV_Target +{ + float4 out = float4(0.0); + out += textures[tex_id].Sample(samplers[samp_id], uv); + out += textures[tex_id].Sample(samplers[samp_id], uv, int2(1.0)); + out += textures[tex_id].Sample(samplers[samp_id], uv, int2(1.0), clamp); + out += textures[tex_id].SampleLevel(samplers[samp_id], uv, lod); + out += textures[tex_id].SampleLevel(samplers[samp_id], uv, lod, int2(1.0)); + out += cmp_texture.SampleCmp(cmp_samplers[samp_id], uv, cmp); + out += cmp_texture.SampleCmp(cmp_samplers[samp_id], uv, cmp, int2(1.0)); + out += cmp_texture.SampleCmpLevelZero(cmp_samplers[samp_id], uv, cmp); + out += cmp_texture.SampleCmpLevelZero(cmp_samplers[samp_id], uv, cmp, int2(1.0)); + out += textures[tex_id].SampleGrad(samplers[samp_id], uv, ddx, ddy); + out += textures[tex_id].SampleGrad(samplers[samp_id], uv, ddx, ddy, int2(1.0)); + out += textures[tex_id].SampleGrad(samplers[samp_id], uv, ddx, ddy, int2(1.0), clamp); + + out += textures[tex_id][location.xy]; + out += textures[tex_id].Load(location, int2(1.0)); + out += rw_textures[tex_id].Load(location.xy); + out += rw_textures[tex_id].Load(location.xy, int2(1.0)); + + rw_textures[tex_id][location.xy] = float4(1.0); + + return out; +}; diff --git a/naga/tests/in/spv/binding-arrays.stress.spv b/naga/tests/in/spv/binding-arrays.stress.spv new file mode 100644 index 000000000..33f0d15bd Binary files /dev/null and b/naga/tests/in/spv/binding-arrays.stress.spv differ diff --git a/naga/tests/in/spv/binding-arrays.stress.spvasm b/naga/tests/in/spv/binding-arrays.stress.spvasm new file mode 100644 index 000000000..7691de1cf --- /dev/null +++ b/naga/tests/in/spv/binding-arrays.stress.spvasm @@ -0,0 +1,269 @@ +; SPIR-V +; Version: 1.5 +; Generator: Khronos; 40 +; Bound: 184 +; Schema: 0 + OpCapability RuntimeDescriptorArray + OpCapability MinLod + OpCapability StorageImageReadWithoutFormat + OpCapability StorageImageWriteWithoutFormat + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %main "main" %textures %samplers %cmp_samplers %cmp_texture %rw_textures %entryPointParam_main %tex_id %samp_id %uv %location %lod %clamp %cmp %ddx %ddy + OpExecutionMode %main OriginUpperLeft + OpSource Slang 1 + OpName %tex_id "tex_id" + OpName %textures "textures" + OpName %samp_id "samp_id" + OpName %samplers "samplers" + OpName %uv "uv" + OpName %sampledImage "sampledImage" + OpName %sampled "sampled" + OpName %sampledImage_0 "sampledImage" + OpName %sampled_0 "sampled" + OpName %out "out" + OpName %clamp "clamp" + OpName %sampledImage_1 "sampledImage" + OpName %sampled_1 "sampled" + OpName %out_0 "out" + OpName %lod "lod" + OpName %sampledImage_2 "sampledImage" + OpName %sampled_2 "sampled" + OpName %out_1 "out" + OpName %sampledImage_3 "sampledImage" + OpName %sampled_3 "sampled" + OpName %out_2 "out" + OpName %cmp_samplers "cmp_samplers" + OpName %cmp_texture "cmp_texture" + OpName %cmp "cmp" + OpName %sampledImage_4 "sampledImage" + OpName %out_3 "out" + OpName %sampledImage_5 "sampledImage" + OpName %out_4 "out" + OpName %sampledImage_6 "sampledImage" + OpName %out_5 "out" + OpName %sampledImage_7 "sampledImage" + OpName %out_6 "out" + OpName %ddx "ddx" + OpName %ddy "ddy" + OpName %sampledImage_8 "sampledImage" + OpName %sampled_4 "sampled" + OpName %out_7 "out" + OpName %sampledImage_9 "sampledImage" + OpName %sampled_5 "sampled" + OpName %out_8 "out" + OpName %sampledImage_10 "sampledImage" + OpName %sampled_6 "sampled" + OpName %out_9 "out" + OpName %location "location" + OpName %sampled_7 "sampled" + OpName %out_10 "out" + OpName %sampled_8 "sampled" + OpName %out_11 "out" + OpName %rw_textures "rw_textures" + OpName %sampled_9 "sampled" + OpName %out_12 "out" + OpName %sampled_10 "sampled" + OpName %out_13 "out" + OpName %entryPointParam_main "entryPointParam_main" + OpName %main "main" + OpDecorate %tex_id Location 0 + OpDecorate %tex_id Flat + OpDecorate %_runtimearr_11 ArrayStride 8 + OpDecorate %textures Binding 0 + OpDecorate %textures DescriptorSet 0 + OpDecorate %samp_id Location 1 + OpDecorate %samp_id Flat + OpDecorate %_runtimearr_20 ArrayStride 8 + OpDecorate %samplers Binding 0 + OpDecorate %samplers DescriptorSet 1 + OpDecorate %uv Location 2 + OpDecorate %clamp Location 5 + OpDecorate %lod Location 4 + OpDecorate %_runtimearr_20_0 ArrayStride 8 + OpDecorate %cmp_samplers Binding 0 + OpDecorate %cmp_samplers DescriptorSet 3 + OpDecorate %cmp_texture Binding 0 + OpDecorate %cmp_texture DescriptorSet 2 + OpDecorate %cmp Location 6 + OpDecorate %ddx Location 7 + OpDecorate %ddy Location 8 + OpDecorate %location Location 3 + OpDecorate %location Flat + OpDecorate %_runtimearr_165 ArrayStride 8 + OpDecorate %rw_textures Binding 0 + OpDecorate %rw_textures DescriptorSet 4 + OpDecorate %entryPointParam_main Location 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint + %float = OpTypeFloat 32 + %11 = OpTypeImage %float 2D 2 0 0 1 Unknown +%_runtimearr_11 = OpTypeRuntimeArray %11 +%_ptr_UniformConstant__runtimearr_11 = OpTypePointer UniformConstant %_runtimearr_11 +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %20 = OpTypeSampler +%_runtimearr_20 = OpTypeRuntimeArray %20 +%_ptr_UniformConstant__runtimearr_20 = OpTypePointer UniformConstant %_runtimearr_20 +%_ptr_UniformConstant_20 = OpTypePointer UniformConstant %20 + %v2float = OpTypeVector %float 2 +%_ptr_Input_v2float = OpTypePointer Input %v2float + %31 = OpTypeSampledImage %11 + %v4float = OpTypeVector %float 4 + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 + %int_1 = OpConstant %int 1 + %43 = OpConstantComposite %v2int %int_1 %int_1 +%_ptr_Input_float = OpTypePointer Input %float +%_runtimearr_20_0 = OpTypeRuntimeArray %20 +%_ptr_UniformConstant__runtimearr_20_0 = OpTypePointer UniformConstant %_runtimearr_20_0 + %float_0 = OpConstant %float 0 + %v3int = OpTypeVector %int 3 +%_ptr_Input_v3int = OpTypePointer Input %v3int + %v2uint = OpTypeVector %uint 2 +%_ptr_Input_int = OpTypePointer Input %int + %int_2 = OpConstant %int 2 + %165 = OpTypeImage %float 2D 2 0 0 2 Unknown +%_runtimearr_165 = OpTypeRuntimeArray %165 +%_ptr_UniformConstant__runtimearr_165 = OpTypePointer UniformConstant %_runtimearr_165 +%_ptr_UniformConstant_165 = OpTypePointer UniformConstant %165 + %float_1 = OpConstant %float 1 + %179 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %tex_id = OpVariable %_ptr_Input_uint Input + %textures = OpVariable %_ptr_UniformConstant__runtimearr_11 UniformConstant + %samp_id = OpVariable %_ptr_Input_uint Input + %samplers = OpVariable %_ptr_UniformConstant__runtimearr_20 UniformConstant + %uv = OpVariable %_ptr_Input_v2float Input + %clamp = OpVariable %_ptr_Input_float Input + %lod = OpVariable %_ptr_Input_float Input +%cmp_samplers = OpVariable %_ptr_UniformConstant__runtimearr_20_0 UniformConstant +%cmp_texture = OpVariable %_ptr_UniformConstant_11 UniformConstant + %cmp = OpVariable %_ptr_Input_float Input + %ddx = OpVariable %_ptr_Input_v2float Input + %ddy = OpVariable %_ptr_Input_v2float Input + %location = OpVariable %_ptr_Input_v3int Input +%rw_textures = OpVariable %_ptr_UniformConstant__runtimearr_165 UniformConstant +%entryPointParam_main = OpVariable %_ptr_Output_v4float Output + %main = OpFunction %void None %3 + %4 = OpLabel + %6 = OpLoad %uint %tex_id + %15 = OpAccessChain %_ptr_UniformConstant_11 %textures %6 + %17 = OpLoad %uint %samp_id + %24 = OpAccessChain %_ptr_UniformConstant_20 %samplers %17 + %25 = OpLoad %11 %15 + %26 = OpLoad %20 %24 + %28 = OpLoad %v2float %uv +%sampledImage = OpSampledImage %31 %25 %26 + %sampled = OpImageSampleImplicitLod %v4float %sampledImage %28 None + %36 = OpLoad %11 %15 + %37 = OpLoad %20 %24 + %38 = OpLoad %v2float %uv +%sampledImage_0 = OpSampledImage %31 %36 %37 + %sampled_0 = OpImageSampleImplicitLod %v4float %sampledImage_0 %38 ConstOffset %43 + %out = OpFAdd %v4float %sampled %sampled_0 + %47 = OpLoad %11 %15 + %48 = OpLoad %20 %24 + %49 = OpLoad %v2float %uv + %50 = OpLoad %float %clamp +%sampledImage_1 = OpSampledImage %31 %47 %48 + %sampled_1 = OpImageSampleImplicitLod %v4float %sampledImage_1 %49 ConstOffset|MinLod %43 %50 + %out_0 = OpFAdd %v4float %out %sampled_1 + %57 = OpLoad %11 %15 + %58 = OpLoad %20 %24 + %59 = OpLoad %v2float %uv + %60 = OpLoad %float %lod +%sampledImage_2 = OpSampledImage %31 %57 %58 + %sampled_2 = OpImageSampleExplicitLod %v4float %sampledImage_2 %59 Lod %60 + %out_1 = OpFAdd %v4float %out_0 %sampled_2 + %66 = OpLoad %11 %15 + %67 = OpLoad %20 %24 + %68 = OpLoad %v2float %uv + %69 = OpLoad %float %lod +%sampledImage_3 = OpSampledImage %31 %66 %67 + %sampled_3 = OpImageSampleExplicitLod %v4float %sampledImage_3 %68 Lod|ConstOffset %69 %43 + %out_2 = OpFAdd %v4float %out_1 %sampled_3 + %74 = OpLoad %uint %samp_id + %78 = OpAccessChain %_ptr_UniformConstant_20 %cmp_samplers %74 + %79 = OpLoad %20 %78 + %80 = OpLoad %11 %cmp_texture + %82 = OpLoad %v2float %uv + %83 = OpLoad %float %cmp +%sampledImage_4 = OpSampledImage %31 %80 %79 + %86 = OpImageSampleDrefImplicitLod %float %sampledImage_4 %82 %83 + %87 = OpCompositeConstruct %v4float %86 %86 %86 %86 + %out_3 = OpFAdd %v4float %out_2 %87 + %89 = OpLoad %20 %78 + %90 = OpLoad %11 %cmp_texture + %91 = OpLoad %v2float %uv + %92 = OpLoad %float %cmp +%sampledImage_5 = OpSampledImage %31 %90 %89 + %94 = OpImageSampleDrefImplicitLod %float %sampledImage_5 %91 %92 ConstOffset %43 + %95 = OpCompositeConstruct %v4float %94 %94 %94 %94 + %out_4 = OpFAdd %v4float %out_3 %95 + %97 = OpLoad %20 %78 + %98 = OpLoad %11 %cmp_texture + %99 = OpLoad %v2float %uv + %100 = OpLoad %float %cmp +%sampledImage_6 = OpSampledImage %31 %98 %97 + %102 = OpImageSampleDrefExplicitLod %float %sampledImage_6 %99 %100 Lod %float_0 + %104 = OpCompositeConstruct %v4float %102 %102 %102 %102 + %out_5 = OpFAdd %v4float %out_4 %104 + %106 = OpLoad %20 %78 + %107 = OpLoad %11 %cmp_texture + %108 = OpLoad %v2float %uv + %109 = OpLoad %float %cmp +%sampledImage_7 = OpSampledImage %31 %107 %106 + %111 = OpImageSampleDrefExplicitLod %float %sampledImage_7 %108 %109 Lod|ConstOffset %float_0 %43 + %112 = OpCompositeConstruct %v4float %111 %111 %111 %111 + %out_6 = OpFAdd %v4float %out_5 %112 + %114 = OpLoad %11 %15 + %115 = OpLoad %20 %24 + %116 = OpLoad %v2float %uv + %117 = OpLoad %v2float %ddx + %119 = OpLoad %v2float %ddy +%sampledImage_8 = OpSampledImage %31 %114 %115 + %sampled_4 = OpImageSampleExplicitLod %v4float %sampledImage_8 %116 Grad %117 %119 + %out_7 = OpFAdd %v4float %out_6 %sampled_4 + %125 = OpLoad %11 %15 + %126 = OpLoad %20 %24 + %127 = OpLoad %v2float %uv + %128 = OpLoad %v2float %ddx + %129 = OpLoad %v2float %ddy +%sampledImage_9 = OpSampledImage %31 %125 %126 + %sampled_5 = OpImageSampleExplicitLod %v4float %sampledImage_9 %127 Grad|ConstOffset %128 %129 %43 + %out_8 = OpFAdd %v4float %out_7 %sampled_5 + %134 = OpLoad %11 %15 + %135 = OpLoad %20 %24 + %136 = OpLoad %v2float %uv + %137 = OpLoad %float %clamp + %138 = OpLoad %v2float %ddx + %139 = OpLoad %v2float %ddy +%sampledImage_10 = OpSampledImage %31 %134 %135 + %sampled_6 = OpImageSampleExplicitLod %v4float %sampledImage_10 %136 Grad|ConstOffset|MinLod %138 %139 %43 %137 + %out_9 = OpFAdd %v4float %out_8 %sampled_6 + %145 = OpLoad %v3int %location + %148 = OpVectorShuffle %v2int %145 %145 0 1 + %150 = OpBitcast %v2uint %148 + %151 = OpLoad %11 %15 + %sampled_7 = OpImageFetch %v4float %151 %150 + %out_10 = OpFAdd %v4float %out_9 %sampled_7 + %156 = OpAccessChain %_ptr_Input_int %location %int_2 + %158 = OpLoad %11 %15 + %159 = OpLoad %int %156 + %sampled_8 = OpImageFetch %v4float %158 %148 Lod|ConstOffset %159 %43 + %out_11 = OpFAdd %v4float %out_10 %sampled_8 + %163 = OpLoad %uint %tex_id + %169 = OpAccessChain %_ptr_UniformConstant_165 %rw_textures %163 + %170 = OpLoad %165 %169 + %sampled_9 = OpImageRead %v4float %170 %148 + %out_12 = OpFAdd %v4float %out_11 %sampled_9 + %174 = OpLoad %165 %169 + %sampled_10 = OpImageRead %v4float %174 %148 ConstOffset %43 + %out_13 = OpFAdd %v4float %out_12 %sampled_10 + %178 = OpLoad %165 %169 + OpImageWrite %178 %150 %179 + OpStore %entryPointParam_main %out_13 + OpReturn + OpFunctionEnd diff --git a/naga/tests/out/ir/binding-arrays.stress.ron b/naga/tests/out/ir/binding-arrays.stress.ron new file mode 100644 index 000000000..3ef6b3c8f --- /dev/null +++ b/naga/tests/out/ir/binding-arrays.stress.ron @@ -0,0 +1,972 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Pointer( + base: 0, + space: Private, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Float, + width: 4, + )), + ), + ( + name: None, + inner: Vector( + size: Bi, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ( + name: None, + inner: Image( + dim: D2, + arrayed: false, + class: Sampled( + kind: Float, + multi: false, + ), + ), + ), + ( + name: None, + inner: BindingArray( + base: 4, + size: Dynamic, + ), + ), + ( + name: None, + inner: Sampler( + comparison: false, + ), + ), + ( + name: None, + inner: BindingArray( + base: 6, + size: Dynamic, + ), + ), + ( + name: None, + inner: Pointer( + base: 3, + space: Private, + ), + ), + ( + name: None, + inner: Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ( + name: None, + inner: Scalar(( + kind: Sint, + width: 4, + )), + ), + ( + name: None, + inner: Vector( + size: Bi, + scalar: ( + kind: Sint, + width: 4, + ), + ), + ), + ( + name: None, + inner: Pointer( + base: 2, + space: Private, + ), + ), + ( + name: None, + inner: Vector( + size: Tri, + scalar: ( + kind: Sint, + width: 4, + ), + ), + ), + ( + name: None, + inner: Pointer( + base: 13, + space: Private, + ), + ), + ( + name: None, + inner: Vector( + size: Bi, + scalar: ( + kind: Uint, + width: 4, + ), + ), + ), + ( + name: None, + inner: Pointer( + base: 10, + space: Private, + ), + ), + ( + name: None, + inner: Pointer( + base: 9, + space: Private, + ), + ), + ( + name: None, + inner: Sampler( + comparison: true, + ), + ), + ( + name: None, + inner: Image( + dim: D2, + arrayed: false, + class: Depth( + multi: false, + ), + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + predeclared_types: {}, + ), + constants: [ + ( + name: None, + ty: 10, + init: 0, + ), + ( + name: None, + ty: 11, + init: 3, + ), + ( + name: None, + ty: 2, + init: 4, + ), + ( + name: None, + ty: 10, + init: 5, + ), + ( + name: None, + ty: 2, + init: 6, + ), + ( + name: None, + ty: 9, + init: 11, + ), + ], + overrides: [], + global_variables: [ + ( + name: Some("tex_id"), + space: Private, + binding: None, + ty: 0, + init: None, + ), + ( + name: Some("textures"), + space: Handle, + binding: Some(( + group: 0, + binding: 0, + )), + ty: 5, + init: None, + ), + ( + name: Some("samp_id"), + space: Private, + binding: None, + ty: 0, + init: None, + ), + ( + name: Some("samplers"), + space: Handle, + binding: Some(( + group: 1, + binding: 0, + )), + ty: 7, + init: None, + ), + ( + name: Some("uv"), + space: Private, + binding: None, + ty: 3, + init: None, + ), + ( + name: Some("clamp"), + space: Private, + binding: None, + ty: 2, + init: None, + ), + ( + name: Some("lod"), + space: Private, + binding: None, + ty: 2, + init: None, + ), + ( + name: Some("cmp_samplers"), + space: Handle, + binding: Some(( + group: 3, + binding: 0, + )), + ty: 18, + init: None, + ), + ( + name: Some("cmp_texture"), + space: Handle, + binding: Some(( + group: 2, + binding: 0, + )), + ty: 19, + init: None, + ), + ( + name: Some("cmp"), + space: Private, + binding: None, + ty: 2, + init: None, + ), + ( + name: Some("ddx"), + space: Private, + binding: None, + ty: 3, + init: None, + ), + ( + name: Some("ddy"), + space: Private, + binding: None, + ty: 3, + init: None, + ), + ( + name: Some("location"), + space: Private, + binding: None, + ty: 13, + init: None, + ), + ( + name: Some("rw_textures"), + space: Handle, + binding: Some(( + group: 4, + binding: 0, + )), + ty: 5, + init: None, + ), + ( + name: Some("entryPointParam_main"), + space: Private, + binding: None, + ty: 9, + init: None, + ), + ], + global_expressions: [ + Literal(I32(1)), + Constant(0), + Constant(0), + Compose( + ty: 11, + components: [ + 1, + 2, + ], + ), + Literal(F32(0.0)), + Literal(I32(2)), + Literal(F32(1.0)), + Constant(4), + Constant(4), + Constant(4), + Constant(4), + Compose( + ty: 9, + components: [ + 7, + 8, + 9, + 10, + ], + ), + Constant(1), + Constant(1), + Constant(1), + Constant(1), + Constant(1), + Constant(1), + Constant(1), + ], + functions: [ + ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(6), + GlobalVariable(14), + GlobalVariable(5), + GlobalVariable(8), + GlobalVariable(9), + GlobalVariable(0), + GlobalVariable(3), + GlobalVariable(12), + GlobalVariable(10), + GlobalVariable(1), + GlobalVariable(7), + GlobalVariable(13), + GlobalVariable(4), + GlobalVariable(11), + GlobalVariable(2), + Constant(5), + Constant(3), + Constant(2), + Constant(0), + Constant(4), + Constant(1), + Load( + pointer: 5, + ), + Access( + base: 9, + index: 21, + ), + Load( + pointer: 14, + ), + Access( + base: 6, + index: 23, + ), + Load( + pointer: 12, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 25, + array_index: None, + offset: None, + level: Auto, + depth_ref: None, + ), + Load( + pointer: 12, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 27, + array_index: None, + offset: Some(12), + level: Auto, + depth_ref: None, + ), + Binary( + op: Add, + left: 26, + right: 28, + ), + Load( + pointer: 12, + ), + Load( + pointer: 2, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 30, + array_index: None, + offset: Some(13), + level: Auto, + depth_ref: None, + ), + Binary( + op: Add, + left: 29, + right: 32, + ), + Load( + pointer: 12, + ), + Load( + pointer: 0, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 34, + array_index: None, + offset: None, + level: Exact(35), + depth_ref: None, + ), + Binary( + op: Add, + left: 33, + right: 36, + ), + Load( + pointer: 12, + ), + Load( + pointer: 0, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 38, + array_index: None, + offset: Some(14), + level: Exact(39), + depth_ref: None, + ), + Binary( + op: Add, + left: 37, + right: 40, + ), + Load( + pointer: 14, + ), + Access( + base: 10, + index: 42, + ), + Load( + pointer: 12, + ), + Load( + pointer: 4, + ), + ImageSample( + image: 3, + sampler: 43, + gather: None, + coordinate: 44, + array_index: None, + offset: None, + level: Auto, + depth_ref: Some(45), + ), + Splat( + size: Quad, + value: 46, + ), + Binary( + op: Add, + left: 41, + right: 47, + ), + Load( + pointer: 12, + ), + Load( + pointer: 4, + ), + ImageSample( + image: 3, + sampler: 43, + gather: None, + coordinate: 49, + array_index: None, + offset: Some(15), + level: Auto, + depth_ref: Some(50), + ), + Splat( + size: Quad, + value: 51, + ), + Binary( + op: Add, + left: 48, + right: 52, + ), + Load( + pointer: 12, + ), + Load( + pointer: 4, + ), + ImageSample( + image: 3, + sampler: 43, + gather: None, + coordinate: 54, + array_index: None, + offset: None, + level: Zero, + depth_ref: Some(55), + ), + Splat( + size: Quad, + value: 56, + ), + Binary( + op: Add, + left: 53, + right: 57, + ), + Load( + pointer: 12, + ), + Load( + pointer: 4, + ), + ImageSample( + image: 3, + sampler: 43, + gather: None, + coordinate: 59, + array_index: None, + offset: Some(16), + level: Zero, + depth_ref: Some(60), + ), + Splat( + size: Quad, + value: 61, + ), + Binary( + op: Add, + left: 58, + right: 62, + ), + Load( + pointer: 12, + ), + Load( + pointer: 8, + ), + Load( + pointer: 13, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 64, + array_index: None, + offset: None, + level: Gradient( + x: 65, + y: 66, + ), + depth_ref: None, + ), + Binary( + op: Add, + left: 63, + right: 67, + ), + Load( + pointer: 12, + ), + Load( + pointer: 8, + ), + Load( + pointer: 13, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 69, + array_index: None, + offset: Some(17), + level: Gradient( + x: 70, + y: 71, + ), + depth_ref: None, + ), + Binary( + op: Add, + left: 68, + right: 72, + ), + Load( + pointer: 12, + ), + Load( + pointer: 2, + ), + Load( + pointer: 8, + ), + Load( + pointer: 13, + ), + ImageSample( + image: 22, + sampler: 24, + gather: None, + coordinate: 74, + array_index: None, + offset: Some(18), + level: Gradient( + x: 76, + y: 77, + ), + depth_ref: None, + ), + Binary( + op: Add, + left: 73, + right: 78, + ), + Load( + pointer: 7, + ), + Swizzle( + size: Bi, + vector: 80, + pattern: (X, Y, X, X), + ), + As( + expr: 81, + kind: Uint, + convert: None, + ), + ImageLoad( + image: 22, + coordinate: 82, + array_index: None, + sample: None, + level: None, + ), + Binary( + op: Add, + left: 79, + right: 83, + ), + Access( + base: 7, + index: 16, + ), + Load( + pointer: 85, + ), + ImageLoad( + image: 22, + coordinate: 81, + array_index: None, + sample: None, + level: Some(86), + ), + Binary( + op: Add, + left: 84, + right: 87, + ), + Load( + pointer: 5, + ), + Access( + base: 11, + index: 89, + ), + ImageLoad( + image: 90, + coordinate: 81, + array_index: None, + sample: None, + level: None, + ), + Binary( + op: Add, + left: 88, + right: 91, + ), + ImageLoad( + image: 90, + coordinate: 81, + array_index: None, + sample: None, + level: None, + ), + Binary( + op: Add, + left: 92, + right: 93, + ), + ], + named_expressions: {}, + body: [ + Emit(( + start: 21, + end: 95, + )), + ImageStore( + image: 90, + coordinate: 82, + array_index: None, + value: 15, + ), + Store( + pointer: 1, + value: 94, + ), + Return( + value: None, + ), + ], + ), + ], + entry_points: [ + ( + name: "main", + stage: Fragment, + early_depth_test: None, + workgroup_size: (0, 0, 0), + function: ( + name: Some("main_wrap"), + arguments: [ + ( + name: Some("tex_id"), + ty: 0, + binding: Some(Location( + location: 0, + second_blend_source: false, + interpolation: Some(Flat), + sampling: None, + )), + ), + ( + name: Some("samp_id"), + ty: 0, + binding: Some(Location( + location: 1, + second_blend_source: false, + interpolation: Some(Flat), + sampling: None, + )), + ), + ( + name: Some("uv"), + ty: 3, + binding: Some(Location( + location: 2, + second_blend_source: false, + interpolation: Some(Perspective), + sampling: Some(Center), + )), + ), + ( + name: Some("location"), + ty: 13, + binding: Some(Location( + location: 3, + second_blend_source: false, + interpolation: Some(Flat), + sampling: None, + )), + ), + ( + name: Some("lod"), + ty: 2, + binding: Some(Location( + location: 4, + second_blend_source: false, + interpolation: Some(Perspective), + sampling: Some(Center), + )), + ), + ( + name: Some("clamp"), + ty: 2, + binding: Some(Location( + location: 5, + second_blend_source: false, + interpolation: Some(Perspective), + sampling: Some(Center), + )), + ), + ( + name: Some("cmp"), + ty: 2, + binding: Some(Location( + location: 6, + second_blend_source: false, + interpolation: Some(Perspective), + sampling: Some(Center), + )), + ), + ( + name: Some("ddx"), + ty: 3, + binding: Some(Location( + location: 7, + second_blend_source: false, + interpolation: Some(Perspective), + sampling: Some(Center), + )), + ), + ( + name: Some("ddy"), + ty: 3, + binding: Some(Location( + location: 8, + second_blend_source: false, + interpolation: Some(Perspective), + sampling: Some(Center), + )), + ), + ], + result: Some(( + ty: 9, + binding: Some(Location( + location: 0, + second_blend_source: false, + interpolation: None, + sampling: None, + )), + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + GlobalVariable(0), + FunctionArgument(1), + GlobalVariable(2), + FunctionArgument(2), + GlobalVariable(4), + FunctionArgument(3), + GlobalVariable(12), + FunctionArgument(4), + GlobalVariable(6), + FunctionArgument(5), + GlobalVariable(5), + FunctionArgument(6), + GlobalVariable(9), + FunctionArgument(7), + GlobalVariable(10), + FunctionArgument(8), + GlobalVariable(11), + GlobalVariable(14), + Load( + pointer: 18, + ), + ], + named_expressions: {}, + body: [ + Store( + pointer: 1, + value: 0, + ), + Store( + pointer: 3, + value: 2, + ), + Store( + pointer: 5, + value: 4, + ), + Store( + pointer: 7, + value: 6, + ), + Store( + pointer: 9, + value: 8, + ), + Store( + pointer: 11, + value: 10, + ), + Store( + pointer: 13, + value: 12, + ), + Store( + pointer: 15, + value: 14, + ), + Store( + pointer: 17, + value: 16, + ), + Call( + function: 0, + arguments: [], + result: None, + ), + Emit(( + start: 19, + end: 20, + )), + Return( + value: Some(19), + ), + ], + ), + ), + ], +) \ No newline at end of file diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 244c9f506..491481b11 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -1043,6 +1043,7 @@ fn convert_spv_all() { Targets::HLSL | Targets::WGSL | Targets::METAL, ); convert_spv("degrees", false, Targets::empty()); + convert_spv("binding-arrays.stress", true, Targets::WGSL | Targets::IR); convert_spv("binding-arrays.dynamic", true, Targets::WGSL); convert_spv("binding-arrays.static", true, Targets::WGSL); convert_spv(