fence.cl
1.15 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
#include <clc/clc.h>
void __clc_amdgcn_s_waitcnt(unsigned flags);
// s_waitcnt takes 16bit argument with a combined number of maximum allowed
// pending operations:
// [12:8] LGKM -- LDS, GDS, Konstant (SMRD), Messages
// [7] -- undefined
// [6:4] -- exports, GDS, and mem write
// [3:0] -- vector memory operations
// Newer clang supports __builtin_amdgcn_s_waitcnt
#if __clang_major__ >= 5
# define __waitcnt(x) __builtin_amdgcn_s_waitcnt(x)
#else
# define __waitcnt(x) __clc_amdgcn_s_waitcnt(x)
_CLC_DEF void __clc_amdgcn_s_waitcnt(unsigned) __asm("llvm.amdgcn.s.waitcnt");
#endif
_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) {
if (flags & CLK_GLOBAL_MEM_FENCE) {
// scalar loads are counted with LGKM but we don't know whether
// the compiler turned any loads to scalar
__waitcnt(0);
} else if (flags & CLK_LOCAL_MEM_FENCE)
__waitcnt(0xff); // LGKM is [12:8]
}
#undef __waitcnt
// We don't have separate mechanism for read and write fences
_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {
mem_fence(flags);
}
_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) {
mem_fence(flags);
}