From f8c7cb127d1a31be82cd298743cc1a7e7153b6d4 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Mon, 3 Apr 2023 12:25:18 -0400 Subject: [PATCH] test: remove atomic ops from `access` --- tests/in/access.wgsl | 17 - tests/out/analysis/access.info.ron | 627 ----------------------------- tests/out/hlsl/access.hlsl | 26 -- tests/out/hlsl/access.hlsl.config | 2 +- tests/out/ir/access.ron | 248 ------------ tests/out/msl/access.msl | 27 -- tests/out/spv/access.spvasm | 106 ++--- tests/out/wgsl/access.wgsl | 25 -- 8 files changed, 33 insertions(+), 1045 deletions(-) diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index a9c27ee78c..5bfd4216bb 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -151,23 +151,6 @@ fn foo_frag() -> @location(0) vec4 { return vec4(0.0); } -@compute @workgroup_size(1) -fn atomics() { - var tmp: i32; - let value = atomicLoad(&bar.atom); - tmp = atomicAdd(&bar.atom, 5); - tmp = atomicSub(&bar.atom, 5); - tmp = atomicAnd(&bar.atom, 5); - tmp = atomicOr(&bar.atom, 5); - tmp = atomicXor(&bar.atom, 5); - tmp = atomicMin(&bar.atom, 5); - tmp = atomicMax(&bar.atom, 5); - tmp = atomicExchange(&bar.atom, 5); - // https://github.com/gpuweb/gpuweb/issues/2021 - // tmp = atomicCompareExchangeWeak(&bar.atom, 5, 5); - atomicStore(&bar.atom, value); -} - var val: u32; fn assign_through_ptr_fn(p: ptr) { diff --git a/tests/out/analysis/access.info.ron b/tests/out/analysis/access.info.ron index 01e08a63ba..8c5097f65a 100644 --- a/tests/out/analysis/access.info.ron +++ b/tests/out/analysis/access.info.ron @@ -5123,633 +5123,6 @@ ], sampling: [], ), - ( - flags: ( - bits: 63, - ), - available_stages: ( - bits: 7, - ), - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - may_kill: false, - sampling_set: [], - global_uses: [ - ( - bits: 0, - ), - ( - bits: 3, - ), - ( - bits: 0, - ), - ( - bits: 0, - ), - ( - bits: 0, - ), - ( - bits: 0, - ), - ], - expressions: [ - ( - uniformity: ( - non_uniform_result: Some(1), - requirements: ( - bits: 0, - ), - ), - ref_count: 8, - assignable_global: None, - ty: Value(Pointer( - base: 3, - space: Function, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(2), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(2), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(2), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(5), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(5), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(8), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(9), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(9), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(12), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(13), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(13), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(16), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(17), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(17), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(20), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(21), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(21), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(24), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(25), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(25), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(28), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(29), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(29), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(32), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(33), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(33), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(36), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(37), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(37), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ], - sampling: [], - ), ( flags: ( bits: 63, diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index f20cb726db..27786717d4 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -311,32 +311,6 @@ float4 foo_frag() : SV_Target0 return (0.0).xxxx; } -[numthreads(1, 1, 1)] -void atomics() -{ - int tmp = (int)0; - - int value_1 = asint(bar.Load(96)); - int _e7; bar.InterlockedAdd(96, 5, _e7); - tmp = _e7; - int _e11; bar.InterlockedAdd(96, -5, _e11); - tmp = _e11; - int _e15; bar.InterlockedAnd(96, 5, _e15); - tmp = _e15; - int _e19; bar.InterlockedOr(96, 5, _e19); - tmp = _e19; - int _e23; bar.InterlockedXor(96, 5, _e23); - tmp = _e23; - int _e27; bar.InterlockedMin(96, 5, _e27); - tmp = _e27; - int _e31; bar.InterlockedMax(96, 5, _e31); - tmp = _e31; - int _e35; bar.InterlockedExchange(96, 5, _e35); - tmp = _e35; - bar.Store(96, asuint(value_1)); - return; -} - [numthreads(1, 1, 1)] void assign_through_ptr(uint3 __local_invocation_id : SV_GroupThreadID) { diff --git a/tests/out/hlsl/access.hlsl.config b/tests/out/hlsl/access.hlsl.config index b68dda72d3..de7719bdff 100644 --- a/tests/out/hlsl/access.hlsl.config +++ b/tests/out/hlsl/access.hlsl.config @@ -1,3 +1,3 @@ vertex=(foo_vert:vs_5_1 ) fragment=(foo_frag:ps_5_1 ) -compute=(atomics:cs_5_1 assign_through_ptr:cs_5_1 ) +compute=(assign_through_ptr:cs_5_1 ) diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 41772b9332..0c1f9df1cb 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -2535,254 +2535,6 @@ ], ), ), - ( - name: "atomics", - stage: Compute, - early_depth_test: None, - workgroup_size: (1, 1, 1), - function: ( - name: Some("atomics"), - arguments: [], - result: None, - local_variables: [ - ( - name: Some("tmp"), - ty: 3, - init: None, - ), - ], - expressions: [ - LocalVariable(1), - GlobalVariable(2), - AccessIndex( - base: 2, - index: 2, - ), - Load( - pointer: 3, - ), - GlobalVariable(2), - AccessIndex( - base: 5, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 9, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 13, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 17, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 21, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 25, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 29, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 33, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 37, - index: 2, - ), - ], - named_expressions: { - 4: "value", - }, - body: [ - Emit(( - start: 2, - end: 4, - )), - Emit(( - start: 5, - end: 6, - )), - Atomic( - pointer: 6, - fun: Add, - value: 7, - result: 8, - ), - Store( - pointer: 1, - value: 8, - ), - Emit(( - start: 9, - end: 10, - )), - Atomic( - pointer: 10, - fun: Subtract, - value: 11, - result: 12, - ), - Store( - pointer: 1, - value: 12, - ), - Emit(( - start: 13, - end: 14, - )), - Atomic( - pointer: 14, - fun: And, - value: 15, - result: 16, - ), - Store( - pointer: 1, - value: 16, - ), - Emit(( - start: 17, - end: 18, - )), - Atomic( - pointer: 18, - fun: InclusiveOr, - value: 19, - result: 20, - ), - Store( - pointer: 1, - value: 20, - ), - Emit(( - start: 21, - end: 22, - )), - Atomic( - pointer: 22, - fun: ExclusiveOr, - value: 23, - result: 24, - ), - Store( - pointer: 1, - value: 24, - ), - Emit(( - start: 25, - end: 26, - )), - Atomic( - pointer: 26, - fun: Min, - value: 27, - result: 28, - ), - Store( - pointer: 1, - value: 28, - ), - Emit(( - start: 29, - end: 30, - )), - Atomic( - pointer: 30, - fun: Max, - value: 31, - result: 32, - ), - Store( - pointer: 1, - value: 32, - ), - Emit(( - start: 33, - end: 34, - )), - Atomic( - pointer: 34, - fun: Exchange( - compare: None, - ), - value: 35, - result: 36, - ), - Store( - pointer: 1, - value: 36, - ), - Emit(( - start: 37, - end: 38, - )), - Store( - pointer: 38, - value: 4, - ), - Return( - value: None, - ), - ], - ), - ), ( name: "assign_through_ptr", stage: Compute, diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index 8ad40973f4..a4967117ff 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -222,33 +222,6 @@ fragment foo_fragOutput foo_frag( } -kernel void atomics( - device Bar& bar [[buffer(0)]] -, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] -) { - int tmp = {}; - int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed); - int _e7 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e7; - int _e11 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e11; - int _e15 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e15; - int _e19 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e19; - int _e23 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e23; - int _e27 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e27; - int _e31 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e31; - int _e35 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e35; - metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed); - return; -} - - kernel void assign_through_ptr( metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] , threadgroup uint& val diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index a882322922..fa828b4d88 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,18 +1,16 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 354 +; Bound: 328 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %247 "foo_vert" %242 %245 OpEntryPoint Fragment %289 "foo_frag" %288 -OpEntryPoint GLCompute %308 "atomics" -OpEntryPoint GLCompute %334 "assign_through_ptr" %337 +OpEntryPoint GLCompute %308 "assign_through_ptr" %311 OpExecutionMode %289 OriginUpperLeft OpExecutionMode %308 LocalSize 1 1 1 -OpExecutionMode %334 LocalSize 1 1 1 OpSource GLSL 450 OpMemberName %36 0 "a" OpMemberName %36 1 "b" @@ -56,10 +54,8 @@ OpName %238 "c2" OpName %242 "vi" OpName %247 "foo_vert" OpName %289 "foo_frag" -OpName %305 "tmp" -OpName %308 "atomics" -OpName %331 "arr" -OpName %334 "assign_through_ptr" +OpName %305 "arr" +OpName %308 "assign_through_ptr" OpMemberDecorate %36 0 Offset 0 OpMemberDecorate %36 1 Offset 16 OpMemberDecorate %36 2 Offset 28 @@ -107,7 +103,7 @@ OpMemberDecorate %83 0 Offset 0 OpDecorate %242 BuiltIn VertexIndex OpDecorate %245 BuiltIn Position OpDecorate %288 Location 0 -OpDecorate %337 BuiltIn LocalInvocationId +OpDecorate %311 BuiltIn LocalInvocationId %2 = OpTypeVoid %4 = OpTypeInt 32 0 %3 = OpConstant %4 0 @@ -233,17 +229,14 @@ OpDecorate %337 BuiltIn LocalInvocationId %270 = OpTypePointer StorageBuffer %37 %271 = OpConstant %4 5 %288 = OpVariable %246 Output -%306 = OpConstantNull %6 -%310 = OpTypePointer StorageBuffer %6 -%313 = OpConstant %4 64 -%332 = OpConstantNull %62 -%336 = OpConstantNull %4 -%338 = OpTypePointer Input %35 -%337 = OpVariable %338 Input -%340 = OpConstantNull %35 -%342 = OpTypeBool -%341 = OpTypeVector %342 3 -%347 = OpConstant %4 264 +%306 = OpConstantNull %62 +%310 = OpConstantNull %4 +%312 = OpTypePointer Input %35 +%311 = OpVariable %312 Input +%314 = OpConstantNull %35 +%316 = OpTypeBool +%315 = OpTypeVector %316 3 +%321 = OpConstant %4 264 %93 = OpFunction %2 None %94 %92 = OpLabel %86 = OpVariable %87 Function %88 @@ -484,61 +477,26 @@ OpReturn OpFunctionEnd %308 = OpFunction %2 None %94 %307 = OpLabel -%305 = OpVariable %87 Function %306 +%305 = OpVariable %63 Function %306 OpBranch %309 %309 = OpLabel -%311 = OpAccessChain %310 %74 %30 -%312 = OpAtomicLoad %6 %311 %9 %313 -%315 = OpAccessChain %310 %74 %30 -%314 = OpAtomicIAdd %6 %315 %9 %313 %26 -OpStore %305 %314 -%317 = OpAccessChain %310 %74 %30 -%316 = OpAtomicISub %6 %317 %9 %313 %26 -OpStore %305 %316 -%319 = OpAccessChain %310 %74 %30 -%318 = OpAtomicAnd %6 %319 %9 %313 %26 -OpStore %305 %318 -%321 = OpAccessChain %310 %74 %30 -%320 = OpAtomicOr %6 %321 %9 %313 %26 -OpStore %305 %320 -%323 = OpAccessChain %310 %74 %30 -%322 = OpAtomicXor %6 %323 %9 %313 %26 -OpStore %305 %322 -%325 = OpAccessChain %310 %74 %30 -%324 = OpAtomicSMin %6 %325 %9 %313 %26 -OpStore %305 %324 -%327 = OpAccessChain %310 %74 %30 -%326 = OpAtomicSMax %6 %327 %9 %313 %26 -OpStore %305 %326 -%329 = OpAccessChain %310 %74 %30 -%328 = OpAtomicExchange %6 %329 %9 %313 %26 -OpStore %305 %328 -%330 = OpAccessChain %310 %74 %30 -OpAtomicStore %330 %9 %313 %312 -OpReturn -OpFunctionEnd -%334 = OpFunction %2 None %94 -%333 = OpLabel -%331 = OpVariable %63 Function %332 -OpBranch %335 -%335 = OpLabel -%339 = OpLoad %35 %337 -%343 = OpIEqual %341 %339 %340 -%344 = OpAll %342 %343 -OpSelectionMerge %345 None -OpBranchConditional %344 %346 %345 -%346 = OpLabel -OpStore %85 %336 -OpBranch %345 -%345 = OpLabel -OpControlBarrier %30 %30 %347 -OpBranch %348 -%348 = OpLabel -%349 = OpCompositeConstruct %57 %14 %14 %14 %14 -%350 = OpCompositeConstruct %57 %25 %25 %25 %25 -%351 = OpCompositeConstruct %62 %349 %350 -OpStore %331 %351 -%352 = OpFunctionCall %2 %225 %85 -%353 = OpFunctionCall %2 %230 %331 +%313 = OpLoad %35 %311 +%317 = OpIEqual %315 %313 %314 +%318 = OpAll %316 %317 +OpSelectionMerge %319 None +OpBranchConditional %318 %320 %319 +%320 = OpLabel +OpStore %85 %310 +OpBranch %319 +%319 = OpLabel +OpControlBarrier %30 %30 %321 +OpBranch %322 +%322 = OpLabel +%323 = OpCompositeConstruct %57 %14 %14 %14 %14 +%324 = OpCompositeConstruct %57 %25 %25 %25 %25 +%325 = OpCompositeConstruct %62 %323 %324 +OpStore %305 %325 +%326 = OpFunctionCall %2 %225 %85 +%327 = OpFunctionCall %2 %230 %305 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 7133f53d69..0bbd1a7759 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -165,31 +165,6 @@ fn foo_frag() -> @location(0) vec4 { return vec4(0.0); } -@compute @workgroup_size(1, 1, 1) -fn atomics() { - var tmp: i32; - - let value_1 = atomicLoad((&bar.atom)); - let _e7 = atomicAdd((&bar.atom), 5); - tmp = _e7; - let _e11 = atomicSub((&bar.atom), 5); - tmp = _e11; - let _e15 = atomicAnd((&bar.atom), 5); - tmp = _e15; - let _e19 = atomicOr((&bar.atom), 5); - tmp = _e19; - let _e23 = atomicXor((&bar.atom), 5); - tmp = _e23; - let _e27 = atomicMin((&bar.atom), 5); - tmp = _e27; - let _e31 = atomicMax((&bar.atom), 5); - tmp = _e31; - let _e35 = atomicExchange((&bar.atom), 5); - tmp = _e35; - atomicStore((&bar.atom), value_1); - return; -} - @compute @workgroup_size(1, 1, 1) fn assign_through_ptr() { var arr: array,2>;