[libclc] Move several 'native' builtins to CLC library

libclc/CMakeLists.txt

@@ -263,6 +263,23 @@ if ( clspv-- IN_LIST LIBCLC_TARGETS_TO_BUILD OR clspv64-- IN_LIST LIBCLC_TARGETS
 endif()
 
 set_source_files_properties(
+  # CLC builtins
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_cos.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_exp2.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_exp.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_log10.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_log2.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_log.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_rsqrt.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_sin.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/generic/math/clc_native_sqrt.cl
+  # Target-specific CLC builtins
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/amdgpu/math/clc_native_exp2.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/amdgpu/math/clc_native_exp.cl
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/amdgpu/math/clc_native_log10.cl
+  # Target-specific OpenCL builtins
+  ${CMAKE_CURRENT_SOURCE_DIR}/clc/lib/r600/math/clc_native_rsqrt.cl
+  # OpenCL builtins
   ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/math/native_cos.cl
   ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/math/native_divide.cl
   ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/math/native_exp.cl
@@ -277,10 +294,6 @@ set_source_files_properties(
   ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/math/native_sin.cl
   ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/math/native_sqrt.cl
   ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/math/native_tan.cl
-  ${CMAKE_CURRENT_SOURCE_DIR}/amdgpu/lib/math/native_exp.cl
-  ${CMAKE_CURRENT_SOURCE_DIR}/amdgpu/lib/math/native_log.cl
-  ${CMAKE_CURRENT_SOURCE_DIR}/amdgpu/lib/math/native_log10.cl
-  ${CMAKE_CURRENT_SOURCE_DIR}/r600/lib/math/native_rsqrt.cl
   PROPERTIES COMPILE_OPTIONS -fapprox-func
 )
 

libclc/amdgpu/lib/SOURCES

@@ -1,7 +1,3 @@
-math/native_exp.cl
-math/native_exp2.cl
-math/native_log.cl
-math/native_log10.cl
 math/half_exp.cl
 math/half_exp10.cl
 math/half_exp2.cl

libclc/clc/include/clc/math/clc_native_cos.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_COS_H__
+#define __CLC_MATH_CLC_NATIVE_COS_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_cos
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_COS_H__

libclc/clc/include/clc/math/clc_native_exp.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_EXP_H__
+#define __CLC_MATH_CLC_NATIVE_EXP_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_exp
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_EXP_H__

libclc/clc/include/clc/math/clc_native_exp2.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_EXP2_H__
+#define __CLC_MATH_CLC_NATIVE_EXP2_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_exp2
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_EXP2_H__

libclc/clc/include/clc/math/clc_native_log.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_LOG_H__
+#define __CLC_MATH_CLC_NATIVE_LOG_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG_H__

libclc/clc/include/clc/math/clc_native_log10.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_LOG10_H__
+#define __CLC_MATH_CLC_NATIVE_LOG10_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log10
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG10_H__

libclc/clc/include/clc/math/clc_native_log2.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_LOG2_H__
+#define __CLC_MATH_CLC_NATIVE_LOG2_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log2
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG2_H__

libclc/clc/include/clc/math/clc_native_rsqrt.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_RSQRT_H__
+#define __CLC_MATH_CLC_NATIVE_RSQRT_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_rsqrt
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_RSQRT_H__

libclc/clc/include/clc/math/clc_native_sin.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_SIN_H__
+#define __CLC_MATH_CLC_NATIVE_SIN_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_sin
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_SIN_H__

libclc/clc/include/clc/math/clc_native_sqrt.h

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_SQRT_H__
+#define __CLC_MATH_CLC_NATIVE_SQRT_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_sqrt
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_SQRT_H__

libclc/clc/lib/amdgpu/SOURCES

@@ -1 +1,4 @@
+math/clc_native_exp2.cl
+math/clc_native_exp.cl
+math/clc_native_log10.cl
 math/clc_sqrt_fp64.cl

libclc/amdgpu/lib/math/native_log10.cl → libclc/clc/lib/amdgpu/math/clc_native_exp.cl

@@ -6,8 +6,10 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/clc.h>
+#include <clc/float/definitions.h>
+#include <clc/internal/clc.h>
+#include <clc/math/clc_native_exp2.h>
 
-#define __CLC_BODY <native_log10.inc>
+#define __CLC_BODY <clc_native_exp.inc>
 #define __FLOAT_ONLY
 #include <clc/math/gentype.inc>

libclc/amdgpu/lib/math/native_exp.inc → libclc/clc/lib/amdgpu/math/clc_native_exp.inc

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_exp(__CLC_GENTYPE val) {
-  return native_exp2(val * M_LOG2E_F);
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_exp(__CLC_GENTYPE val) {
+  return __clc_native_exp2(val * M_LOG2E_F);
 }

libclc/amdgpu/lib/math/native_exp2.cl → libclc/clc/lib/amdgpu/math/clc_native_exp2.cl

@@ -6,11 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/clc.h>
 #include <clc/clcmacro.h>
+#include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DEF float native_exp2(float val) {
+_CLC_OVERLOAD _CLC_DEF float __clc_native_exp2(float val) {
   return __builtin_amdgcn_exp2f(val);
 }
 
-_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, native_exp2, float)
+_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_native_exp2, float)

libclc/amdgpu/lib/math/native_log.cl → libclc/clc/lib/amdgpu/math/clc_native_log10.cl

@@ -6,8 +6,10 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/clc.h>
+#include <clc/float/definitions.h>
+#include <clc/internal/clc.h>
+#include <clc/math/clc_native_log2.h>
 
-#define __CLC_BODY <native_log.inc>
+#define __CLC_BODY <clc_native_log10.inc>
 #define __FLOAT_ONLY
 #include <clc/math/gentype.inc>

libclc/amdgpu/lib/math/native_log10.inc → libclc/clc/lib/amdgpu/math/clc_native_log10.inc

@@ -6,6 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_log10(__CLC_GENTYPE val) {
-  return native_log2(val) * (M_LN2_F / M_LN10_F);
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_log10(__CLC_GENTYPE val) {
+  return __clc_native_log2(val) * (M_LN2_F / M_LN10_F);
 }

libclc/clc/lib/generic/SOURCES

@@ -45,6 +45,15 @@ math/clc_log2.cl
 math/clc_mad.cl
 math/clc_modf.cl
 math/clc_nan.cl
+math/clc_native_cos.cl
+math/clc_native_exp.cl
+math/clc_native_exp2.cl
+math/clc_native_log.cl
+math/clc_native_log10.cl
+math/clc_native_log2.cl
+math/clc_native_rsqrt.cl
+math/clc_native_sin.cl
+math/clc_native_sqrt.cl
 math/clc_nextafter.cl
 math/clc_pow.cl
 math/clc_pown.cl