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

test(naga): add struct-layout test

Browse Source
This commit is contained in:
Erich Gubler 2023-11-21 09:51:08 -07:00
parent 06b0a67551
commit 1cf0fe841a
13 changed files with 664 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

50
naga/tests/in/struct-layout.wgsl Normal file
View File

@ -0,0 +1,50 @@
// Create several type definitions to test `align` and `size` layout.
struct NoPadding {
@location(0)
v3: vec3f, // align 16, size 12; no start padding needed
@location(1)
f3: f32, // align 4, size 4; no start padding needed
}
@fragment
fn no_padding_frag(input: NoPadding) -> @location(0) vec4f {
_ = input;
return vec4f(0.0);
}
@vertex
fn no_padding_vert(input: NoPadding) -> @builtin(position) vec4f {
_ = input;
return vec4f(0.0);
}
@group(0) @binding(0) var<uniform> no_padding_uniform: NoPadding;
@group(0) @binding(1) var<storage, read_write> no_padding_storage: NoPadding;
@compute @workgroup_size(16,1,1)
fn no_padding_comp() {
var x: NoPadding;
x = no_padding_uniform;
x = no_padding_storage;
}
struct NeedsPadding {
@location(0) f3_forces_padding: f32, // align 4, size 4; no start padding needed
@location(1) v3_needs_padding: vec3f, // align 16, size 12; needs 12 bytes of padding
@location(2) f3: f32, // align 4, size 4; no start padding needed
}
@fragment
fn needs_padding_frag(input: NeedsPadding) -> @location(0) vec4f {
_ = input;
return vec4f(0.0);
}
@vertex
fn needs_padding_vert(input: NeedsPadding) -> @builtin(position) vec4f {
_ = input;
return vec4f(0.0);
}
@group(0) @binding(2) var<uniform> needs_padding_uniform: NeedsPadding;
@group(0) @binding(3) var<storage, read_write> needs_padding_storage: NeedsPadding;
@compute @workgroup_size(16,1,1)
fn needs_padding_comp() {
var x: NeedsPadding;
x = needs_padding_uniform;
x = needs_padding_storage;
}

30
naga/tests/out/glsl/struct-layout.needs_padding_comp.Compute.glsl Normal file
View File

@ -0,0 +1,30 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
struct NoPadding {
vec3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
vec3 v3_needs_padding;
float f3_;
};
uniform NeedsPadding_block_0Compute { NeedsPadding _group_0_binding_2_cs; };
layout(std430) buffer NeedsPadding_block_1Compute { NeedsPadding _group_0_binding_3_cs; };
void main() {
NeedsPadding x_1 = NeedsPadding(0.0, vec3(0.0), 0.0);
NeedsPadding _e2 = _group_0_binding_2_cs;
x_1 = _e2;
NeedsPadding _e4 = _group_0_binding_3_cs;
x_1 = _e4;
return;
}

25
naga/tests/out/glsl/struct-layout.needs_padding_frag.Fragment.glsl Normal file
View File

@ -0,0 +1,25 @@
#version 310 es
precision highp float;
precision highp int;
struct NoPadding {
vec3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
vec3 v3_needs_padding;
float f3_;
};
layout(location = 0) smooth in float _vs2fs_location0;
layout(location = 1) smooth in vec3 _vs2fs_location1;
layout(location = 2) smooth in float _vs2fs_location2;
layout(location = 0) out vec4 _fs2p_location0;
void main() {
NeedsPadding input_2 = NeedsPadding(_vs2fs_location0, _vs2fs_location1, _vs2fs_location2);
_fs2p_location0 = vec4(0.0);
return;
}

25
naga/tests/out/glsl/struct-layout.needs_padding_vert.Vertex.glsl Normal file
View File

@ -0,0 +1,25 @@
#version 310 es
precision highp float;
precision highp int;
struct NoPadding {
vec3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
vec3 v3_needs_padding;
float f3_;
};
layout(location = 0) in float _p2vs_location0;
layout(location = 1) in vec3 _p2vs_location1;
layout(location = 2) in float _p2vs_location2;
void main() {
NeedsPadding input_3 = NeedsPadding(_p2vs_location0, _p2vs_location1, _p2vs_location2);
gl_Position = vec4(0.0);
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;
}

30
naga/tests/out/glsl/struct-layout.no_padding_comp.Compute.glsl Normal file
View File

