forked from OSchip/llvm-project
40 lines
1.0 KiB
Common Lisp
40 lines
1.0 KiB
Common Lisp
|
#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)
|
||
|
|
#endif
|
||
|
|
|
||
|
|
_CLC_DEF 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 void read_mem_fence(cl_mem_fence_flags flags)
|
||
|
|
{
|
||
|
|
mem_fence(flags);
|
||
|
|
}
|
||
|
|
|
||
|
|
_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags)
|
||
|
|
{
|
||
|
|
mem_fence(flags);
|
||
|
|
}
|