mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-22 23:04:07 +00:00
hlsl: storage writes
This commit is contained in:
Dzmitry Malyshau
Dzmitry Malyshau
parent
4b6846d5da
commit
284cfcce47
@ -24,7 +24,7 @@ Back-end | Status | Feature | Notes |
|
||||
SPIR-V | :white_check_mark: | spv-out | |
|
||||
WGSL | :ok: | wgsl-out | |
|
||||
Metal | :white_check_mark: | msl-out | |
|
||||
HLSL | :construction: | hlsl-out | Shader Model 5.0+ (DirectX 11+) |
|
||||
HLSL | :ok: | hlsl-out | Shader Model 5.0+ (DirectX 11+) |
|
||||
GLSL | :ok: | glsl-out | |
|
||||
AIR | | | |
|
||||
DXIL/DXIR | | | |
|
||||
|
||||
@ -10,6 +10,7 @@ use crate::{
|
||||
use std::{fmt, mem};
|
||||
|
||||
const LOCATION_SEMANTIC: &str = "LOC";
|
||||
const STORE_TEMP_NAME: &str = "_value";
|
||||
|
||||
/// Shorthand result used internally by the backend
|
||||
pub(super) type BackendResult = Result<(), Error>;
|
||||
@ -28,6 +29,20 @@ enum SubAccess {
|
||||
},
|
||||
}
|
||||
|
||||
enum StoreValue {
|
||||
Expression(Handle<crate::Expression>),
|
||||
TempIndex {
|
||||
depth: usize,
|
||||
index: u32,
|
||||
ty: proc::TypeResolution,
|
||||
},
|
||||
TempAccess {
|
||||
depth: usize,
|
||||
base: Handle<crate::Type>,
|
||||
member_index: u32,
|
||||
},
|
||||
}
|
||||
|
||||
/// Structure contains information required for generating
|
||||
/// wrapped structure of all entry points arguments
|
||||
struct EntryPointBinding {
|
||||
@ -521,11 +536,7 @@ impl<'a, W: fmt::Write> Writer<'a, W> {
|
||||
let size = module.constants[const_handle].to_array_length().unwrap();
|
||||
write!(self.out, "{}", size)?;
|
||||
}
|
||||
crate::ArraySize::Dynamic => {
|
||||
//TODO: https://github.com/gfx-rs/naga/issues/1127
|
||||
log::warn!("Dynamically sized arrays are not properly supported yet");
|
||||
write!(self.out, "1")?;
|
||||
}
|
||||
crate::ArraySize::Dynamic => unreachable!(),
|
||||
}
|
||||
|
||||
write!(self.out, "]")?;
|
||||
@ -965,7 +976,14 @@ impl<'a, W: fmt::Write> Writer<'a, W> {
|
||||
.pointer_class()
|
||||
== Some(crate::StorageClass::Storage)
|
||||
{
|
||||
return Err(Error::Unimplemented("Storage stores".to_string()));
|
||||
let var_handle = self.fill_access_chain(module, pointer, func_ctx)?;
|
||||
self.write_storage_store(
|
||||
module,
|
||||
var_handle,
|
||||
StoreValue::Expression(value),
|
||||
func_ctx,
|
||||
indent,
|
||||
)?;
|
||||
} else if let Some((const_handle, base_ty)) = array_info {
|
||||
let size = module.constants[const_handle].to_array_length().unwrap();
|
||||
writeln!(self.out, "{}{{", INDENT.repeat(indent))?;
|
||||
@ -1813,7 +1831,7 @@ impl<'a, W: fmt::Write> Writer<'a, W> {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Helper function to write down the load operation on a `ByteAddressBuffer`.
|
||||
/// Helper function to write down the Load operation on a `ByteAddressBuffer`.
|
||||
fn write_storage_load(
|
||||
&mut self,
|
||||
module: &Module,
|
||||
@ -1895,6 +1913,203 @@ impl<'a, W: fmt::Write> Writer<'a, W> {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
fn write_store_value(
|
||||
&mut self,
|
||||
module: &Module,
|
||||
value: &StoreValue,
|
||||
func_ctx: &back::FunctionCtx,
|
||||
) -> BackendResult {
|
||||
match *value {
|
||||
StoreValue::Expression(expr) => self.write_expr(module, expr, &func_ctx)?,
|
||||
StoreValue::TempIndex {
|
||||
depth,
|
||||
index,
|
||||
ty: _,
|
||||
} => write!(self.out, "{}{}[{}]", STORE_TEMP_NAME, depth, index)?,
|
||||
StoreValue::TempAccess {
|
||||
depth,
|
||||
base,
|
||||
member_index,
|
||||
} => {
|
||||
let name = &self.names[&NameKey::StructMember(base, member_index)];
|
||||
write!(self.out, "{}{}.{}", STORE_TEMP_NAME, depth, name)?
|
||||
}
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Helper function to write down the Store operation on a `ByteAddressBuffer`.
|
||||
fn write_storage_store(
|
||||
&mut self,
|
||||
module: &Module,
|
||||
var_handle: Handle<crate::GlobalVariable>,
|
||||
value: StoreValue,
|
||||
func_ctx: &back::FunctionCtx<'_>,
|
||||
indent: usize,
|
||||
) -> BackendResult {
|
||||
let temp_resolution;
|
||||
let ty_resolution = match value {
|
||||
StoreValue::Expression(expr) => &func_ctx.info[expr].ty,
|
||||
StoreValue::TempIndex {
|
||||
depth: _,
|
||||
index: _,
|
||||
ref ty,
|
||||
} => ty,
|
||||
StoreValue::TempAccess {
|
||||
depth: _,
|
||||
base,
|
||||
member_index,
|
||||
} => {
|
||||
let ty_handle = match module.types[base].inner {
|
||||
TypeInner::Struct { ref members, .. } => members[member_index as usize].ty,
|
||||
_ => unreachable!(),
|
||||
};
|
||||
temp_resolution = proc::TypeResolution::Handle(ty_handle);
|
||||
&temp_resolution
|
||||
}
|
||||
};
|
||||
match *ty_resolution.inner_with(&module.types) {
|
||||
TypeInner::Scalar { .. } => {
|
||||
// working around the borrow checker in `self.write_expr`
|
||||
let chain = mem::take(&mut self.temp_access_chain);
|
||||
let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
|
||||
write!(
|
||||
self.out,
|
||||
"{}{}.Store(",
|
||||
back::INDENT.repeat(indent),
|
||||
var_name
|
||||
)?;
|
||||
self.write_storage_address(module, &chain, func_ctx)?;
|
||||
write!(self.out, ", asuint(")?;
|
||||
self.write_store_value(module, &value, func_ctx)?;
|
||||
writeln!(self.out, "));")?;
|
||||
self.temp_access_chain = chain;
|
||||
}
|
||||
TypeInner::Vector { size, .. } => {
|
||||
// working around the borrow checker in `self.write_expr`
|
||||
let chain = mem::take(&mut self.temp_access_chain);
|
||||
let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
|
||||
write!(
|
||||
self.out,
|
||||
"{}{}.Store{}(",
|
||||
back::INDENT.repeat(indent),
|
||||
var_name,
|
||||
size as u8
|
||||
)?;
|
||||
self.write_storage_address(module, &chain, func_ctx)?;
|
||||
write!(self.out, ", asuint(")?;
|
||||
self.write_store_value(module, &value, func_ctx)?;
|
||||
writeln!(self.out, "));")?;
|
||||
self.temp_access_chain = chain;
|
||||
}
|
||||
TypeInner::Matrix {
|
||||
columns,
|
||||
rows,
|
||||
width,
|
||||
} => {
|
||||
// first, assign the value to a temporary
|
||||
writeln!(self.out, "{}{{", back::INDENT.repeat(indent))?;
|
||||
let depth = indent + 1;
|
||||
write!(
|
||||
self.out,
|
||||
"{}{}{}x{} {}{} = ",
|
||||
back::INDENT.repeat(indent + 1),
|
||||
crate::ScalarKind::Float.to_hlsl_str(width)?,
|
||||
columns as u8,
|
||||
rows as u8,
|
||||
STORE_TEMP_NAME,
|
||||
depth,
|
||||
)?;
|
||||
self.write_store_value(module, &value, func_ctx)?;
|
||||
writeln!(self.out, ";")?;
|
||||
// then iterate the stores
|
||||
let row_stride = width as u32 * rows as u32;
|
||||
for i in 0..columns as u32 {
|
||||
self.temp_access_chain
|
||||
.push(SubAccess::Offset(i * row_stride));
|
||||
let ty_inner = TypeInner::Vector {
|
||||
size: rows,
|
||||
kind: crate::ScalarKind::Float,
|
||||
width,
|
||||
};
|
||||
let sv = StoreValue::TempIndex {
|
||||
depth,
|
||||
index: i,
|
||||
ty: proc::TypeResolution::Value(ty_inner),
|
||||
};
|
||||
self.write_storage_store(module, var_handle, sv, func_ctx, indent + 1)?;
|
||||
self.temp_access_chain.pop();
|
||||
}
|
||||
// done
|
||||
writeln!(self.out, "{}}}", back::INDENT.repeat(indent))?;
|
||||
}
|
||||
TypeInner::Array {
|
||||
base,
|
||||
size: crate::ArraySize::Constant(const_handle),
|
||||
..
|
||||
} => {
|
||||
// first, assign the value to a temporary
|
||||
writeln!(self.out, "{}{{", back::INDENT.repeat(indent))?;
|
||||
write!(self.out, "{}", back::INDENT.repeat(indent + 1))?;
|
||||
self.write_value_type(module, &module.types[base].inner)?;
|
||||
let depth = indent + 1;
|
||||
write!(self.out, " {}{}", STORE_TEMP_NAME, depth)?;
|
||||
self.write_array_size(module, crate::ArraySize::Constant(const_handle))?;
|
||||
write!(self.out, " = ")?;
|
||||
self.write_store_value(module, &value, func_ctx)?;
|
||||
writeln!(self.out, ";")?;
|
||||
// then iterate the stores
|
||||
let count = module.constants[const_handle].to_array_length().unwrap();
|
||||
let stride = module.types[base].inner.span(&module.constants);
|
||||
for i in 0..count {
|
||||
self.temp_access_chain.push(SubAccess::Offset(i * stride));
|
||||
let sv = StoreValue::TempIndex {
|
||||
depth,
|
||||
index: i,
|
||||
ty: proc::TypeResolution::Handle(base),
|
||||
};
|
||||
self.write_storage_store(module, var_handle, sv, func_ctx, indent + 1)?;
|
||||
self.temp_access_chain.pop();
|
||||
}
|
||||
// done
|
||||
writeln!(self.out, "{}}}", back::INDENT.repeat(indent))?;
|
||||
}
|
||||
TypeInner::Struct { ref members, .. } => {
|
||||
// first, assign the value to a temporary
|
||||
writeln!(self.out, "{}{{", back::INDENT.repeat(indent))?;
|
||||
let depth = indent + 1;
|
||||
let struct_ty = ty_resolution.handle().unwrap();
|
||||
let struct_name = &self.names[&NameKey::Type(struct_ty)];
|
||||
write!(
|
||||
self.out,
|
||||
"{}{} {}{} = ",
|
||||
back::INDENT.repeat(indent + 1),
|
||||
struct_name,
|
||||
STORE_TEMP_NAME,
|
||||
depth
|
||||
)?;
|
||||
self.write_store_value(module, &value, func_ctx)?;
|
||||
writeln!(self.out, ";")?;
|
||||
// then iterate the stores
|
||||
for (i, member) in members.iter().enumerate() {
|
||||
self.temp_access_chain
|
||||
.push(SubAccess::Offset(member.offset));
|
||||
let sv = StoreValue::TempAccess {
|
||||
depth,
|
||||
base: struct_ty,
|
||||
member_index: i as u32,
|
||||
};
|
||||
self.write_storage_store(module, var_handle, sv, func_ctx, indent + 1)?;
|
||||
self.temp_access_chain.pop();
|
||||
}
|
||||
// done
|
||||
writeln!(self.out, "{}}}", back::INDENT.repeat(indent))?;
|
||||
}
|
||||
_ => unreachable!(),
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
fn fill_access_chain(
|
||||
&mut self,
|
||||
module: &Module,
|
||||
|
||||
@ -17,14 +17,19 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
|
||||
let baz: f32 = foo;
|
||||
foo = 1.0;
|
||||
|
||||
// test storage loads
|
||||
let matrix = bar.matrix;
|
||||
let arr = bar.arr;
|
||||
|
||||
let index = 3u;
|
||||
let b = bar.matrix[index].x;
|
||||
|
||||
let a = bar.data[arrayLength(&bar.data) - 2u];
|
||||
|
||||
// test storage stores
|
||||
bar.matrix[1].z = 1.0;
|
||||
bar.matrix = mat4x4<f32>(vec4<f32>(0.0), vec4<f32>(1.0), vec4<f32>(2.0), vec4<f32>(3.0));
|
||||
bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
|
||||
|
||||
// test array indexing
|
||||
var c = array<i32, 5>(a, i32(b), 3, 4, 5);
|
||||
c[vi + 1u] = 42;
|
||||
let value = c[vi];
|
||||
|
||||
@ -24,6 +24,19 @@ float4 foo(VertexInput_foo vertexinput_foo) : SV_Position
|
||||
float4 _expr13 = asfloat(bar.Load4(12+0));
|
||||
float b = _expr13.x;
|
||||
int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 80) / 4) - 2u)*4+8));
|
||||
bar.Store(8+4+0, asuint(1.0));
|
||||
{
|
||||
float4x4 _value2 = float4x4(float4(0.0.xxxx), float4(1.0.xxxx), float4(2.0.xxxx), float4(3.0.xxxx));
|
||||
bar.Store4(0+0, asuint(_value2[0]));
|
||||
bar.Store4(0+16, asuint(_value2[1]));
|
||||
bar.Store4(0+32, asuint(_value2[2]));
|
||||
bar.Store4(0+48, asuint(_value2[3]));
|
||||
}
|
||||
{
|
||||
uint2 _value2[2] = { uint2(0u.xx), uint2(1u.xx) };
|
||||
bar.Store2(4+0, asuint(_value2[0]));
|
||||
bar.Store2(4+8, asuint(_value2[1]));
|
||||
}
|
||||
{
|
||||
int _result[5]={ a, int(b), 3, 4, 5 };
|
||||
for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
|
||||
|
||||
147
tests/out/hlsl/boids.hlsl
Normal file
147
tests/out/hlsl/boids.hlsl
Normal file
@ -0,0 +1,147 @@
|
||||
static const uint NUM_PARTICLES = 1500;
|
||||
|
||||
struct Particle {
|
||||
float2 pos;
|
||||
float2 vel;
|
||||
};
|
||||
|
||||
struct SimParams {
|
||||
float deltaT;
|
||||
float rule1Distance;
|
||||
float rule2Distance;
|
||||
float rule3Distance;
|
||||
float rule1Scale;
|
||||
float rule2Scale;
|
||||
float rule3Scale;
|
||||
};
|
||||
|
||||
cbuffer params : register(b0) { SimParams params; }
|
||||
ByteAddressBuffer particlesSrc : register(t1);
|
||||
RWByteAddressBuffer particlesDst : register(u2);
|
||||
|
||||
struct ComputeInput_main {
|
||||
uint3 global_invocation_id1 : SV_DispatchThreadID;
|
||||
};
|
||||
|
||||
[numthreads(64, 1, 1)]
|
||||
void main(ComputeInput_main computeinput_main)
|
||||
{
|
||||
float2 vPos = (float2)0;
|
||||
float2 vVel = (float2)0;
|
||||
float2 cMass = (float2)0;
|
||||
float2 cVel = (float2)0;
|
||||
float2 colVel = (float2)0;
|
||||
int cMassCount = 0;
|
||||
int cVelCount = 0;
|
||||
float2 pos = (float2)0;
|
||||
float2 vel = (float2)0;
|
||||
uint i = 0u;
|
||||
|
||||
uint index = computeinput_main.global_invocation_id1.x;
|
||||
if ((index >= NUM_PARTICLES)) {
|
||||
return;
|
||||
}
|
||||
float2 _expr10 = asfloat(particlesSrc.Load2(0+index*4+0));
|
||||
vPos = _expr10;
|
||||
float2 _expr15 = asfloat(particlesSrc.Load2(4+index*4+0));
|
||||
vVel = _expr15;
|
||||
cMass = float2(0.0, 0.0);
|
||||
cVel = float2(0.0, 0.0);
|
||||
colVel = float2(0.0, 0.0);
|
||||
while(true) {
|
||||
uint _expr37 = i;
|
||||
if ((_expr37 >= NUM_PARTICLES)) {
|
||||
break;
|
||||
}
|
||||
uint _expr39 = i;
|
||||
if ((_expr39 == index)) {
|
||||
continue;
|
||||
}
|
||||
uint _expr42 = i;
|
||||
float2 _expr45 = asfloat(particlesSrc.Load2(0+_expr42*4+0));
|
||||
pos = _expr45;
|
||||
uint _expr47 = i;
|
||||
float2 _expr50 = asfloat(particlesSrc.Load2(4+_expr47*4+0));
|
||||
vel = _expr50;
|
||||
float2 _expr51 = pos;
|
||||
float2 _expr52 = vPos;
|
||||
float _expr55 = params.rule1Distance;
|
||||
if ((distance(_expr51, _expr52) < _expr55)) {
|
||||
float2 _expr57 = cMass;
|
||||
float2 _expr58 = pos;
|
||||
cMass = (_expr57 + _expr58);
|
||||
int _expr60 = cMassCount;
|
||||
cMassCount = (_expr60 + 1);
|
||||
}
|
||||
float2 _expr63 = pos;
|
||||
float2 _expr64 = vPos;
|
||||
float _expr67 = params.rule2Distance;
|
||||
if ((distance(_expr63, _expr64) < _expr67)) {
|
||||
float2 _expr69 = colVel;
|
||||
float2 _expr70 = pos;
|
||||
float2 _expr71 = vPos;
|
||||
colVel = (_expr69 - (_expr70 - _expr71));
|
||||
}
|
||||
float2 _expr74 = pos;
|
||||
float2 _expr75 = vPos;
|
||||
float _expr78 = params.rule3Distance;
|
||||
if ((distance(_expr74, _expr75) < _expr78)) {
|
||||
float2 _expr80 = cVel;
|
||||
float2 _expr81 = vel;
|
||||
cVel = (_expr80 + _expr81);
|
||||
int _expr83 = cVelCount;
|
||||
cVelCount = (_expr83 + 1);
|
||||
}
|
||||
uint _expr86 = i;
|
||||
i = (_expr86 + 1u);
|
||||
}
|
||||
int _expr89 = cMassCount;
|
||||
if ((_expr89 > 0)) {
|
||||
float2 _expr92 = cMass;
|
||||
int _expr93 = cMassCount;
|
||||
float2 _expr97 = vPos;
|
||||
cMass = ((_expr92 / float2(float(_expr93).xx)) - _expr97);
|
||||
}
|
||||
int _expr99 = cVelCount;
|
||||
if ((_expr99 > 0)) {
|
||||
float2 _expr102 = cVel;
|
||||
int _expr103 = cVelCount;
|
||||
cVel = (_expr102 / float2(float(_expr103).xx));
|
||||
}
|
||||
float2 _expr107 = vVel;
|
||||
float2 _expr108 = cMass;
|
||||
float _expr110 = params.rule1Scale;
|
||||
float2 _expr113 = colVel;
|
||||
float _expr115 = params.rule2Scale;
|
||||
float2 _expr118 = cVel;
|
||||
float _expr120 = params.rule3Scale;
|
||||
vVel = (((_expr107 + (_expr108 * _expr110)) + (_expr113 * _expr115)) + (_expr118 * _expr120));
|
||||
float2 _expr123 = vVel;
|
||||
float2 _expr125 = vVel;
|
||||
vVel = (normalize(_expr123) * clamp(length(_expr125), 0.0, 0.1));
|
||||
float2 _expr131 = vPos;
|
||||
float2 _expr132 = vVel;
|
||||
float _expr134 = params.deltaT;
|
||||
vPos = (_expr131 + (_expr132 * _expr134));
|
||||
float2 _expr137 = vPos;
|
||||
if ((_expr137.x < -1.0)) {
|
||||
vPos.x = 1.0;
|
||||
}
|
||||
float2 _expr143 = vPos;
|
||||
if ((_expr143.x > 1.0)) {
|
||||
vPos.x = -1.0;
|
||||
}
|
||||
float2 _expr149 = vPos;
|
||||
if ((_expr149.y < -1.0)) {
|
||||
vPos.y = 1.0;
|
||||
}
|
||||
float2 _expr155 = vPos;
|
||||
if ((_expr155.y > 1.0)) {
|
||||
vPos.y = -1.0;
|
||||
}
|
||||
float2 _expr164 = vPos;
|
||||
particlesDst.Store2(0+index*4+0, asuint(_expr164));
|
||||
float2 _expr168 = vVel;
|
||||
particlesDst.Store2(4+index*4+0, asuint(_expr168));
|
||||
return;
|
||||
}
|
||||
3
tests/out/hlsl/boids.hlsl.config
Normal file
3
tests/out/hlsl/boids.hlsl.config
Normal file
@ -0,0 +1,3 @@
|
||||
vertex=()
|
||||
fragment=()
|
||||
compute=(main:cs_5_0 )
|
||||
41
tests/out/hlsl/collatz.hlsl
Normal file
41
tests/out/hlsl/collatz.hlsl
Normal file
@ -0,0 +1,41 @@
|
||||
|
||||
RWByteAddressBuffer v_indices : register(u0);
|
||||
|
||||
struct ComputeInput_main {
|
||||
uint3 global_id1 : SV_DispatchThreadID;
|
||||
};
|
||||
|
||||
uint collatz_iterations(uint n_base)
|
||||
{
|
||||
uint n = (uint)0;
|
||||
uint i = 0u;
|
||||
|
||||
n = n_base;
|
||||
while(true) {
|
||||
uint _expr5 = n;
|
||||
if ((_expr5 <= 1u)) {
|
||||
break;
|
||||
}
|
||||
uint _expr8 = n;
|
||||
if (((_expr8 % 2u) == 0u)) {
|
||||
uint _expr13 = n;
|
||||
n = (_expr13 / 2u);
|
||||
} else {
|
||||
uint _expr17 = n;
|
||||
n = ((3u * _expr17) + 1u);
|
||||
}
|
||||
uint _expr21 = i;
|
||||
i = (_expr21 + 1u);
|
||||
}
|
||||
uint _expr24 = i;
|
||||
return _expr24;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main(ComputeInput_main computeinput_main)
|
||||
{
|
||||
uint _expr8 = asuint(v_indices.Load(computeinput_main.global_id1.x*4+0));
|
||||
const uint _e9 = collatz_iterations(_expr8);
|
||||
v_indices.Store(computeinput_main.global_id1.x*4+0, asuint(_e9));
|
||||
return;
|
||||
}
|
||||
3
tests/out/hlsl/collatz.hlsl.config
Normal file
3
tests/out/hlsl/collatz.hlsl.config
Normal file
@ -0,0 +1,3 @@
|
||||
vertex=()
|
||||
fragment=()
|
||||
compute=(main:cs_5_0 )
|
||||
@ -38,6 +38,9 @@ vertex fooOutput foo(
|
||||
metal::float4 _e13 = bar.matrix[3];
|
||||
float b = _e13.x;
|
||||
int a = bar.data[(1 + (_buffer_sizes.size0 - 80 - 4) / 4) - 2u];
|
||||
bar.matrix[1].z = 1.0;
|
||||
bar.matrix = metal::float4x4(metal::float4(0.0), metal::float4(1.0), metal::float4(2.0), metal::float4(3.0));
|
||||
for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type2 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
|
||||
for(int _i=0; _i<5; ++_i) c.inner[_i] = type8 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
|
||||
c.inner[vi + 1u] = 42;
|
||||
int value = c.inner[vi];
|
||||
|
||||
@ -1,35 +1,35 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 67
|
||||
; Bound: 82
|
||||
OpCapability Shader
|
||||
OpExtension "SPV_KHR_storage_buffer_storage_class"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %35 "foo" %30 %33
|
||||
OpEntryPoint Vertex %39 "foo" %34 %37
|
||||
OpSource GLSL 450
|
||||
OpName %21 "Bar"
|
||||
OpMemberName %21 0 "matrix"
|
||||
OpMemberName %21 1 "arr"
|
||||
OpMemberName %21 2 "data"
|
||||
OpName %23 "bar"
|
||||
OpName %25 "foo"
|
||||
OpName %27 "c"
|
||||
OpName %30 "vi"
|
||||
OpName %35 "foo"
|
||||
OpDecorate %19 ArrayStride 8
|
||||
OpDecorate %20 ArrayStride 4
|
||||
OpDecorate %21 Block
|
||||
OpMemberDecorate %21 0 Offset 0
|
||||
OpMemberDecorate %21 0 ColMajor
|
||||
OpMemberDecorate %21 0 MatrixStride 16
|
||||
OpMemberDecorate %21 1 Offset 64
|
||||
OpMemberDecorate %21 2 Offset 80
|
||||
OpDecorate %22 ArrayStride 4
|
||||
OpDecorate %23 DescriptorSet 0
|
||||
OpDecorate %23 Binding 0
|
||||
OpDecorate %30 BuiltIn VertexIndex
|
||||
OpDecorate %33 BuiltIn Position
|
||||
OpName %25 "Bar"
|
||||
OpMemberName %25 0 "matrix"
|
||||
OpMemberName %25 1 "arr"
|
||||
OpMemberName %25 2 "data"
|
||||
OpName %27 "bar"
|
||||
OpName %29 "foo"
|
||||
OpName %31 "c"
|
||||
OpName %34 "vi"
|
||||
OpName %39 "foo"
|
||||
OpDecorate %23 ArrayStride 8
|
||||
OpDecorate %24 ArrayStride 4
|
||||
OpDecorate %25 Block
|
||||
OpMemberDecorate %25 0 Offset 0
|
||||
OpMemberDecorate %25 0 ColMajor
|
||||
OpMemberDecorate %25 0 MatrixStride 16
|
||||
OpMemberDecorate %25 1 Offset 64
|
||||
OpMemberDecorate %25 2 Offset 80
|
||||
OpDecorate %26 ArrayStride 4
|
||||
OpDecorate %27 DescriptorSet 0
|
||||
OpDecorate %27 Binding 0
|
||||
OpDecorate %34 BuiltIn VertexIndex
|
||||
OpDecorate %37 BuiltIn Position
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeInt 32 1
|
||||
%3 = OpConstant %4 2
|
||||
@ -39,66 +39,84 @@ OpDecorate %33 BuiltIn Position
|
||||
%9 = OpTypeInt 32 0
|
||||
%8 = OpConstant %9 3
|
||||
%10 = OpConstant %9 2
|
||||
%11 = OpConstant %4 5
|
||||
%12 = OpConstant %4 3
|
||||
%13 = OpConstant %4 4
|
||||
%14 = OpConstant %9 1
|
||||
%15 = OpConstant %4 42
|
||||
%17 = OpTypeVector %6 4
|
||||
%16 = OpTypeMatrix %17 4
|
||||
%18 = OpTypeVector %9 2
|
||||
%19 = OpTypeArray %18 %3
|
||||
%20 = OpTypeRuntimeArray %4
|
||||
%21 = OpTypeStruct %16 %19 %20
|
||||
%22 = OpTypeArray %4 %11
|
||||
%24 = OpTypePointer StorageBuffer %21
|
||||
%23 = OpVariable %24 StorageBuffer
|
||||
%26 = OpTypePointer Function %6
|
||||
%28 = OpTypePointer Function %22
|
||||
%31 = OpTypePointer Input %9
|
||||
%30 = OpVariable %31 Input
|
||||
%34 = OpTypePointer Output %17
|
||||
%33 = OpVariable %34 Output
|
||||
%36 = OpTypeFunction %2
|
||||
%39 = OpTypePointer StorageBuffer %16
|
||||
%40 = OpConstant %9 0
|
||||
%43 = OpTypePointer StorageBuffer %19
|
||||
%46 = OpTypePointer StorageBuffer %17
|
||||
%50 = OpTypePointer StorageBuffer %20
|
||||
%53 = OpTypePointer StorageBuffer %4
|
||||
%59 = OpTypePointer Function %4
|
||||
%63 = OpTypeVector %4 4
|
||||
%35 = OpFunction %2 None %36
|
||||
%29 = OpLabel
|
||||
%25 = OpVariable %26 Function %5
|
||||
%27 = OpVariable %28 Function
|
||||
%32 = OpLoad %9 %30
|
||||
OpBranch %37
|
||||
%37 = OpLabel
|
||||
%38 = OpLoad %6 %25
|
||||
OpStore %25 %7
|
||||
%41 = OpAccessChain %39 %23 %40
|
||||
%42 = OpLoad %16 %41
|
||||
%44 = OpAccessChain %43 %23 %14
|
||||
%45 = OpLoad %19 %44
|
||||
%47 = OpAccessChain %46 %23 %40 %8
|
||||
%48 = OpLoad %17 %47
|
||||
%49 = OpCompositeExtract %6 %48 0
|
||||
%51 = OpArrayLength %9 %23 2
|
||||
%52 = OpISub %9 %51 %10
|
||||
%54 = OpAccessChain %53 %23 %10 %52
|
||||
%55 = OpLoad %4 %54
|
||||
%56 = OpConvertFToS %4 %49
|
||||
%57 = OpCompositeConstruct %22 %55 %56 %12 %13 %11
|
||||
OpStore %27 %57
|
||||
%58 = OpIAdd %9 %32 %14
|
||||
%60 = OpAccessChain %59 %27 %58
|
||||
OpStore %60 %15
|
||||
%61 = OpAccessChain %59 %27 %32
|
||||
%62 = OpLoad %4 %61
|
||||
%64 = OpCompositeConstruct %63 %62 %62 %62 %62
|
||||
%65 = OpConvertSToF %17 %64
|
||||
%66 = OpMatrixTimesVector %17 %42 %65
|
||||
OpStore %33 %66
|
||||
%11 = OpConstant %4 1
|
||||
%12 = OpConstant %6 2.0
|
||||
%13 = OpConstant %6 3.0
|
||||
%14 = OpConstant %9 0
|
||||
%15 = OpConstant %9 1
|
||||
%16 = OpConstant %4 5
|
||||
%17 = OpConstant %4 3
|
||||
%18 = OpConstant %4 4
|
||||
%19 = OpConstant %4 42
|
||||
%21 = OpTypeVector %6 4
|
||||
%20 = OpTypeMatrix %21 4
|
||||
%22 = OpTypeVector %9 2
|
||||
%23 = OpTypeArray %22 %3
|
||||
%24 = OpTypeRuntimeArray %4
|
||||
%25 = OpTypeStruct %20 %23 %24
|
||||
%26 = OpTypeArray %4 %16
|
||||
%28 = OpTypePointer StorageBuffer %25
|
||||
%27 = OpVariable %28 StorageBuffer
|
||||
%30 = OpTypePointer Function %6
|
||||
%32 = OpTypePointer Function %26
|
||||
%35 = OpTypePointer Input %9
|
||||
%34 = OpVariable %35 Input
|
||||
%38 = OpTypePointer Output %21
|
||||
%37 = OpVariable %38 Output
|
||||
%40 = OpTypeFunction %2
|
||||
%43 = OpTypePointer StorageBuffer %20
|
||||
%46 = OpTypePointer StorageBuffer %23
|
||||
%49 = OpTypePointer StorageBuffer %21
|
||||
%53 = OpTypePointer StorageBuffer %24
|
||||
%56 = OpTypePointer StorageBuffer %4
|
||||
%59 = OpTypePointer StorageBuffer %6
|
||||
%74 = OpTypePointer Function %4
|
||||
%78 = OpTypeVector %4 4
|
||||
%39 = OpFunction %2 None %40
|
||||
%33 = OpLabel
|
||||
%29 = OpVariable %30 Function %5
|
||||
%31 = OpVariable %32 Function
|
||||
%36 = OpLoad %9 %34
|
||||
OpBranch %41
|
||||
%41 = OpLabel
|
||||
%42 = OpLoad %6 %29
|
||||
OpStore %29 %7
|
||||
%44 = OpAccessChain %43 %27 %14
|
||||
%45 = OpLoad %20 %44
|
||||
%47 = OpAccessChain %46 %27 %15
|
||||
%48 = OpLoad %23 %47
|
||||
%50 = OpAccessChain %49 %27 %14 %8
|
||||
%51 = OpLoad %21 %50
|
||||
%52 = OpCompositeExtract %6 %51 0
|
||||
%54 = OpArrayLength %9 %27 2
|
||||
%55 = OpISub %9 %54 %10
|
||||
%57 = OpAccessChain %56 %27 %10 %55
|
||||
%58 = OpLoad %4 %57
|
||||
%60 = OpAccessChain %59 %27 %14 %15 %10
|
||||
OpStore %60 %7
|
||||
%61 = OpCompositeConstruct %21 %5 %5 %5 %5
|
||||
%62 = OpCompositeConstruct %21 %7 %7 %7 %7
|
||||
%63 = OpCompositeConstruct %21 %12 %12 %12 %12
|
||||
%64 = OpCompositeConstruct %21 %13 %13 %13 %13
|
||||
%65 = OpCompositeConstruct %20 %61 %62 %63 %64
|
||||
%66 = OpAccessChain %43 %27 %14
|
||||
OpStore %66 %65
|
||||
%67 = OpCompositeConstruct %22 %14 %14
|
||||
%68 = OpCompositeConstruct %22 %15 %15
|
||||
%69 = OpCompositeConstruct %23 %67 %68
|
||||
%70 = OpAccessChain %46 %27 %15
|
||||
OpStore %70 %69
|
||||
%71 = OpConvertFToS %4 %52
|
||||
%72 = OpCompositeConstruct %26 %58 %71 %17 %18 %16
|
||||
OpStore %31 %72
|
||||
%73 = OpIAdd %9 %36 %15
|
||||
%75 = OpAccessChain %74 %31 %73
|
||||
OpStore %75 %19
|
||||
%76 = OpAccessChain %74 %31 %36
|
||||
%77 = OpLoad %4 %76
|
||||
%79 = OpCompositeConstruct %78 %77 %77 %77 %77
|
||||
%80 = OpConvertSToF %21 %79
|
||||
%81 = OpMatrixTimesVector %21 %45 %80
|
||||
OpStore %37 %81
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -20,6 +20,9 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
|
||||
let _e13: vec4<f32> = bar.matrix[3];
|
||||
let b: f32 = _e13.x;
|
||||
let a: i32 = bar.data[(arrayLength(&bar.data) - 2u)];
|
||||
bar.matrix[1][2] = 1.0;
|
||||
bar.matrix = mat4x4<f32>(vec4<f32>(0.0), vec4<f32>(1.0), vec4<f32>(2.0), vec4<f32>(3.0));
|
||||
bar.arr = array<vec2<u32>,2>(vec2<u32>(0u), vec2<u32>(1u));
|
||||
c = array<i32,5>(a, i32(b), 3, 4, 5);
|
||||
c[(vi + 1u)] = 42;
|
||||
let value: i32 = c[vi];
|
||||
|
||||
@ -396,7 +396,7 @@ fn convert_wgsl() {
|
||||
),
|
||||
(
|
||||
"boids",
|
||||
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
|
||||
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
|
||||
),
|
||||
(
|
||||
"skybox",
|
||||
@ -404,7 +404,12 @@ fn convert_wgsl() {
|
||||
),
|
||||
(
|
||||
"collatz",
|
||||
Targets::SPIRV | Targets::METAL | Targets::IR | Targets::ANALYSIS | Targets::WGSL,
|
||||
Targets::SPIRV
|
||||
| Targets::METAL
|
||||
| Targets::IR
|
||||
| Targets::ANALYSIS
|
||||
| Targets::HLSL
|
||||
| Targets::WGSL,
|
||||
),
|
||||
(
|
||||
"shadow",
|
||||
|
||||
Loading…
Reference in New Issue
Block a user