brintos

brintos / llvm-project-archived public Read only

0
0
Text · 3.3 KiB · 6d2a096 Raw
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