@ -0,0 +1,30 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
struct NoPadding {
vec3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
vec3 v3_needs_padding;
float f3_;
};
uniform NoPadding_block_0Compute { NoPadding _group_0_binding_0_cs; };
layout(std430) buffer NoPadding_block_1Compute { NoPadding _group_0_binding_1_cs; };
void main() {
NoPadding x = NoPadding(vec3(0.0), 0.0);
NoPadding _e2 = _group_0_binding_0_cs;
x = _e2;
NoPadding _e4 = _group_0_binding_1_cs;
x = _e4;
return;
}

24
naga/tests/out/glsl/struct-layout.no_padding_frag.Fragment.glsl Normal file
View File

@ -0,0 +1,24 @@
#version 310 es
precision highp float;
precision highp int;
struct NoPadding {
vec3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
vec3 v3_needs_padding;
float f3_;
};
layout(location = 0) smooth in vec3 _vs2fs_location0;
layout(location = 1) smooth in float _vs2fs_location1;
layout(location = 0) out vec4 _fs2p_location0;
void main() {
NoPadding input_ = NoPadding(_vs2fs_location0, _vs2fs_location1);
_fs2p_location0 = vec4(0.0);
return;
}

24
naga/tests/out/glsl/struct-layout.no_padding_vert.Vertex.glsl Normal file
View File

@ -0,0 +1,24 @@
#version 310 es
precision highp float;
precision highp int;
struct NoPadding {
vec3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
vec3 v3_needs_padding;
float f3_;
};
layout(location = 0) in vec3 _p2vs_location0;
layout(location = 1) in float _p2vs_location1;
void main() {
NoPadding input_1 = NoPadding(_p2vs_location0, _p2vs_location1);
gl_Position = vec4(0.0);
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;
}

87
naga/tests/out/hlsl/struct-layout.hlsl Normal file
View File

@ -0,0 +1,87 @@
struct NoPadding {
float3 v3_ : LOC0;
float f3_ : LOC1;
};
struct NeedsPadding {
float f3_forces_padding : LOC0;
float3 v3_needs_padding : LOC1;
float f3_ : LOC2;
};
cbuffer no_padding_uniform : register(b0) { NoPadding no_padding_uniform; }
RWByteAddressBuffer no_padding_storage : register(u1);
cbuffer needs_padding_uniform : register(b2) { NeedsPadding needs_padding_uniform; }
RWByteAddressBuffer needs_padding_storage : register(u3);
struct FragmentInput_no_padding_frag {
float3 v3_ : LOC0;
float f3_ : LOC1;
};
struct FragmentInput_needs_padding_frag {
float f3_forces_padding : LOC0;
float3 v3_needs_padding : LOC1;
float f3_1 : LOC2;
};
float4 no_padding_frag(FragmentInput_no_padding_frag fragmentinput_no_padding_frag) : SV_Target0
{
NoPadding input = { fragmentinput_no_padding_frag.v3_, fragmentinput_no_padding_frag.f3_ };
return (0.0).xxxx;
}
float4 no_padding_vert(NoPadding input_1) : SV_Position
{
return (0.0).xxxx;
}
NoPadding ConstructNoPadding(float3 arg0, float arg1) {
NoPadding ret = (NoPadding)0;
ret.v3_ = arg0;
ret.f3_ = arg1;
return ret;
}
[numthreads(16, 1, 1)]
void no_padding_comp()
{
NoPadding x = (NoPadding)0;
NoPadding _expr2 = no_padding_uniform;
x = _expr2;
NoPadding _expr4 = ConstructNoPadding(asfloat(no_padding_storage.Load3(0)), asfloat(no_padding_storage.Load(12)));
x = _expr4;
return;
}
float4 needs_padding_frag(FragmentInput_needs_padding_frag fragmentinput_needs_padding_frag) : SV_Target0
{
NeedsPadding input_2 = { fragmentinput_needs_padding_frag.f3_forces_padding, fragmentinput_needs_padding_frag.v3_needs_padding, fragmentinput_needs_padding_frag.f3_1 };
return (0.0).xxxx;
}
float4 needs_padding_vert(NeedsPadding input_3) : SV_Position
{
return (0.0).xxxx;
}
NeedsPadding ConstructNeedsPadding(float arg0, float3 arg1, float arg2) {
NeedsPadding ret = (NeedsPadding)0;
ret.f3_forces_padding = arg0;
ret.v3_needs_padding = arg1;
ret.f3_ = arg2;
return ret;
}
[numthreads(16, 1, 1)]
void needs_padding_comp()
{
NeedsPadding x_1 = (NeedsPadding)0;
NeedsPadding _expr2 = needs_padding_uniform;
x_1 = _expr2;
NeedsPadding _expr4 = ConstructNeedsPadding(asfloat(needs_padding_storage.Load(0)), asfloat(needs_padding_storage.Load3(16)), asfloat(needs_padding_storage.Load(28)));
x_1 = _expr4;
return;
}

