7d048674a4
This commit bulk updates all '.h', '.cl', '.inc', and '.cpp' files to add any missing license headers. The remaining files are generally CMake, SOURCES, scripts, markdown, etc. There are still some '.ll' files which may benefit from a license header. I can't find an example of an LLVM IR file with a license header in the rest of LLVM, but unlike most other (sub)projects, libclc has examples of LLVM IR as source files, compiled and built into the library.
46 lines
1.5 KiB
Common Lisp
46 lines
1.5 KiB
Common Lisp
//===----------------------------------------------------------------------===//
|
|
//
|
|
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
// See https://llvm.org/LICENSE.txt for license information.
|
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
#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);
|
|
}
|