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

[spv-in] sampling from depth textures returns a scalar and is splatted (#6384)

Browse Source
This commit is contained in:
Schell Carl Scivally 2024-10-12 07:24:16 +13:00 committed by GitHub
parent 1047fa57f0
commit 76b1605090
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
13 changed files with 744 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
CHANGELOG.md
View File

@ -104,6 +104,7 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
#### Naga
- SPIR-V frontend splats depth texture sample and load results. Fixes [issue #4551](https://github.com/gfx-rs/wgpu/issues/4551). By @schell in [#6384](https://github.com/gfx-rs/wgpu/pull/6384).
- Accept only `vec3` (not `vecN`) for the `cross` built-in. By @ErichDonGubler in [#6171](https://github.com/gfx-rs/wgpu/pull/6171).
- Configure `SourceLanguage` when enabling debug info in SPV-out. By @kvark in [#6256](https://github.com/gfx-rs/wgpu/pull/6256).
- Per-polygon and flat inputs should not be considered subgroup uniform. By @magcius in [#6276](https://github.com/gfx-rs/wgpu/pull/6276).

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

@ -377,36 +377,61 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
let coord_handle =
self.get_expr_handle(coordinate_id, coord_lexp, ctx, emitter, block, body_idx);
let coord_type_handle = self.lookup_type.lookup(coord_lexp.type_id)?.handle;
let (coordinate, array_index) = match ctx.type_arena[image_ty].inner {
let (coordinate, array_index, is_depth) = match ctx.type_arena[image_ty].inner {
crate::TypeInner::Image {
dim,
arrayed,
class: _,
} => extract_image_coordinates(
dim,
if arrayed {
ExtraCoordinate::ArrayLayer
} else {
ExtraCoordinate::Garbage
},
coord_handle,
coord_type_handle,
ctx,
),
class,
} => {
let (coord, array_index) = extract_image_coordinates(
dim,
if arrayed {
ExtraCoordinate::ArrayLayer
} else {
ExtraCoordinate::Garbage
},
coord_handle,
coord_type_handle,
ctx,
);
(coord, array_index, class.is_depth())
}
_ => return Err(Error::InvalidImage(image_ty)),
};
let expr = crate::Expression::ImageLoad {
let image_load_expr = crate::Expression::ImageLoad {
image: image_lexp.handle,
coordinate,
array_index,
sample,
level,
};
let image_load_handle = ctx
.expressions
.append(image_load_expr, self.span_from_with_op(start));
let handle = if is_depth {
let result_ty = self.lookup_type.lookup(result_type_id)?;
// The return type of `OpImageRead` can be a scalar or vector.
match ctx.type_arena[result_ty.handle].inner {
crate::TypeInner::Vector { size, .. } => {
let splat_expr = crate::Expression::Splat {
size,
value: image_load_handle,
};
ctx.expressions
.append(splat_expr, self.span_from_with_op(start))
}
_ => image_load_handle,
}
} else {
image_load_handle
};
self.lookup_expression.insert(
result_id,
LookupExpression {
handle: ctx.expressions.append(expr, self.span_from_with_op(start)),
handle,
type_id: result_type_id,
block_id,
},
@ -593,11 +618,12 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
ref other => return Err(Error::InvalidGlobalVar(other.clone())),
}
let ((coordinate, array_index), depth_ref) = match ctx.type_arena[image_ty].inner {
let ((coordinate, array_index), depth_ref, is_depth) = match ctx.type_arena[image_ty].inner
{
crate::TypeInner::Image {
dim,
arrayed,
class: _,
class,
} => (
extract_image_coordinates(
dim,
@ -642,6 +668,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
None => None,
}
},
class.is_depth(),
),
_ => return Err(Error::InvalidImage(image_ty)),
};
@ -656,10 +683,21 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
level,
depth_ref,
};
let image_sample_handle = ctx.expressions.append(expr, self.span_from_with_op(start));
let handle = if is_depth && depth_ref.is_none() {
let splat_expr = crate::Expression::Splat {
size: crate::VectorSize::Quad,
value: image_sample_handle,
};
ctx.expressions
.append(splat_expr, self.span_from_with_op(start))
} else {
image_sample_handle
};
self.lookup_expression.insert(
result_id,
LookupExpression {
handle: ctx.expressions.append(expr, self.span_from_with_op(start)),
handle,
type_id: result_type_id,
block_id,
},

4
naga/src/proc/mod.rs
View File

@ -622,6 +622,10 @@ impl super::ImageClass {
crate::ImageClass::Storage { .. } => false,
}
}
pub const fn is_depth(self) -> bool {
matches!(self, crate::ImageClass::Depth { .. })
}
}
impl crate::Module {

BIN
naga/tests/in/spv/fetch_depth.spv Normal file
View File

Binary file not shown.

52
naga/tests/in/spv/fetch_depth.spvasm Normal file
View File

@ -0,0 +1,52 @@
; SPIR-V
; Version: 1.5
; Generator: Google rspirv; 0
; Bound: 56
; Schema: 0
OpCapability Shader
OpCapability VulkanMemoryModel
OpMemoryModel Logical Vulkan
OpEntryPoint GLCompute %1 "cull::fetch_depth" %2 %3 %4
OpExecutionMode %1 LocalSize 32 1 1
OpDecorate %_struct_10 Block
OpMemberDecorate %_struct_10 0 Offset 0
OpDecorate %_struct_11 Block
OpMemberDecorate %_struct_11 0 Offset 0
OpDecorate %2 Binding 0
OpDecorate %2 DescriptorSet 0
OpDecorate %3 NonWritable
OpDecorate %3 Binding 1
OpDecorate %3 DescriptorSet 0
OpDecorate %4 Binding 2
OpDecorate %4 DescriptorSet 0
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%void = OpTypeVoid
%16 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%_struct_10 = OpTypeStruct %float
%_ptr_StorageBuffer__struct_10 = OpTypePointer StorageBuffer %_struct_10
%v2uint = OpTypeVector %uint 2
%_struct_11 = OpTypeStruct %v2uint
%_ptr_StorageBuffer__struct_11 = OpTypePointer StorageBuffer %_struct_11
%24 = OpTypeImage %float 2D 1 0 0 1 Unknown
%_ptr_UniformConstant_24 = OpTypePointer UniformConstant %24
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%2 = OpVariable %_ptr_StorageBuffer__struct_10 StorageBuffer
%_ptr_StorageBuffer_v2uint = OpTypePointer StorageBuffer %v2uint
%3 = OpVariable %_ptr_StorageBuffer__struct_11 StorageBuffer
%4 = OpVariable %_ptr_UniformConstant_24 UniformConstant
%1 = OpFunction %void None %16
%32 = OpLabel
%33 = OpInBoundsAccessChain %_ptr_StorageBuffer_float %2 %uint_0
%34 = OpInBoundsAccessChain %_ptr_StorageBuffer_v2uint %3 %uint_0
%35 = OpLoad %v2uint %34
%54 = OpLoad %24 %4
%55 = OpImageFetch %v4float %54 %35 Lod %int_0
%38 = OpCompositeExtract %float %55 0
OpStore %33 %38
OpReturn
OpFunctionEnd

25
naga/tests/out/hlsl/fetch_depth.hlsl Normal file
View File

@ -0,0 +1,25 @@
struct type_2 {
float member;
};
struct type_4 {
uint2 member;
};
RWByteAddressBuffer global : register(u0);
ByteAddressBuffer global_1 : register(t1);
Texture2D<float> global_2 : register(t2);
void function()
{
uint2 _e6 = asuint(global_1.Load2(0));
float _e7 = global_2.Load(int3(_e6, 0)).x;
global.Store(0, asuint((_e7).xxxx.x));
return;
}
[numthreads(32, 1, 1)]
void cullfetch_depth()
{
function();
}

12
naga/tests/out/hlsl/fetch_depth.ron Normal file
View File

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

195
naga/tests/out/ir/fetch_depth.compact.ron Normal file
View File

@ -0,0 +1,195 @@
(
types: [
(
name: None,
inner: Scalar((
kind: Float,
width: 4,
)),
),
(
name: None,
inner: Scalar((
kind: Sint,
width: 4,
)),
),
(
name: None,
inner: Struct(
members: [
(
name: None,
ty: 0,
binding: None,
offset: 0,
),
],
span: 4,
),
),
(
name: None,
inner: Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
),
),
(
name: None,
inner: Struct(
members: [
(
name: None,
ty: 3,
binding: None,
offset: 0,
),
],
span: 8,
),
),
(
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: 1,
init: 0,
),
],
overrides: [],
global_variables: [
(
name: None,
space: Storage(
access: ("LOAD | STORE"),
),
binding: Some((
group: 0,
binding: 0,
)),
ty: 2,
init: None,
),
(
name: None,
space: Storage(
access: ("LOAD"),
),
binding: Some((
group: 0,
binding: 1,
)),
ty: 4,
init: None,
),
(
name: None,
space: Handle,
binding: Some((
group: 0,
binding: 2,
)),
ty: 5,
init: None,
),
],
global_expressions: [
Literal(I32(0)),
],
functions: [
(
name: None,
arguments: [],
result: None,
local_variables: [],
expressions: [
GlobalVariable(2),
GlobalVariable(0),
GlobalVariable(1),
Constant(0),
AccessIndex(
base: 1,
index: 0,
),
AccessIndex(
base: 2,
index: 0,
),
Load(
pointer: 5,
),
ImageLoad(
image: 0,
coordinate: 6,
array_index: None,
sample: None,
level: Some(3),
),
Splat(
size: Quad,
value: 7,
),
AccessIndex(
base: 8,
index: 0,
),
],
named_expressions: {},
body: [
Emit((
start: 4,
end: 10,
)),
Store(
pointer: 4,
value: 9,
),
Return(
value: None,
),
],
),
],
entry_points: [
(
name: "cull::fetch_depth",
stage: Compute,
early_depth_test: None,
workgroup_size: (32, 1, 1),
function: (
name: Some("cull::fetch_depth_wrap"),
arguments: [],
result: None,
local_variables: [],
expressions: [],
named_expressions: {},
body: [
Call(
function: 0,
arguments: [],
result: None,
),
],
),
),
],
)

265
naga/tests/out/ir/fetch_depth.ron Normal file
View File

@ -0,0 +1,265 @@
(
types: [
(
name: None,
inner: Scalar((
kind: Uint,
width: 4,
)),
),
(
name: None,
inner: Scalar((
kind: Float,
width: 4,
)),
),
(
name: None,
inner: Vector(
size: Quad,
scalar: (
kind: Float,
width: 4,
),
),
),
(
name: None,
inner: Scalar((
kind: Sint,
width: 4,
)),
),
(
name: None,
inner: Struct(
members: [
(
name: None,
ty: 1,
binding: None,
offset: 0,
),
],
span: 4,
),
),
(
name: None,
inner: Pointer(
base: 4,
space: Storage(
access: ("LOAD | STORE"),
),
),
),
(
name: None,
inner: Vector(
size: Bi,
scalar: (
kind: Uint,
width: 4,
),
),
),
(
name: None,
inner: Struct(
members: [
(
name: None,
ty: 6,
binding: None,
offset: 0,
),
],
span: 8,
),
),
(
name: None,
inner: Pointer(
base: 7,
space: Storage(
access: ("LOAD | STORE"),
),
),
),
(
name: None,
inner: Vector(
size: Bi,
scalar: (
kind: Float,
width: 4,
),
),
),
(
name: None,
inner: Image(
dim: D2,
arrayed: false,
class: Depth(
multi: false,
),
),
),
(
name: None,
inner: Pointer(
base: 1,
space: Storage(
access: ("LOAD | STORE"),
),
),
),
(
name: None,
inner: Pointer(
base: 6,
space: Storage(
access: ("LOAD | STORE"),
),
),
),
],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [
(
name: None,
ty: 0,
init: 0,
),
(
name: None,
ty: 3,
init: 1,
),
],
overrides: [],
global_variables: [
(
name: None,
space: Storage(
access: ("LOAD | STORE"),
),
binding: Some((
group: 0,
binding: 0,
)),
ty: 4,
init: None,
),
(
name: None,
space: Storage(
access: ("LOAD"),
),
binding: Some((
group: 0,
binding: 1,
)),
ty: 7,
init: None,
),
(
name: None,
space: Handle,
binding: Some((
group: 0,
binding: 2,
)),
ty: 10,
init: None,
),
],
global_expressions: [
Literal(U32(0)),
Literal(I32(0)),
],
functions: [
(
name: None,
arguments: [],
result: None,
local_variables: [],
expressions: [
GlobalVariable(2),
GlobalVariable(0),
GlobalVariable(1),
Constant(1),
Constant(0),
AccessIndex(
base: 1,
index: 0,
),
AccessIndex(
base: 2,
index: 0,
),
Load(
pointer: 6,
),
ImageLoad(
image: 0,
coordinate: 7,
array_index: None,
sample: None,
level: Some(3),
),
Splat(
size: Quad,
value: 8,
),
AccessIndex(
base: 9,
index: 0,
),
],
named_expressions: {},
body: [
Emit((
start: 5,
end: 11,
)),
Store(
pointer: 5,
value: 10,
),
Return(
value: None,
),
],
),
],
entry_points: [
(
name: "cull::fetch_depth",
stage: Compute,
early_depth_test: None,
workgroup_size: (32, 1, 1),
function: (
name: Some("cull::fetch_depth_wrap"),
arguments: [],
result: None,
local_variables: [],
expressions: [],
named_expressions: {},
body: [
Call(
function: 0,
arguments: [],
result: None,
),
],
),
),
],
)

31
naga/tests/out/msl/fetch_depth.msl Normal file
View File

@ -0,0 +1,31 @@
// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
struct type_2 {
float member;
};
struct type_4 {
metal::uint2 member;
};
void function(
device type_2& global,
device type_4 const& global_1,
metal::depth2d<float, metal::access::sample> global_2
) {
metal::uint2 _e6 = global_1.member;
float _e7 = global_2.read(metal::uint2(_e6), 0);
global.member = metal::float4(_e7).x;
return;
}
kernel void cullfetch_depth(
device type_2& global [[user(fake0)]]
, device type_4 const& global_1 [[user(fake0)]]
, metal::depth2d<float, metal::access::sample> global_2 [[user(fake0)]]
) {
function(global, global_1, global_2);
}

74
naga/tests/out/spv/fetch_depth.spvasm Normal file
View File

@ -0,0 +1,74 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 46
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %40 "cull::fetch_depth"
OpExecutionMode %40 LocalSize 32 1 1
OpMemberDecorate %5 0 Offset 0
OpMemberDecorate %8 0 Offset 0
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 0
OpDecorate %12 Block
OpMemberDecorate %12 0 Offset 0
OpDecorate %14 NonWritable
OpDecorate %14 DescriptorSet 0
OpDecorate %14 Binding 1
OpDecorate %15 Block
OpMemberDecorate %15 0 Offset 0
OpDecorate %17 DescriptorSet 0
OpDecorate %17 Binding 2
%2 = OpTypeVoid
%3 = OpTypeFloat 32
%4 = OpTypeInt 32 1
%5 = OpTypeStruct %3
%7 = OpTypeInt 32 0
%6 = OpTypeVector %7 2
%8 = OpTypeStruct %6
%9 = OpTypeImage %3 2D 1 0 0 1 Unknown
%10 = OpConstant %4 0
%12 = OpTypeStruct %5
%13 = OpTypePointer StorageBuffer %12
%11 = OpVariable %13 StorageBuffer
%15 = OpTypeStruct %8
%16 = OpTypePointer StorageBuffer %15
%14 = OpVariable %16 StorageBuffer
%18 = OpTypePointer UniformConstant %9
%17 = OpVariable %18 UniformConstant
%21 = OpTypeFunction %2
%22 = OpTypePointer StorageBuffer %5
%23 = OpConstant %7 0
%25 = OpTypePointer StorageBuffer %8
%29 = OpTypePointer StorageBuffer %3
%30 = OpTypePointer StorageBuffer %6
%33 = OpTypeVector %3 4
%20 = OpFunction %2 None %21
%19 = OpLabel
%24 = OpAccessChain %22 %11 %23
%26 = OpAccessChain %25 %14 %23
%27 = OpLoad %9 %17
OpBranch %28
%28 = OpLabel
%31 = OpAccessChain %30 %26 %23
%32 = OpLoad %6 %31
%34 = OpImageFetch %33 %27 %32 Lod %10
%35 = OpCompositeExtract %3 %34 0
%36 = OpCompositeConstruct %33 %35 %35 %35 %35
%37 = OpCompositeExtract %3 %36 0
%38 = OpAccessChain %29 %24 %23
OpStore %38 %37
OpReturn
OpFunctionEnd
%40 = OpFunction %2 None %21
%39 = OpLabel
%41 = OpAccessChain %22 %11 %23
%42 = OpAccessChain %25 %14 %23
%43 = OpLoad %9 %17
OpBranch %44
%44 = OpLabel
%45 = OpFunctionCall %2 %20
OpReturn
OpFunctionEnd

26
naga/tests/out/wgsl/fetch_depth.wgsl Normal file
View File

@ -0,0 +1,26 @@
struct type_2 {
member: f32,
}
struct type_4 {
member: vec2<u32>,
}
@group(0) @binding(0)
var<storage, read_write> global: type_2;
@group(0) @binding(1)
var<storage> global_1: type_4;
@group(0) @binding(2)
var global_2: texture_depth_2d;
fn function() {
let _e6 = global_1.member;
let _e7 = textureLoad(global_2, _e6, 0i);
global.member = vec4(_e7).x;
return;
}
@compute @workgroup_size(32, 1, 1)
fn cullfetch_depth() {
function();
}

6
naga/tests/snapshots.rs
View File

@ -1063,11 +1063,11 @@ fn convert_spv_all() {
false,
Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
);
convert_spv("atomic_i_increment", false, Targets::IR);
convert_spv(
"atomic_i_increment",
"fetch_depth",
false,
// TODO(@schell): remove Targets::NO_VALIDATION when OpAtomicIIncrement lands
Targets::IR | Targets::NO_VALIDATION,
Targets::IR | Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL,
);
}