mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-26 08:44:08 +00:00
More tests for non-struct global types.
This commit is contained in:
Jim Blandy
parent
5d631d9e2d
commit
27d0fee2e7
@ -923,13 +923,17 @@ impl<'a, W: Write> Writer<'a, W> {
|
||||
crate::AddressSpace::PushConstant => {
|
||||
self.write_simple_global(handle, global)?;
|
||||
}
|
||||
crate::AddressSpace::Uniform | crate::AddressSpace::Handle => {
|
||||
crate::AddressSpace::Uniform => {
|
||||
self.write_interface_block(handle, global)?;
|
||||
}
|
||||
crate::AddressSpace::Storage { .. } => {
|
||||
self.write_interface_block(handle, global)?;
|
||||
}
|
||||
// A global variable in the `Function` address space is a
|
||||
// contradiction in terms.
|
||||
crate::AddressSpace::Function => unreachable!(),
|
||||
// Textures and samplers are handled directly in `Writer::write`.
|
||||
crate::AddressSpace::Handle => unreachable!(),
|
||||
}
|
||||
|
||||
Ok(())
|
||||
|
||||
@ -19,6 +19,9 @@ var<storage> dummy: array<vec2<f32>>;
|
||||
@group(0) @binding(3)
|
||||
var<uniform> float_vecs: array<vec4<f32>, 20>;
|
||||
|
||||
@group(0) @binding(4)
|
||||
var<uniform> global_vec: vec4<f32>;
|
||||
|
||||
fn test_msl_packed_vec3_as_arg(arg: vec3<f32>) {}
|
||||
|
||||
fn test_msl_packed_vec3() {
|
||||
@ -50,6 +53,9 @@ fn test_msl_packed_vec3() {
|
||||
fn main() {
|
||||
test_msl_packed_vec3();
|
||||
|
||||
wg[6] = global_vec.x;
|
||||
wg[5] = dummy[1].y;
|
||||
wg[4] = float_vecs[0].w;
|
||||
wg[3] = alignment.v1;
|
||||
wg[2] = alignment.v3.x;
|
||||
alignment.v1 = 4.0;
|
||||
|
||||
@ -17,6 +17,10 @@ layout(std430) buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; };
|
||||
|
||||
layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; };
|
||||
|
||||
uniform type_8_block_2Compute { vec4 _group_0_binding_3_cs[20]; };
|
||||
|
||||
uniform type_7_block_3Compute { vec4 _group_0_binding_4_cs; };
|
||||
|
||||
|
||||
void test_msl_packed_vec3_as_arg(vec3 arg) {
|
||||
return;
|
||||
@ -27,8 +31,8 @@ void test_msl_packed_vec3_() {
|
||||
_group_0_binding_1_cs.v3_ = vec3(1.0);
|
||||
_group_0_binding_1_cs.v3_.x = 1.0;
|
||||
_group_0_binding_1_cs.v3_.x = 2.0;
|
||||
int _e19 = idx;
|
||||
_group_0_binding_1_cs.v3_[_e19] = 3.0;
|
||||
int _e20 = idx;
|
||||
_group_0_binding_1_cs.v3_[_e20] = 3.0;
|
||||
Foo data = _group_0_binding_1_cs;
|
||||
vec3 unnamed = data.v3_;
|
||||
vec2 unnamed_1 = data.v3_.zx;
|
||||
@ -43,10 +47,16 @@ void main() {
|
||||
float Foo_1 = 1.0;
|
||||
bool at = true;
|
||||
test_msl_packed_vec3_();
|
||||
float _e9 = _group_0_binding_1_cs.v1_;
|
||||
wg[3] = _e9;
|
||||
float _e14 = _group_0_binding_1_cs.v3_.x;
|
||||
wg[2] = _e14;
|
||||
float _e10 = _group_0_binding_4_cs.x;
|
||||
wg[6] = _e10;
|
||||
float _e16 = _group_0_binding_2_cs[1].y;
|
||||
wg[5] = _e16;
|
||||
float _e22 = _group_0_binding_3_cs[0].w;
|
||||
wg[4] = _e22;
|
||||
float _e26 = _group_0_binding_1_cs.v1_;
|
||||
wg[3] = _e26;
|
||||
float _e31 = _group_0_binding_1_cs.v3_.x;
|
||||
wg[2] = _e31;
|
||||
_group_0_binding_1_cs.v1_ = 4.0;
|
||||
wg[1] = float(uint(_group_0_binding_2_cs.length()));
|
||||
at_1 = 2u;
|
||||
|
||||
@ -10,6 +10,7 @@ groupshared uint at_1;
|
||||
RWByteAddressBuffer alignment : register(u1);
|
||||
ByteAddressBuffer dummy : register(t2);
|
||||
cbuffer float_vecs : register(b3) { float4 float_vecs[20]; }
|
||||
cbuffer global_vec : register(b4) { float4 global_vec; }
|
||||
|
||||
void test_msl_packed_vec3_as_arg(float3 arg)
|
||||
{
|
||||
@ -23,8 +24,8 @@ void test_msl_packed_vec3_()
|
||||
alignment.Store3(0, asuint(float3(1.0.xxx)));
|
||||
alignment.Store(0+0, asuint(1.0));
|
||||
alignment.Store(0+0, asuint(2.0));
|
||||
int _expr19 = idx;
|
||||
alignment.Store(_expr19*4+0, asuint(3.0));
|
||||
int _expr20 = idx;
|
||||
alignment.Store(_expr20*4+0, asuint(3.0));
|
||||
Foo data = {asfloat(alignment.Load3(0)), asfloat(alignment.Load(12))};
|
||||
float3 unnamed = data.v3_;
|
||||
float2 unnamed_1 = data.v3_.zx;
|
||||
@ -49,10 +50,16 @@ void main()
|
||||
bool at = true;
|
||||
|
||||
test_msl_packed_vec3_();
|
||||
float _expr9 = asfloat(alignment.Load(12));
|
||||
wg[3] = _expr9;
|
||||
float _expr14 = asfloat(alignment.Load(0+0));
|
||||
wg[2] = _expr14;
|
||||
float _expr10 = global_vec.x;
|
||||
wg[6] = _expr10;
|
||||
float _expr16 = asfloat(dummy.Load(4+8));
|
||||
wg[5] = _expr16;
|
||||
float _expr22 = float_vecs[0].w;
|
||||
wg[4] = _expr22;
|
||||
float _expr26 = asfloat(alignment.Load(12));
|
||||
wg[3] = _expr26;
|
||||
float _expr31 = asfloat(alignment.Load(0+0));
|
||||
wg[2] = _expr31;
|
||||
alignment.Store(12, asuint(4.0));
|
||||
wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
|
||||
at_1 = 2u;
|
||||
|
||||
@ -36,8 +36,8 @@ void test_msl_packed_vec3_(
|
||||
alignment.v3_ = metal::float3(1.0);
|
||||
alignment.v3_[0] = 1.0;
|
||||
alignment.v3_[0] = 2.0;
|
||||
int _e19 = idx;
|
||||
alignment.v3_[_e19] = 3.0;
|
||||
int _e20 = idx;
|
||||
alignment.v3_[_e20] = 3.0;
|
||||
Foo data = alignment;
|
||||
metal::float3 unnamed = data.v3_;
|
||||
metal::float2 unnamed_1 = metal::float3(data.v3_).zx;
|
||||
@ -53,15 +53,23 @@ kernel void main_(
|
||||
, threadgroup metal::atomic_uint& at_1
|
||||
, device Foo& alignment [[user(fake0)]]
|
||||
, device type_6 const& dummy [[user(fake0)]]
|
||||
, constant type_8& float_vecs [[user(fake0)]]
|
||||
, constant metal::float4& global_vec [[user(fake0)]]
|
||||
, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
|
||||
) {
|
||||
float Foo_1 = 1.0;
|
||||
bool at = true;
|
||||
test_msl_packed_vec3_(alignment);
|
||||
float _e9 = alignment.v1_;
|
||||
wg.inner[3] = _e9;
|
||||
float _e14 = alignment.v3_[0];
|
||||
wg.inner[2] = _e14;
|
||||
float _e10 = global_vec.x;
|
||||
wg.inner[6] = _e10;
|
||||
float _e16 = dummy[1].y;
|
||||
wg.inner[5] = _e16;
|
||||
float _e22 = float_vecs.inner[0].w;
|
||||
wg.inner[4] = _e22;
|
||||
float _e26 = alignment.v1_;
|
||||
wg.inner[3] = _e26;
|
||||
float _e31 = alignment.v3_[0];
|
||||
wg.inner[2] = _e31;
|
||||
alignment.v1_ = 4.0;
|
||||
wg.inner[1] = static_cast<float>(1 + (_buffer_sizes.size3 - 0 - 8) / 8);
|
||||
metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);
|
||||
|
||||
@ -1,31 +1,35 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 107
|
||||
; Bound: 130
|
||||
OpCapability Shader
|
||||
OpExtension "SPV_KHR_storage_buffer_storage_class"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %87 "main"
|
||||
OpExecutionMode %87 LocalSize 1 1 1
|
||||
OpDecorate %21 ArrayStride 4
|
||||
OpMemberDecorate %23 0 Offset 0
|
||||
OpMemberDecorate %23 1 Offset 12
|
||||
OpDecorate %25 ArrayStride 8
|
||||
OpDecorate %27 ArrayStride 16
|
||||
OpDecorate %35 DescriptorSet 0
|
||||
OpDecorate %35 Binding 1
|
||||
OpDecorate %36 Block
|
||||
OpMemberDecorate %36 0 Offset 0
|
||||
OpDecorate %38 NonWritable
|
||||
OpEntryPoint GLCompute %94 "main"
|
||||
OpExecutionMode %94 LocalSize 1 1 1
|
||||
OpDecorate %24 ArrayStride 4
|
||||
OpMemberDecorate %26 0 Offset 0
|
||||
OpMemberDecorate %26 1 Offset 12
|
||||
OpDecorate %28 ArrayStride 8
|
||||
OpDecorate %30 ArrayStride 16
|
||||
OpDecorate %38 DescriptorSet 0
|
||||
OpDecorate %38 Binding 2
|
||||
OpDecorate %38 Binding 1
|
||||
OpDecorate %39 Block
|
||||
OpMemberDecorate %39 0 Offset 0
|
||||
OpDecorate %41 NonWritable
|
||||
OpDecorate %41 DescriptorSet 0
|
||||
OpDecorate %41 Binding 3
|
||||
OpDecorate %41 Binding 2
|
||||
OpDecorate %42 Block
|
||||
OpMemberDecorate %42 0 Offset 0
|
||||
OpDecorate %44 DescriptorSet 0
|
||||
OpDecorate %44 Binding 3
|
||||
OpDecorate %45 Block
|
||||
OpMemberDecorate %45 0 Offset 0
|
||||
OpDecorate %47 DescriptorSet 0
|
||||
OpDecorate %47 Binding 4
|
||||
OpDecorate %48 Block
|
||||
OpMemberDecorate %48 0 Offset 0
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeBool
|
||||
%3 = OpConstantTrue %4
|
||||
@ -40,112 +44,138 @@ OpMemberDecorate %42 0 Offset 0
|
||||
%13 = OpConstant %10 2.0
|
||||
%14 = OpConstant %10 3.0
|
||||
%15 = OpConstant %10 0.0
|
||||
%16 = OpConstant %8 3
|
||||
%17 = OpConstant %8 2
|
||||
%18 = OpConstant %10 4.0
|
||||
%19 = OpConstant %6 2
|
||||
%20 = OpConstantTrue %4
|
||||
%21 = OpTypeArray %10 %5
|
||||
%22 = OpTypeVector %10 3
|
||||
%23 = OpTypeStruct %22 %10
|
||||
%24 = OpTypeVector %10 2
|
||||
%25 = OpTypeRuntimeArray %24
|
||||
%26 = OpTypeVector %10 4
|
||||
%27 = OpTypeArray %26 %7
|
||||
%28 = OpTypeMatrix %22 3
|
||||
%29 = OpConstantComposite %22 %15 %15 %15
|
||||
%30 = OpConstantComposite %28 %29 %29 %29
|
||||
%32 = OpTypePointer Workgroup %21
|
||||
%31 = OpVariable %32 Workgroup
|
||||
%34 = OpTypePointer Workgroup %6
|
||||
%33 = OpVariable %34 Workgroup
|
||||
%36 = OpTypeStruct %23
|
||||
%37 = OpTypePointer StorageBuffer %36
|
||||
%35 = OpVariable %37 StorageBuffer
|
||||
%39 = OpTypeStruct %25
|
||||
%16 = OpConstant %8 6
|
||||
%17 = OpConstant %8 5
|
||||
%18 = OpConstant %8 4
|
||||
%19 = OpConstant %8 3
|
||||
%20 = OpConstant %8 2
|
||||
%21 = OpConstant %10 4.0
|
||||
%22 = OpConstant %6 2
|
||||
%23 = OpConstantTrue %4
|
||||
%24 = OpTypeArray %10 %5
|
||||
%25 = OpTypeVector %10 3
|
||||
%26 = OpTypeStruct %25 %10
|
||||
%27 = OpTypeVector %10 2
|
||||
%28 = OpTypeRuntimeArray %27
|
||||
%29 = OpTypeVector %10 4
|
||||
%30 = OpTypeArray %29 %7
|
||||
%31 = OpTypeMatrix %25 3
|
||||
%32 = OpConstantComposite %25 %15 %15 %15
|
||||
%33 = OpConstantComposite %31 %32 %32 %32
|
||||
%35 = OpTypePointer Workgroup %24
|
||||
%34 = OpVariable %35 Workgroup
|
||||
%37 = OpTypePointer Workgroup %6
|
||||
%36 = OpVariable %37 Workgroup
|
||||
%39 = OpTypeStruct %26
|
||||
%40 = OpTypePointer StorageBuffer %39
|
||||
%38 = OpVariable %40 StorageBuffer
|
||||
%42 = OpTypeStruct %27
|
||||
%43 = OpTypePointer Uniform %42
|
||||
%41 = OpVariable %43 Uniform
|
||||
%47 = OpTypeFunction %2 %22
|
||||
%48 = OpTypePointer StorageBuffer %23
|
||||
%49 = OpTypePointer StorageBuffer %25
|
||||
%50 = OpTypePointer Uniform %27
|
||||
%53 = OpTypePointer Function %8
|
||||
%56 = OpTypeFunction %2
|
||||
%57 = OpConstant %6 0
|
||||
%60 = OpTypePointer StorageBuffer %22
|
||||
%63 = OpTypePointer StorageBuffer %10
|
||||
%83 = OpTypePointer Function %10
|
||||
%85 = OpTypePointer Function %4
|
||||
%92 = OpTypePointer Workgroup %10
|
||||
%93 = OpTypePointer StorageBuffer %10
|
||||
%94 = OpConstant %6 1
|
||||
%97 = OpConstant %6 3
|
||||
%106 = OpConstant %6 256
|
||||
%46 = OpFunction %2 None %47
|
||||
%45 = OpFunctionParameter %22
|
||||
%44 = OpLabel
|
||||
OpBranch %51
|
||||
%51 = OpLabel
|
||||
%42 = OpTypeStruct %28
|
||||
%43 = OpTypePointer StorageBuffer %42
|
||||
%41 = OpVariable %43 StorageBuffer
|
||||
%45 = OpTypeStruct %30
|
||||
%46 = OpTypePointer Uniform %45
|
||||
%44 = OpVariable %46 Uniform
|
||||
%48 = OpTypeStruct %29
|
||||
%49 = OpTypePointer Uniform %48
|
||||
%47 = OpVariable %49 Uniform
|
||||
%53 = OpTypeFunction %2 %25
|
||||
%54 = OpTypePointer StorageBuffer %26
|
||||
%55 = OpTypePointer StorageBuffer %28
|
||||
%56 = OpTypePointer Uniform %29
|
||||
%57 = OpTypePointer Uniform %30
|
||||
%60 = OpTypePointer Function %8
|
||||
%63 = OpTypeFunction %2
|
||||
%64 = OpConstant %6 0
|
||||
%67 = OpTypePointer StorageBuffer %25
|
||||
%70 = OpTypePointer StorageBuffer %10
|
||||
%90 = OpTypePointer Function %10
|
||||
%92 = OpTypePointer Function %4
|
||||
%101 = OpTypePointer Workgroup %10
|
||||
%102 = OpTypePointer Uniform %10
|
||||
%105 = OpConstant %6 6
|
||||
%107 = OpTypePointer StorageBuffer %27
|
||||
%108 = OpConstant %6 1
|
||||
%111 = OpConstant %6 5
|
||||
%113 = OpConstant %6 3
|
||||
%116 = OpConstant %6 4
|
||||
%118 = OpTypePointer StorageBuffer %10
|
||||
%129 = OpConstant %6 256
|
||||
%52 = OpFunction %2 None %53
|
||||
%51 = OpFunctionParameter %25
|
||||
%50 = OpLabel
|
||||
OpBranch %58
|
||||
%58 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%55 = OpFunction %2 None %56
|
||||
%54 = OpLabel
|
||||
%52 = OpVariable %53 Function %11
|
||||
%58 = OpAccessChain %48 %35 %57
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
%61 = OpCompositeConstruct %22 %9 %9 %9
|
||||
%62 = OpAccessChain %60 %58 %57
|
||||
OpStore %62 %61
|
||||
%64 = OpAccessChain %63 %58 %57 %57
|
||||
OpStore %64 %9
|
||||
%65 = OpAccessChain %63 %58 %57 %57
|
||||
OpStore %65 %13
|
||||
%66 = OpLoad %8 %52
|
||||
%67 = OpAccessChain %63 %58 %57 %66
|
||||
OpStore %67 %14
|
||||
%68 = OpLoad %23 %58
|
||||
%69 = OpCompositeExtract %22 %68 0
|
||||
%70 = OpCompositeExtract %22 %68 0
|
||||
%71 = OpVectorShuffle %24 %70 %70 2 0
|
||||
%72 = OpCompositeExtract %22 %68 0
|
||||
%73 = OpFunctionCall %2 %46 %72
|
||||
%74 = OpCompositeExtract %22 %68 0
|
||||
%75 = OpVectorTimesMatrix %22 %74 %30
|
||||
%76 = OpCompositeExtract %22 %68 0
|
||||
%77 = OpMatrixTimesVector %22 %30 %76
|
||||
%78 = OpCompositeExtract %22 %68 0
|
||||
%79 = OpVectorTimesScalar %22 %78 %13
|
||||
%80 = OpCompositeExtract %22 %68 0
|
||||
%81 = OpVectorTimesScalar %22 %80 %13
|
||||
%62 = OpFunction %2 None %63
|
||||
%61 = OpLabel
|
||||
%59 = OpVariable %60 Function %11
|
||||
%65 = OpAccessChain %54 %38 %64
|
||||
OpBranch %66
|
||||
%66 = OpLabel
|
||||
%68 = OpCompositeConstruct %25 %9 %9 %9
|
||||
%69 = OpAccessChain %67 %65 %64
|
||||
OpStore %69 %68
|
||||
%71 = OpAccessChain %70 %65 %64 %64
|
||||
OpStore %71 %9
|
||||
%72 = OpAccessChain %70 %65 %64 %64
|
||||
OpStore %72 %13
|
||||
%73 = OpLoad %8 %59
|
||||
%74 = OpAccessChain %70 %65 %64 %73
|
||||
OpStore %74 %14
|
||||
%75 = OpLoad %26 %65
|
||||
%76 = OpCompositeExtract %25 %75 0
|
||||
%77 = OpCompositeExtract %25 %75 0
|
||||
%78 = OpVectorShuffle %27 %77 %77 2 0
|
||||
%79 = OpCompositeExtract %25 %75 0
|
||||
%80 = OpFunctionCall %2 %52 %79
|
||||
%81 = OpCompositeExtract %25 %75 0
|
||||
%82 = OpVectorTimesMatrix %25 %81 %33
|
||||
%83 = OpCompositeExtract %25 %75 0
|
||||
%84 = OpMatrixTimesVector %25 %33 %83
|
||||
%85 = OpCompositeExtract %25 %75 0
|
||||
%86 = OpVectorTimesScalar %25 %85 %13
|
||||
%87 = OpCompositeExtract %25 %75 0
|
||||
%88 = OpVectorTimesScalar %25 %87 %13
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%87 = OpFunction %2 None %56
|
||||
%86 = OpLabel
|
||||
%82 = OpVariable %83 Function %9
|
||||
%84 = OpVariable %85 Function %20
|
||||
%88 = OpAccessChain %48 %35 %57
|
||||
%89 = OpAccessChain %49 %38 %57
|
||||
OpBranch %90
|
||||
%90 = OpLabel
|
||||
%91 = OpFunctionCall %2 %55
|
||||
%95 = OpAccessChain %93 %88 %94
|
||||
%96 = OpLoad %10 %95
|
||||
%98 = OpAccessChain %92 %31 %97
|
||||
OpStore %98 %96
|
||||
%99 = OpAccessChain %63 %88 %57 %57
|
||||
%100 = OpLoad %10 %99
|
||||
%101 = OpAccessChain %92 %31 %19
|
||||
OpStore %101 %100
|
||||
%102 = OpAccessChain %93 %88 %94
|
||||
OpStore %102 %18
|
||||
%103 = OpArrayLength %6 %38 0
|
||||
%104 = OpConvertUToF %10 %103
|
||||
%105 = OpAccessChain %92 %31 %94
|
||||
OpStore %105 %104
|
||||
OpAtomicStore %33 %17 %106 %19
|
||||
%94 = OpFunction %2 None %63
|
||||
%93 = OpLabel
|
||||
%89 = OpVariable %90 Function %9
|
||||
%91 = OpVariable %92 Function %23
|
||||
%95 = OpAccessChain %54 %38 %64
|
||||
%96 = OpAccessChain %55 %41 %64
|
||||
%97 = OpAccessChain %57 %44 %64
|
||||
%98 = OpAccessChain %56 %47 %64
|
||||
OpBranch %99
|
||||
%99 = OpLabel
|
||||
%100 = OpFunctionCall %2 %62
|
||||
%103 = OpAccessChain %102 %98 %64
|
||||
%104 = OpLoad %10 %103
|
||||
%106 = OpAccessChain %101 %34 %105
|
||||
OpStore %106 %104
|
||||
%109 = OpAccessChain %70 %96 %108 %108
|
||||
%110 = OpLoad %10 %109
|
||||
%112 = OpAccessChain %101 %34 %111
|
||||
OpStore %112 %110
|
||||
%114 = OpAccessChain %102 %97 %64 %113
|
||||
%115 = OpLoad %10 %114
|
||||
%117 = OpAccessChain %101 %34 %116
|
||||
OpStore %117 %115
|
||||
%119 = OpAccessChain %118 %95 %108
|
||||
%120 = OpLoad %10 %119
|
||||
%121 = OpAccessChain %101 %34 %113
|
||||
OpStore %121 %120
|
||||
%122 = OpAccessChain %70 %95 %64 %64
|
||||
%123 = OpLoad %10 %122
|
||||
%124 = OpAccessChain %101 %34 %22
|
||||
OpStore %124 %123
|
||||
%125 = OpAccessChain %118 %95 %108
|
||||
OpStore %125 %21
|
||||
%126 = OpArrayLength %6 %41 0
|
||||
%127 = OpConvertUToF %10 %126
|
||||
%128 = OpAccessChain %101 %34 %108
|
||||
OpStore %128 %127
|
||||
OpAtomicStore %36 %20 %129 %22
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -13,6 +13,8 @@ var<storage, read_write> alignment: Foo;
|
||||
var<storage> dummy: array<vec2<f32>>;
|
||||
@group(0) @binding(3)
|
||||
var<uniform> float_vecs: array<vec4<f32>,20>;
|
||||
@group(0) @binding(4)
|
||||
var<uniform> global_vec: vec4<f32>;
|
||||
|
||||
fn test_msl_packed_vec3_as_arg(arg: vec3<f32>) {
|
||||
return;
|
||||
@ -24,8 +26,8 @@ fn test_msl_packed_vec3_() {
|
||||
alignment.v3_ = vec3<f32>(1.0);
|
||||
alignment.v3_.x = 1.0;
|
||||
alignment.v3_.x = 2.0;
|
||||
let _e19 = idx;
|
||||
alignment.v3_[_e19] = 3.0;
|
||||
let _e20 = idx;
|
||||
alignment.v3_[_e20] = 3.0;
|
||||
let data = alignment;
|
||||
let unnamed = data.v3_;
|
||||
let unnamed_1 = data.v3_.zx;
|
||||
@ -42,10 +44,16 @@ fn main() {
|
||||
var at: bool = true;
|
||||
|
||||
test_msl_packed_vec3_();
|
||||
let _e9 = alignment.v1_;
|
||||
wg[3] = _e9;
|
||||
let _e14 = alignment.v3_.x;
|
||||
wg[2] = _e14;
|
||||
let _e10 = global_vec.x;
|
||||
wg[6] = _e10;
|
||||
let _e16 = dummy[1].y;
|
||||
wg[5] = _e16;
|
||||
let _e22 = float_vecs[0].w;
|
||||
wg[4] = _e22;
|
||||
let _e26 = alignment.v1_;
|
||||
wg[3] = _e26;
|
||||
let _e31 = alignment.v3_.x;
|
||||
wg[2] = _e31;
|
||||
alignment.v1_ = 4.0;
|
||||
wg[1] = f32(arrayLength((&dummy)));
|
||||
atomicStore((&at_1), 2u);
|
||||
|
||||
Loading…
Reference in New Issue
Block a user