58 lines · plain
1//===----------------------------------------------------------------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8 9#include <clc/mem_fence/clc_mem_fence.h>10 11#define BUILTIN_FENCE_ORDER(memory_order, ...) \12 switch (memory_order) { \13 case __ATOMIC_ACQUIRE: \14 __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, __VA_ARGS__); \15 break; \16 case __ATOMIC_RELEASE: \17 __builtin_amdgcn_fence(__ATOMIC_RELEASE, __VA_ARGS__); \18 break; \19 case __ATOMIC_ACQ_REL: \20 __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, __VA_ARGS__); \21 break; \22 case __ATOMIC_SEQ_CST: \23 __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, __VA_ARGS__); \24 break; \25 default: \26 __builtin_unreachable(); \27 } \28 break;29 30#define BUILTIN_FENCE(memory_scope, memory_order, ...) \31 switch (memory_scope) { \32 case __MEMORY_SCOPE_DEVICE: \33 BUILTIN_FENCE_ORDER(memory_order, "agent", ##__VA_ARGS__) \34 case __MEMORY_SCOPE_WRKGRP: \35 BUILTIN_FENCE_ORDER(memory_order, "workgroup", ##__VA_ARGS__) \36 case __MEMORY_SCOPE_WVFRNT: \37 BUILTIN_FENCE_ORDER(memory_order, "wavefront", ##__VA_ARGS__) \38 case __MEMORY_SCOPE_SINGLE: \39 BUILTIN_FENCE_ORDER(memory_order, "singlethread", ##__VA_ARGS__) \40 case __MEMORY_SCOPE_SYSTEM: \41 default: \42 BUILTIN_FENCE_ORDER(memory_order, "", ##__VA_ARGS__) \43 }44 45_CLC_OVERLOAD _CLC_DEF void46__clc_mem_fence(int memory_scope, int memory_order,47 __CLC_MemorySemantics memory_semantics) {48 if (memory_semantics == __CLC_MEMORY_LOCAL) {49 BUILTIN_FENCE(memory_scope, memory_order, "local")50 } else if (memory_semantics == __CLC_MEMORY_GLOBAL) {51 BUILTIN_FENCE(memory_scope, memory_order, "global")52 } else if (memory_semantics == (__CLC_MEMORY_LOCAL | __CLC_MEMORY_GLOBAL)) {53 BUILTIN_FENCE(memory_scope, memory_order, "local", "global")54 } else {55 BUILTIN_FENCE(memory_scope, memory_order)56 }57}58