32
naga/tests/out/hlsl/struct-layout.ron Normal file
View File

@ -0,0 +1,32 @@
(
vertex:[
(
entry_point:"no_padding_vert",
target_profile:"vs_5_1",
),
(
entry_point:"needs_padding_vert",
target_profile:"vs_5_1",
),
],
fragment:[
(
entry_point:"no_padding_frag",
target_profile:"ps_5_1",
),
(
entry_point:"needs_padding_frag",
target_profile:"ps_5_1",
),
],
compute:[
(
entry_point:"no_padding_comp",
target_profile:"cs_5_1",
),
(
entry_point:"needs_padding_comp",
target_profile:"cs_5_1",
),
],
)

103
naga/tests/out/msl/struct-layout.msl Normal file
View File

@ -0,0 +1,103 @@
// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
struct NoPadding {
metal::packed_float3 v3_;
float f3_;
};
struct NeedsPadding {
float f3_forces_padding;
char _pad1[12];
metal::packed_float3 v3_needs_padding;
float f3_;
};
struct no_padding_fragInput {
metal::float3 v3_ [[user(loc0), center_perspective]];
float f3_ [[user(loc1), center_perspective]];
};
struct no_padding_fragOutput {
metal::float4 member [[color(0)]];
};
fragment no_padding_fragOutput no_padding_frag(
no_padding_fragInput varyings [[stage_in]]
) {
const NoPadding input = { varyings.v3_, varyings.f3_ };
return no_padding_fragOutput { metal::float4(0.0) };
}
struct no_padding_vertInput {
metal::float3 v3_ [[attribute(0)]];
float f3_ [[attribute(1)]];
};
struct no_padding_vertOutput {
metal::float4 member_1 [[position]];
};
vertex no_padding_vertOutput no_padding_vert(
no_padding_vertInput varyings_1 [[stage_in]]
) {
const NoPadding input_1 = { varyings_1.v3_, varyings_1.f3_ };
return no_padding_vertOutput { metal::float4(0.0) };
}
kernel void no_padding_comp(
constant NoPadding& no_padding_uniform [[user(fake0)]]
, device NoPadding const& no_padding_storage [[user(fake0)]]
) {
NoPadding x = {};
NoPadding _e2 = no_padding_uniform;
x = _e2;
NoPadding _e4 = no_padding_storage;
x = _e4;
return;
}
struct needs_padding_fragInput {
float f3_forces_padding [[user(loc0), center_perspective]];
metal::float3 v3_needs_padding [[user(loc1), center_perspective]];
float f3_ [[user(loc2), center_perspective]];
};
struct needs_padding_fragOutput {
metal::float4 member_3 [[color(0)]];
};
fragment needs_padding_fragOutput needs_padding_frag(
needs_padding_fragInput varyings_3 [[stage_in]]
) {
const NeedsPadding input_2 = { varyings_3.f3_forces_padding, {}, varyings_3.v3_needs_padding, varyings_3.f3_ };
return needs_padding_fragOutput { metal::float4(0.0) };
}
struct needs_padding_vertInput {
float f3_forces_padding [[attribute(0)]];
metal::float3 v3_needs_padding [[attribute(1)]];
float f3_ [[attribute(2)]];
};
struct needs_padding_vertOutput {
metal::float4 member_4 [[position]];
};
vertex needs_padding_vertOutput needs_padding_vert(
needs_padding_vertInput varyings_4 [[stage_in]]
) {
const NeedsPadding input_3 = { varyings_4.f3_forces_padding, {}, varyings_4.v3_needs_padding, varyings_4.f3_ };
return needs_padding_vertOutput { metal::float4(0.0) };
}
kernel void needs_padding_comp(
constant NeedsPadding& needs_padding_uniform [[user(fake0)]]
, device NeedsPadding const& needs_padding_storage [[user(fake0)]]
) {
NeedsPadding x_1 = {};
NeedsPadding _e2 = needs_padding_uniform;
x_1 = _e2;
NeedsPadding _e4 = needs_padding_storage;
x_1 = _e4;
return;
}

