github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

[spv-out] implement array value indexing (#723)

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-04-16 17:11:18 -04:00
committed by GitHub
parent daf4c82376
commit bb716f9c10
16 changed files with 188 additions and 37 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
src/back/spv/instructions.rs
View File

@@ -433,12 +433,12 @@ impl super::Instruction {
}
pub(super) fn store(
pointer_type_id: Word,
pointer_id: Word,
object_id: Word,
memory_access: Option<spirv::MemoryAccess>,
) -> Self {
let mut instruction = Self::new(Op::Store);
instruction.add_operand(pointer_type_id);
instruction.add_operand(pointer_id);
instruction.add_operand(object_id);
if let Some(memory_access) = memory_access {

112
src/back/spv/writer.rs
View File

@@ -70,6 +70,7 @@ struct Function {
signature: Option<Instruction>,
parameters: Vec<Instruction>,
variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
internal_variables: Vec<LocalVariable>,
blocks: Vec<Block>,
entry_point_context: Option<EntryPointContext>,
}
@@ -86,6 +87,9 @@ impl Function {
for local_var in self.variables.values() {
local_var.instruction.to_words(sink);
}
for internal_var in self.internal_variables.iter() {
internal_var.instruction.to_words(sink);
}
}
for instruction in block.body.iter() {
instruction.to_words(sink);
@@ -339,6 +343,20 @@ impl Writer {
}
}
fn get_expression_type_id(
&mut self,
arena: &Arena<crate::Type>,
tr: &TypeResolution,
) -> Result<Word, Error> {
let lookup_ty = match *tr {
TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
TypeResolution::Value(ref inner) => {
LookupType::Local(self.physical_layout.make_local(inner).unwrap())
}
};
self.get_type_id(arena, lookup_ty)
}
fn get_pointer_id(
&mut self,
arena: &Arena<crate::Type>,
@@ -649,11 +667,6 @@ impl Writer {
};
self.check(exec_model.required_capabilities())?;
if self.flags.contains(WriterFlags::DEBUG) {
self.debugs
.push(Instruction::name(function_id, &entry_point.name));
}
Ok(Instruction::entry_point(
exec_model,
function_id,
@@ -1288,23 +1301,74 @@ impl Writer {
})
}
/// Cache an expression for a value.
fn cache_expression_value<'a>(
#[allow(clippy::too_many_arguments)]
fn promote_access_expression_to_variable(
&mut self,
ir_module: &'a crate::Module,
ir_types: &Arena<crate::Type>,
result_type_id: Word,
container_id: Word,
container_resolution: &TypeResolution,
index_id: Word,
element_ty: Handle<crate::Type>,
block: &mut Block,
) -> Result<(Word, LocalVariable), Error> {
let container_type_id = self.get_expression_type_id(ir_types, container_resolution)?;
let pointer_type_id = self.id_gen.next();
Instruction::type_pointer(
pointer_type_id,
spirv::StorageClass::Function,
container_type_id,
)
.to_words(&mut self.logical_layout.declarations);
let variable = {
let id = self.id_gen.next();
LocalVariable {
id,
instruction: Instruction::variable(
pointer_type_id,
id,
spirv::StorageClass::Function,
None,
),
}
};
block
.body
.push(Instruction::store(variable.id, container_id, None));
let element_pointer_id = self.id_gen.next();
let element_pointer_type_id =
self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?;
block.body.push(Instruction::access_chain(
element_pointer_type_id,
element_pointer_id,
variable.id,
&[index_id],
));
let id = self.id_gen.next();
block.body.push(Instruction::load(
result_type_id,
id,
element_pointer_id,
None,
));
Ok((id, variable))
}
/// Cache an expression for a value.
fn cache_expression_value(
&mut self,
ir_module: &crate::Module,
ir_function: &crate::Function,
fun_info: &FunctionInfo,
expr_handle: Handle<crate::Expression>,
block: &mut Block,
function: &mut Function,
) -> Result<(), Error> {
let result_lookup_ty = match fun_info[expr_handle].ty {
TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
TypeResolution::Value(ref inner) => {
LookupType::Local(self.physical_layout.make_local(inner).unwrap())
}
};
let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?;
let result_type_id =
self.get_expression_type_id(&ir_module.types, &fun_info[expr_handle].ty)?;
let id = match ir_function.expressions[expr_handle] {
crate::Expression::Access { base, index } => {
@@ -1318,10 +1382,10 @@ impl Writer {
0
} else {
let index_id = self.cached[index];
let base_id = self.cached[base];
match *fun_info[base].ty.inner_with(&ir_module.types) {
crate::TypeInner::Vector { .. } => {
let id = self.id_gen.next();
let base_id = self.cached[base];
block.body.push(Instruction::vector_extract_dynamic(
result_type_id,
id,
@@ -1330,7 +1394,21 @@ impl Writer {
));
id
}
//TODO: support `crate::TypeInner::Array { .. }` ?
crate::TypeInner::Array {
base: ty_element, ..
} => {
let (id, variable) = self.promote_access_expression_to_variable(
&ir_module.types,
result_type_id,
base_id,
&fun_info[base].ty,
index_id,
ty_element,
block,
)?;
function.internal_variables.push(variable);
id
}
ref other => {
log::error!("Unable to access {:?}", other);
return Err(Error::FeatureNotImplemented("access for type"));

7
tests/in/access.param.ron Normal file
View File

@@ -0,0 +1,7 @@
(
spv_version: (1, 1),
spv_capabilities: [ Shader, Image1D, Sampled1D ],
spv_debug: true,
spv_adjust_coordinate_space: false,
msl_custom: false,
)

8
tests/in/access.wgsl Normal file
View File

@@ -0,0 +1,8 @@
// This snapshot tests accessing various containers, dereferencing pointers.
[[stage(vertex)]]
fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
let array = array<i32, 5>(1, 2, 3, 4, 5);
let value = array[vi];
return vec4<f32>(vec4<i32>(value));
}

2
tests/in/interpolate.param.ron
View File

@@ -4,5 +4,5 @@
spv_debug: true,
spv_adjust_coordinate_space: true,
msl_custom: false,
glsl_desktop_version: Some(400)
glsl_desktop_version: Some(400)
)

6
tests/in/operators.wgsl
View File

@@ -1,6 +1,6 @@
[[stage(vertex)]]
fn splat() -> [[builtin(position)]] vec4<f32> {
let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0;
let b = vec4<i32>(5) % 2;
return a.xyxy + vec4<f32>(b);
let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0;
let b = vec4<i32>(5) % 2;
return a.xyxy + vec4<f32>(b);
}

12
tests/in/shadow.param.ron
View File

@@ -1,8 +1,8 @@
(
spv_flow_dump_prefix: "",
spv_version: (1, 2),
spv_capabilities: [ Shader ],
spv_debug: true,
spv_adjust_coordinate_space: true,
msl_custom: false,
spv_flow_dump_prefix: "",
spv_version: (1, 2),
spv_capabilities: [ Shader ],
spv_debug: true,
spv_adjust_coordinate_space: true,
msl_custom: false,
)

15
tests/out/access.msl Normal file
View File

@@ -0,0 +1,15 @@
#include <metal_stdlib>
#include <simd/simd.h>
typedef int type3[5];
struct fooInput {
};
struct fooOutput {
metal::float4 member [[position]];
};
vertex fooOutput foo(
metal::uint vi [[vertex_id]]
) {
return fooOutput { static_cast<float4>(type3 {1, 2, 3, 4, 5}[vi]) };
}

50
tests/out/access.spvasm Normal file
View File

@@ -0,0 +1,50 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 31
OpCapability Image1D
OpCapability Shader
OpCapability Sampled1D
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %19 "foo" %14 %17
OpSource GLSL 450
OpName %14 "vi"
OpName %19 "foo"
OpDecorate %12 ArrayStride 4
OpDecorate %14 BuiltIn VertexIndex
OpDecorate %17 BuiltIn Position
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 5
%5 = OpConstant %4 1
%6 = OpConstant %4 2
%7 = OpConstant %4 3
%8 = OpConstant %4 4
%9 = OpTypeInt 32 0
%11 = OpTypeFloat 32
%10 = OpTypeVector %11 4
%12 = OpTypeArray %4 %3
%15 = OpTypePointer Input %9
%14 = OpVariable %15 Input
%18 = OpTypePointer Output %10
%17 = OpVariable %18 Output
%20 = OpTypeFunction %2
%23 = OpTypePointer Function %12
%26 = OpTypePointer Function %4
%28 = OpTypeVector %4 4
%19 = OpFunction %2 None %20
%13 = OpLabel
%24 = OpVariable %23 Function
%16 = OpLoad %9 %14
OpBranch %21
%21 = OpLabel
%22 = OpCompositeConstruct %12 %5 %6 %7 %8 %3
OpStore %24 %22
%25 = OpAccessChain %26 %24 %16
%27 = OpLoad %4 %25
%29 = OpCompositeConstruct %28 %27 %27 %27 %27
%30 = OpConvertSToF %10 %29
OpStore %17 %30
OpReturn
OpFunctionEnd

1
tests/out/boids.spvasm
View File

@@ -38,7 +38,6 @@ OpName %36 "vel"
OpName %37 "i"
OpName %40 "global_invocation_id"
OpName %43 "main"
OpName %43 "main"
OpMemberDecorate %16 0 Offset 0
OpMemberDecorate %16 1 Offset 8
OpDecorate %17 Block

1
tests/out/collatz.spvasm
View File

@@ -17,7 +17,6 @@ OpName %15 "i"
OpName %18 "collatz_iterations"
OpName %45 "global_id"
OpName %48 "main"
OpName %48 "main"
OpDecorate %8 ArrayStride 4
OpDecorate %9 Block
OpMemberDecorate %9 0 Offset 0

2
tests/out/interpolate.spvasm
View File

@@ -25,7 +25,6 @@ OpName %33 "centroid"
OpName %35 "sample"
OpName %37 "perspective"
OpName %38 "main"
OpName %38 "main"
OpName %76 "position"
OpName %79 "flat"
OpName %82 "linear"
@@ -33,7 +32,6 @@ OpName %85 "centroid"
OpName %88 "sample"
OpName %91 "perspective"
OpName %93 "main"
OpName %93 "main"
OpMemberDecorate %23 0 Offset 0
OpMemberDecorate %23 1 Offset 16
OpMemberDecorate %23 2 Offset 20

2
tests/out/quad.spvasm
View File

@@ -21,10 +21,8 @@ OpName %22 "uv"
OpName %24 "uv"
OpName %26 "position"
OpName %28 "main"
OpName %28 "main"
OpName %48 "uv"
OpName %51 "main"
OpName %51 "main"
OpMemberDecorate %9 0 Offset 0
OpMemberDecorate %9 1 Offset 16
OpDecorate %12 DescriptorSet 0

1
tests/out/shadow.spvasm
View File

@@ -29,7 +29,6 @@ OpName %71 "i"
OpName %74 "raw_normal"
OpName %77 "position"
OpName %82 "fs_main"
OpName %82 "fs_main"
OpDecorate %14 Block
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %17 0 Offset 0

1
tests/out/texture-array.spvasm
View File

@@ -16,7 +16,6 @@ OpName %14 "sampler"
OpName %16 "pc"
OpName %19 "tex_coord"
OpName %24 "main"
OpName %24 "main"
OpDecorate %8 Block
OpMemberDecorate %8 0 Offset 0
OpDecorate %11 DescriptorSet 0

1
tests/snapshots.rs
View File

@@ -229,6 +229,7 @@ fn convert_wgsl() {
"interpolate",
Targets::SPIRV | Targets::METAL | Targets::GLSL,
),
("access", Targets::SPIRV | Targets::METAL),
];
for &(name, targets) in inputs.iter() {