1; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV 2; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} 3 4;; This test checks that the backend is capable to correctly translate 5;; atomic_work_item_fence OpenCL C 2.0 built-in function [1] into corresponding 6;; SPIR-V instruction [2]. 7 8;; __kernel void test_mem_fence_const_flags() { 9;; atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_relaxed, memory_scope_work_item); 10;; atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_acquire, memory_scope_work_group); 11;; atomic_work_item_fence(CLK_IMAGE_MEM_FENCE, memory_order_release, memory_scope_device); 12;; atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acq_rel, memory_scope_all_svm_devices); 13;; atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_seq_cst, memory_scope_sub_group); 14;; atomic_work_item_fence(CLK_IMAGE_MEM_FENCE | CLK_LOCAL_MEM_FENCE, memory_order_acquire, memory_scope_sub_group); 15;; } 16 17;; __kernel void test_mem_fence_non_const_flags(cl_mem_fence_flags flags, memory_order order, memory_scope scope) { 18;; // FIXME: OpenCL spec doesn't require flags to be compile-time known 19;; // atomic_work_item_fence(flags, order, scope); 20;; } 21 22; CHECK-SPIRV: OpName %[[#TEST_CONST_FLAGS:]] "test_mem_fence_const_flags" 23; CHECK-SPIRV: %[[#UINT:]] = OpTypeInt 32 0 24 25;; 0x0 Relaxed + 0x100 WorkgroupMemory 26; CHECK-SPIRV-DAG: %[[#LOCAL_RELAXED:]] = OpConstant %[[#UINT]] 256 27;; 0x2 Acquire + 0x200 CrossWorkgroupMemory 28; CHECK-SPIRV-DAG: %[[#GLOBAL_ACQUIRE:]] = OpConstant %[[#UINT]] 514 29;; 0x4 Release + 0x800 ImageMemory 30; CHECK-SPIRV-DAG: %[[#IMAGE_RELEASE:]] = OpConstant %[[#UINT]] 2052 31;; 0x8 AcquireRelease + 0x100 WorkgroupMemory 32; CHECK-SPIRV-DAG: %[[#LOCAL_ACQ_REL:]] = OpConstant %[[#UINT]] 264 33;; 0x10 SequentiallyConsistent + 0x200 CrossWorkgroupMemory 34; CHECK-SPIRV-DAG: %[[#GLOBAL_SEQ_CST:]] = OpConstant %[[#UINT]] 528 35;; 0x2 Acquire + 0x100 WorkgroupMemory + 0x800 ImageMemory 36; CHECK-SPIRV-DAG: %[[#LOCAL_IMAGE_ACQUIRE:]] = OpConstant %[[#UINT]] 2306 37 38;; Scopes [4]: 39;; 4 Invocation 40; CHECK-SPIRV-DAG: %[[#SCOPE_INVOCATION:]] = OpConstant %[[#UINT]] 4 41;; 2 Workgroup 42; CHECK-SPIRV-DAG: %[[#SCOPE_WORK_GROUP:]] = OpConstant %[[#UINT]] 2 43;; 1 Device 44; CHECK-SPIRV-DAG: %[[#SCOPE_DEVICE:]] = OpConstant %[[#UINT]] 1 45;; 0 CrossDevice 46; CHECK-SPIRV-DAG: %[[#SCOPE_CROSS_DEVICE:]] = OpConstant %[[#UINT]] 0 47;; 3 Subgroup 48; CHECK-SPIRV-DAG: %[[#SCOPE_SUBGROUP:]] = OpConstant %[[#UINT]] 3 49 50; CHECK-SPIRV: %[[#TEST_CONST_FLAGS]] = OpFunction %[[#]] 51; CHECK-SPIRV: OpMemoryBarrier %[[#SCOPE_INVOCATION]] %[[#LOCAL_RELAXED]] 52; CHECK-SPIRV: OpMemoryBarrier %[[#SCOPE_WORK_GROUP]] %[[#GLOBAL_ACQUIRE]] 53; CHECK-SPIRV: OpMemoryBarrier %[[#SCOPE_DEVICE]] %[[#IMAGE_RELEASE]] 54; CHECK-SPIRV: OpMemoryBarrier %[[#SCOPE_CROSS_DEVICE]] %[[#LOCAL_ACQ_REL]] 55; CHECK-SPIRV: OpMemoryBarrier %[[#SCOPE_SUBGROUP]] %[[#GLOBAL_SEQ_CST]] 56; CHECK-SPIRV: OpMemoryBarrier %[[#SCOPE_SUBGROUP]] %[[#LOCAL_IMAGE_ACQUIRE]] 57 58define dso_local spir_kernel void @test_mem_fence_const_flags() local_unnamed_addr { 59entry: 60 tail call spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef 1, i32 noundef 0, i32 noundef 0) 61 tail call spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef 2, i32 noundef 2, i32 noundef 1) 62 tail call spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef 4, i32 noundef 3, i32 noundef 2) 63 tail call spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef 1, i32 noundef 4, i32 noundef 3) 64 tail call spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef 2, i32 noundef 5, i32 noundef 4) 65 tail call spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef 5, i32 noundef 2, i32 noundef 4) 66 ret void 67} 68 69declare spir_func void @_Z22atomic_work_item_fencej12memory_order12memory_scope(i32 noundef, i32 noundef, i32 noundef) local_unnamed_addr 70 71define dso_local spir_kernel void @test_mem_fence_non_const_flags(i32 noundef %flags, i32 noundef %order, i32 noundef %scope) local_unnamed_addr { 72entry: 73 ret void 74} 75 76;; References: 77;; [1]: https://www.khronos.org/registry/OpenCL/sdk/2.0/docs/man/xhtml/atomic_work_item_fence.html 78;; [2]: https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#OpMemoryBarrier 79