169
naga/tests/out/spv/struct-layout.spvasm Normal file
View File

@ -0,0 +1,169 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 92
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %30 "no_padding_frag" %22 %25 %28
OpEntryPoint Vertex %42 "no_padding_vert" %37 %39 %41
OpEntryPoint GLCompute %45 "no_padding_comp"
OpEntryPoint Fragment %67 "needs_padding_frag" %60 %62 %64 %66
OpEntryPoint Vertex %78 "needs_padding_vert" %71 %73 %75 %77
OpEntryPoint GLCompute %81 "needs_padding_comp"
OpExecutionMode %30 OriginUpperLeft
OpExecutionMode %45 LocalSize 16 1 1
OpExecutionMode %67 OriginUpperLeft
OpExecutionMode %81 LocalSize 16 1 1
OpMemberDecorate %5 0 Offset 0
OpMemberDecorate %5 1 Offset 12
OpMemberDecorate %7 0 Offset 0
OpMemberDecorate %7 1 Offset 16
OpMemberDecorate %7 2 Offset 28
OpDecorate %8 DescriptorSet 0
OpDecorate %8 Binding 0
OpDecorate %9 Block
OpMemberDecorate %9 0 Offset 0
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 1
OpDecorate %12 Block
OpMemberDecorate %12 0 Offset 0
OpDecorate %14 DescriptorSet 0
OpDecorate %14 Binding 2
OpDecorate %15 Block
OpMemberDecorate %15 0 Offset 0
OpDecorate %17 DescriptorSet 0
OpDecorate %17 Binding 3
OpDecorate %18 Block
OpMemberDecorate %18 0 Offset 0
OpDecorate %22 Location 0
OpDecorate %25 Location 1
OpDecorate %28 Location 0
OpDecorate %37 Location 0
OpDecorate %39 Location 1
OpDecorate %41 BuiltIn Position
OpDecorate %60 Location 0
OpDecorate %62 Location 1
OpDecorate %64 Location 2
OpDecorate %66 Location 0
OpDecorate %71 Location 0
OpDecorate %73 Location 1
OpDecorate %75 Location 2
OpDecorate %77 BuiltIn Position
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpTypeVector %4 3
%5 = OpTypeStruct %3 %4
%6 = OpTypeVector %4 4
%7 = OpTypeStruct %4 %3 %4
%9 = OpTypeStruct %5
%10 = OpTypePointer Uniform %9
%8 = OpVariable %10 Uniform
%12 = OpTypeStruct %5
%13 = OpTypePointer StorageBuffer %12
%11 = OpVariable %13 StorageBuffer
%15 = OpTypeStruct %7
%16 = OpTypePointer Uniform %15
%14 = OpVariable %16 Uniform
%18 = OpTypeStruct %7
%19 = OpTypePointer StorageBuffer %18
%17 = OpVariable %19 StorageBuffer
%23 = OpTypePointer Input %3
%22 = OpVariable %23 Input
%26 = OpTypePointer Input %4
%25 = OpVariable %26 Input
%29 = OpTypePointer Output %6
%28 = OpVariable %29 Output
%31 = OpTypeFunction %2
%32 = OpConstant %4 0.0
%33 = OpConstantComposite %6 %32 %32 %32 %32
%37 = OpVariable %23 Input
%39 = OpVariable %26 Input
%41 = OpVariable %29 Output
%46 = OpTypePointer Uniform %5
%48 = OpTypeInt 32 0
%47 = OpConstant %48 0
%50 = OpTypePointer StorageBuffer %5
%53 = OpTypePointer Function %5
%54 = OpConstantNull %5
%60 = OpVariable %26 Input
%62 = OpVariable %23 Input
%64 = OpVariable %26 Input
%66 = OpVariable %29 Output
%71 = OpVariable %26 Input
%73 = OpVariable %23 Input
%75 = OpVariable %26 Input
%77 = OpVariable %29 Output
%82 = OpTypePointer Uniform %7
%84 = OpTypePointer StorageBuffer %7
%87 = OpTypePointer Function %7
%88 = OpConstantNull %7
%30 = OpFunction %2 None %31
%20 = OpLabel
%24 = OpLoad %3 %22
%27 = OpLoad %4 %25
%21 = OpCompositeConstruct %5 %24 %27
OpBranch %34
%34 = OpLabel
OpStore %28 %33
OpReturn
OpFunctionEnd
%42 = OpFunction %2 None %31
%35 = OpLabel
%38 = OpLoad %3 %37
%40 = OpLoad %4 %39
%36 = OpCompositeConstruct %5 %38 %40
OpBranch %43
%43 = OpLabel
OpStore %41 %33
OpReturn
OpFunctionEnd
%45 = OpFunction %2 None %31
%44 = OpLabel
%52 = OpVariable %53 Function %54
%49 = OpAccessChain %46 %8 %47
%51 = OpAccessChain %50 %11 %47
OpBranch %55
%55 = OpLabel
%56 = OpLoad %5 %49
OpStore %52 %56
%57 = OpLoad %5 %51
OpStore %52 %57
OpReturn
OpFunctionEnd
%67 = OpFunction %2 None %31
%58 = OpLabel
%61 = OpLoad %4 %60
%63 = OpLoad %3 %62
%65 = OpLoad %4 %64
%59 = OpCompositeConstruct %7 %61 %63 %65
OpBranch %68
%68 = OpLabel
OpStore %66 %33
OpReturn
OpFunctionEnd
%78 = OpFunction %2 None %31
%69 = OpLabel
%72 = OpLoad %4 %71
%74 = OpLoad %3 %73
%76 = OpLoad %4 %75
%70 = OpCompositeConstruct %7 %72 %74 %76
OpBranch %79
%79 = OpLabel
OpStore %77 %33
OpReturn
OpFunctionEnd
%81 = OpFunction %2 None %31
%80 = OpLabel
%86 = OpVariable %87 Function %88
%83 = OpAccessChain %82 %14 %47
%85 = OpAccessChain %84 %17 %47
OpBranch %89
%89 = OpLabel
%90 = OpLoad %7 %83
OpStore %86 %90
%91 = OpLoad %7 %85
OpStore %86 %91
OpReturn
OpFunctionEnd

