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

Slang Binding Array Test

Browse Source
This commit is contained in:
Connor Fitzgerald 2024-10-20 23:50:47 -04:00
parent e06f10e0a4
commit 57d787eb53
No known key found for this signature in database
GPG Key ID: 7D3FA445BAD3008D
7 changed files with 1315 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

26
naga/src/front/spv/image.rs
View File

@ -30,6 +30,16 @@ impl<'function> super::BlockContext<'function> {
match self.expressions[handle] { match self.expressions[handle] {
crate::Expression::GlobalVariable(handle) => Ok(self.global_arena[handle].ty), crate::Expression::GlobalVariable(handle) => Ok(self.global_arena[handle].ty),
crate::Expression::FunctionArgument(i) => Ok(self.arguments[i as usize].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())), ref other => Err(Error::InvalidImageExpression(other.clone())),
} }
} }
@ -210,6 +220,22 @@ pub(super) fn patch_comparison_type(
arrayed, arrayed,
}, },
crate::TypeInner::Sampler { .. } => crate::TypeInner::Sampler { comparison: true }, 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), ref other => unreachable!("Unexpected type for comparison mutation: {:?}", other),
}; };

3
naga/tests/in/spv/binding-arrays.stress.param.ron Normal file
View File

@ -0,0 +1,3 @@
(
god_mode: true,
)

44
naga/tests/in/spv/binding-arrays.stress.slang Normal file
View File

@ -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;
};

BIN
naga/tests/in/spv/binding-arrays.stress.spv Normal file
View File

Binary file not shown.

269
naga/tests/in/spv/binding-arrays.stress.spvasm Normal file
View File

@ -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

972
naga/tests/out/ir/binding-arrays.stress.ron Normal file
View File

@ -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),
),
],
),
),
],
)

1
naga/tests/snapshots.rs
View File

@ -1043,6 +1043,7 @@ fn convert_spv_all() {
Targets::HLSL | Targets::WGSL | Targets::METAL, Targets::HLSL | Targets::WGSL | Targets::METAL,
); );
convert_spv("degrees", false, Targets::empty()); 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.dynamic", true, Targets::WGSL);
convert_spv("binding-arrays.static", true, Targets::WGSL); convert_spv("binding-arrays.static", true, Targets::WGSL);
convert_spv( convert_spv(