Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

hlsl: support arrays of matrices #1629

Merged
merged 1 commit into from
Dec 27, 2021
Merged

hlsl: support arrays of matrices #1629

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 4 additions & 15 deletions src/back/hlsl/writer.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -715,22 +715,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
stride: _,
} => {
// HLSL arrays are written as `type name[size]`
let (ty_name, vec_size) = match module.types[base].inner {
// Write scalar type by backend so as not to depend on the front-end implementation
// Name returned from frontend can be generated (type1, float1, etc.)
TypeInner::Scalar { kind, width } => (kind.to_hlsl_str(width)?, None),
// Similarly, write vector types directly.
TypeInner::Vector { size, kind, width } => {
(kind.to_hlsl_str(width)?, Some(size))
}
_ => (self.names[&NameKey::Type(base)].as_str(), None),
};

// Write `type` and `name`
write!(self.out, "{}", ty_name)?;
if let Some(s) = vec_size {
write!(self.out, "{}", s as usize)?;
if let TypeInner::Matrix { .. } = module.types[base].inner {
write!(self.out, "row_major ")?;
}
self.write_type(module, base)?;
// Write `name`
write!(
self.out,
" {}",
Expand Down
1 change: 1 addition & 0 deletions tests/in/access.wgsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,7 @@

struct Bar {
matrix: mat4x4<f32>;
matrix_array: array<mat2x2<f32>, 2>;
atom: atomic<i32>;
arr: [[stride(8)]] array<vec2<u32>, 2>;
data: [[stride(8)]] array<i32>;
Expand Down
1 change: 1 addition & 0 deletions tests/out/glsl/access.atomics.Compute.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,7 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

layout(std430) buffer Bar_block_0Compute {
mat4x4 matrix;
mat2x2 matrix_array[2];
int atom;
uvec2 arr[2];
int data[];
Expand Down
1 change: 1 addition & 0 deletions tests/out/glsl/access.foo.Vertex.glsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@ precision highp int;

layout(std430) buffer Bar_block_0Vertex {
mat4x4 matrix;
mat2x2 matrix_array[2];
int atom;
uvec2 arr[2];
int data[];
Expand Down
30 changes: 15 additions & 15 deletions tests/out/hlsl/access.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,9 +22,9 @@ float4 foo(uint vi : SV_VertexID) : SV_Position
float baz = foo_1;
foo_1 = 1.0;
float4x4 matrix_ = float4x4(asfloat(bar.Load4(0+0)), asfloat(bar.Load4(0+16)), asfloat(bar.Load4(0+32)), asfloat(bar.Load4(0+48)));
uint2 arr[2] = {asuint(bar.Load2(72+0)), asuint(bar.Load2(72+8))};
uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))};
float b = asfloat(bar.Load(0+48+0));
int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 88) / 8) - 2u)*8+88));
int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120));
const float _e25 = read_from_private(foo_1);
bar.Store(8+16+0, asuint(1.0));
{
Expand All @@ -36,10 +36,10 @@ float4 foo(uint vi : SV_VertexID) : SV_Position
}
{
uint2 _value2[2] = { uint2(0u.xx), uint2(1u.xx) };
bar.Store2(72+0, asuint(_value2[0]));
bar.Store2(72+8, asuint(_value2[1]));
bar.Store2(104+0, asuint(_value2[0]));
bar.Store2(104+8, asuint(_value2[1]));
}
bar.Store(8+88, asuint(1));
bar.Store(8+120, asuint(1));
{
int _result[5]={ a, int(b), 3, 4, 5 };
for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
Expand All @@ -54,23 +54,23 @@ void atomics()
{
int tmp = (int)0;

int value_1 = asint(bar.Load(64));
int _e6; bar.InterlockedAdd(64, 5, _e6);
int value_1 = asint(bar.Load(96));
int _e6; bar.InterlockedAdd(96, 5, _e6);
tmp = _e6;
int _e9; bar.InterlockedAdd(64, -5, _e9);
int _e9; bar.InterlockedAdd(96, -5, _e9);
tmp = _e9;
int _e12; bar.InterlockedAnd(64, 5, _e12);
int _e12; bar.InterlockedAnd(96, 5, _e12);
tmp = _e12;
int _e15; bar.InterlockedOr(64, 5, _e15);
int _e15; bar.InterlockedOr(96, 5, _e15);
tmp = _e15;
int _e18; bar.InterlockedXor(64, 5, _e18);
int _e18; bar.InterlockedXor(96, 5, _e18);
tmp = _e18;
int _e21; bar.InterlockedMin(64, 5, _e21);
int _e21; bar.InterlockedMin(96, 5, _e21);
tmp = _e21;
int _e24; bar.InterlockedMax(64, 5, _e24);
int _e24; bar.InterlockedMax(96, 5, _e24);
tmp = _e24;
int _e27; bar.InterlockedExchange(64, 5, _e27);
int _e27; bar.InterlockedExchange(96, 5, _e27);
tmp = _e27;
bar.Store(64, asuint(value_1));
bar.Store(96, asuint(value_1));
return;
}
26 changes: 15 additions & 11 deletions tests/out/msl/access.msl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,18 +6,22 @@ struct _mslBufferSizes {
metal::uint size0;
};

struct type_3 {
struct type_2 {
metal::float2x2 inner[2];
};
struct type_5 {
metal::uint2 inner[2];
};
typedef int type_5[1];
typedef int type_7[1];
struct Bar {
metal::float4x4 matrix;
type_2 matrix_array;
metal::atomic_int atom;
char _pad2[4];
type_3 arr;
type_5 data;
char _pad3[4];
type_5 arr;
type_7 data;
};
struct type_11 {
struct type_13 {
int inner[5];
};

Expand All @@ -39,19 +43,19 @@ vertex fooOutput foo(
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
float foo_1 = 0.0;
type_11 c;
type_13 c;
float baz = foo_1;
foo_1 = 1.0;
metal::float4x4 matrix = bar.matrix;
type_3 arr = bar.arr;
type_5 arr = bar.arr;
float b = bar.matrix[3].x;
int a = bar.data[(1 + (_buffer_sizes.size0 - 88 - 4) / 8) - 2u];
int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 4) / 8) - 2u];
float _e25 = read_from_private(foo_1);
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] = type_3 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_5 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
bar.data[1] = 1;
for(int _i=0; _i<5; ++_i) c.inner[_i] = type_11 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
c.inner[vi + 1u] = 42;
int value = c.inner[vi];
return fooOutput { matrix * static_cast<metal::float4>(metal::int4(value)) };
Expand Down
Loading