61
naga/tests/out/wgsl/struct-layout.wgsl Normal file
View File

@ -0,0 +1,61 @@
struct NoPadding {
@location(0) v3_: vec3<f32>,
@location(1) f3_: f32,
}
struct NeedsPadding {
@location(0) f3_forces_padding: f32,
@location(1) v3_needs_padding: vec3<f32>,
@location(2) f3_: f32,
}
@group(0) @binding(0)
var<uniform> no_padding_uniform: NoPadding;
@group(0) @binding(1)
var<storage, read_write> no_padding_storage: NoPadding;
@group(0) @binding(2)
var<uniform> needs_padding_uniform: NeedsPadding;
@group(0) @binding(3)
var<storage, read_write> needs_padding_storage: NeedsPadding;
@fragment
fn no_padding_frag(input: NoPadding) -> @location(0) vec4<f32> {
return vec4(0.0);
}
@vertex
fn no_padding_vert(input_1: NoPadding) -> @builtin(position) vec4<f32> {
return vec4(0.0);
}
@compute @workgroup_size(16, 1, 1)
fn no_padding_comp() {
var x: NoPadding;
let _e2 = no_padding_uniform;
x = _e2;
let _e4 = no_padding_storage;
x = _e4;
return;
}
@fragment
fn needs_padding_frag(input_2: NeedsPadding) -> @location(0) vec4<f32> {
return vec4(0.0);
}
@vertex
fn needs_padding_vert(input_3: NeedsPadding) -> @builtin(position) vec4<f32> {
return vec4(0.0);
}
@compute @workgroup_size(16, 1, 1)
fn needs_padding_comp() {
var x_1: NeedsPadding;
let _e2 = needs_padding_uniform;
x_1 = _e2;
let _e4 = needs_padding_storage;
x_1 = _e4;
return;
}

4
naga/tests/snapshots.rs
View File

@ -775,6 +775,10 @@ fn convert_wgsl() {
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
),
("separate-entry-points", Targets::SPIRV | Targets::GLSL),
(
"struct-layout",
Targets::WGSL | Targets::GLSL | Targets::SPIRV | Targets::HLSL | Targets::METAL,
),
];
for &(name, targets) in inputs.iter() {