diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 59a70a200c95..efe7f5804e8f 100644 --- a/libclc/CMakeLists.txt +++ b/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 ) diff --git a/libclc/amdgpu/lib/SOURCES b/libclc/amdgpu/lib/SOURCES index ed5e45a37c18..ab5da40711aa 100644 --- a/libclc/amdgpu/lib/SOURCES +++ b/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 diff --git a/libclc/amdgpu/lib/math/native_log.inc b/libclc/amdgpu/lib/math/native_log.inc deleted file mode 100644 index 820e4929f02c..000000000000 --- a/libclc/amdgpu/lib/math/native_log.inc +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 -// -//===----------------------------------------------------------------------===// - -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_log(__CLC_GENTYPE val) { - return native_log2(val) * (1.0f / M_LOG2E_F); -} diff --git a/libclc/clc/include/clc/math/clc_native_cos.h b/libclc/clc/include/clc/math/clc_native_cos.h new file mode 100644 index 000000000000..8a580f13ad2a --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_COS_H__ diff --git a/libclc/clc/include/clc/math/clc_native_exp.h b/libclc/clc/include/clc/math/clc_native_exp.h new file mode 100644 index 000000000000..48a1be616ea3 --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_EXP_H__ diff --git a/libclc/clc/include/clc/math/clc_native_exp2.h b/libclc/clc/include/clc/math/clc_native_exp2.h new file mode 100644 index 000000000000..bc0b32d6212b --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_EXP2_H__ diff --git a/libclc/clc/include/clc/math/clc_native_log.h b/libclc/clc/include/clc/math/clc_native_log.h new file mode 100644 index 000000000000..ea0362503f67 --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_LOG_H__ diff --git a/libclc/clc/include/clc/math/clc_native_log10.h b/libclc/clc/include/clc/math/clc_native_log10.h new file mode 100644 index 000000000000..c5cceeeba295 --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_LOG10_H__ diff --git a/libclc/clc/include/clc/math/clc_native_log2.h b/libclc/clc/include/clc/math/clc_native_log2.h new file mode 100644 index 000000000000..25375970cedc --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_LOG2_H__ diff --git a/libclc/clc/include/clc/math/clc_native_rsqrt.h b/libclc/clc/include/clc/math/clc_native_rsqrt.h new file mode 100644 index 000000000000..59fd2134107d --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_RSQRT_H__ diff --git a/libclc/clc/include/clc/math/clc_native_sin.h b/libclc/clc/include/clc/math/clc_native_sin.h new file mode 100644 index 000000000000..878e19882ae6 --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_SIN_H__ diff --git a/libclc/clc/include/clc/math/clc_native_sqrt.h b/libclc/clc/include/clc/math/clc_native_sqrt.h new file mode 100644 index 000000000000..afff77c89b22 --- /dev/null +++ b/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 + +#include + +#undef __CLC_BODY +#undef __CLC_FUNCTION +#undef __FLOAT_ONLY + +#endif // __CLC_MATH_CLC_NATIVE_SQRT_H__ diff --git a/libclc/clc/include/clc/math/unary_intrin.inc b/libclc/clc/include/clc/math/unary_intrin.inc deleted file mode 100644 index 8028470114b8..000000000000 --- a/libclc/clc/include/clc/math/unary_intrin.inc +++ /dev/null @@ -1,42 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 -#include - -_CLC_OVERLOAD float __CLC_FUNCTION(float f) __asm(__CLC_INTRINSIC ".f32"); -_CLC_OVERLOAD float2 __CLC_FUNCTION(float2 f) __asm(__CLC_INTRINSIC ".v2f32"); -_CLC_OVERLOAD float3 __CLC_FUNCTION(float3 f) __asm(__CLC_INTRINSIC ".v3f32"); -_CLC_OVERLOAD float4 __CLC_FUNCTION(float4 f) __asm(__CLC_INTRINSIC ".v4f32"); -_CLC_OVERLOAD float8 __CLC_FUNCTION(float8 f) __asm(__CLC_INTRINSIC ".v8f32"); -_CLC_OVERLOAD float16 __CLC_FUNCTION(float16 f) __asm(__CLC_INTRINSIC - ".v16f32"); - -#ifdef cl_khr_fp64 -#pragma OPENCL EXTENSION cl_khr_fp64 : enable -_CLC_OVERLOAD double __CLC_FUNCTION(double d) __asm(__CLC_INTRINSIC ".f64"); -_CLC_OVERLOAD double2 __CLC_FUNCTION(double2 d) __asm(__CLC_INTRINSIC ".v2f64"); -_CLC_OVERLOAD double3 __CLC_FUNCTION(double3 d) __asm(__CLC_INTRINSIC ".v3f64"); -_CLC_OVERLOAD double4 __CLC_FUNCTION(double4 d) __asm(__CLC_INTRINSIC ".v4f64"); -_CLC_OVERLOAD double8 __CLC_FUNCTION(double8 d) __asm(__CLC_INTRINSIC ".v8f64"); -_CLC_OVERLOAD double16 __CLC_FUNCTION(double16 d) __asm(__CLC_INTRINSIC - ".v16f64"); -#endif - -#ifdef cl_khr_fp16 -#pragma OPENCL EXTENSION cl_khr_fp16 : enable -_CLC_OVERLOAD half __CLC_FUNCTION(half d) __asm(__CLC_INTRINSIC ".f16"); -_CLC_OVERLOAD half2 __CLC_FUNCTION(half2 d) __asm(__CLC_INTRINSIC ".v2f16"); -_CLC_OVERLOAD half3 __CLC_FUNCTION(half3 d) __asm(__CLC_INTRINSIC ".v3f16"); -_CLC_OVERLOAD half4 __CLC_FUNCTION(half4 d) __asm(__CLC_INTRINSIC ".v4f16"); -_CLC_OVERLOAD half8 __CLC_FUNCTION(half8 d) __asm(__CLC_INTRINSIC ".v8f16"); -_CLC_OVERLOAD half16 __CLC_FUNCTION(half16 d) __asm(__CLC_INTRINSIC ".v16f16"); -#endif - -#undef __CLC_FUNCTION -#undef __CLC_INTRINSIC diff --git a/libclc/clc/lib/amdgpu/SOURCES b/libclc/clc/lib/amdgpu/SOURCES index fd64a862021e..31e07b608c4c 100644 --- a/libclc/clc/lib/amdgpu/SOURCES +++ b/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 diff --git a/libclc/amdgpu/lib/math/native_log10.cl b/libclc/clc/lib/amdgpu/math/clc_native_exp.cl similarity index 74% rename from libclc/amdgpu/lib/math/native_log10.cl rename to libclc/clc/lib/amdgpu/math/clc_native_exp.cl index 7cbe1f98988d..591ecb0ac00b 100644 --- a/libclc/amdgpu/lib/math/native_log10.cl +++ b/libclc/clc/lib/amdgpu/math/clc_native_exp.cl @@ -6,8 +6,10 @@ // //===----------------------------------------------------------------------===// -#include +#include +#include +#include -#define __CLC_BODY +#define __CLC_BODY #define __FLOAT_ONLY #include diff --git a/libclc/amdgpu/lib/math/native_exp.inc b/libclc/clc/lib/amdgpu/math/clc_native_exp.inc similarity index 75% rename from libclc/amdgpu/lib/math/native_exp.inc rename to libclc/clc/lib/amdgpu/math/clc_native_exp.inc index d7dbd888d298..cf2d48c4054e 100644 --- a/libclc/amdgpu/lib/math/native_exp.inc +++ b/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); } diff --git a/libclc/amdgpu/lib/math/native_exp2.cl b/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl similarity index 72% rename from libclc/amdgpu/lib/math/native_exp2.cl rename to libclc/clc/lib/amdgpu/math/clc_native_exp2.cl index 39ae914b1963..76b1850fce57 100644 --- a/libclc/amdgpu/lib/math/native_exp2.cl +++ b/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl @@ -6,11 +6,11 @@ // //===----------------------------------------------------------------------===// -#include #include +#include -_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) diff --git a/libclc/amdgpu/lib/math/native_log.cl b/libclc/clc/lib/amdgpu/math/clc_native_log10.cl similarity index 74% rename from libclc/amdgpu/lib/math/native_log.cl rename to libclc/clc/lib/amdgpu/math/clc_native_log10.cl index ce3258bb5b37..0668a635d24d 100644 --- a/libclc/amdgpu/lib/math/native_log.cl +++ b/libclc/clc/lib/amdgpu/math/clc_native_log10.cl @@ -6,8 +6,10 @@ // //===----------------------------------------------------------------------===// -#include +#include +#include +#include -#define __CLC_BODY +#define __CLC_BODY #define __FLOAT_ONLY #include diff --git a/libclc/amdgpu/lib/math/native_log10.inc b/libclc/clc/lib/amdgpu/math/clc_native_log10.inc similarity index 73% rename from libclc/amdgpu/lib/math/native_log10.inc rename to libclc/clc/lib/amdgpu/math/clc_native_log10.inc index 2aef43a8ed1a..c91d69860979 100644 --- a/libclc/amdgpu/lib/math/native_log10.inc +++ b/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); } diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES index 6a1ba9536222..c31963c59e95 100644 --- a/libclc/clc/lib/generic/SOURCES +++ b/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 diff --git a/libclc/clc/lib/generic/math/clc_native_cos.cl b/libclc/clc/lib/generic/math/clc_native_cos.cl new file mode 100644 index 000000000000..de56fdec48d2 --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_cos.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_cos +#define __CLC_FUNCTION(x) __builtin_elementwise_cos +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/generic/math/clc_native_exp.cl b/libclc/clc/lib/generic/math/clc_native_exp.cl new file mode 100644 index 000000000000..400270a6163a --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_exp.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_exp +#define __CLC_FUNCTION(x) __builtin_elementwise_exp +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/generic/math/clc_native_exp2.cl b/libclc/clc/lib/generic/math/clc_native_exp2.cl new file mode 100644 index 000000000000..427d901fcdb1 --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_exp2.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_exp2 +#define __CLC_FUNCTION(x) __builtin_elementwise_exp2 +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/generic/math/clc_native_log.cl b/libclc/clc/lib/generic/math/clc_native_log.cl new file mode 100644 index 000000000000..85f188b65428 --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_log.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_log +#define __CLC_FUNCTION(x) __builtin_elementwise_log +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/generic/math/clc_native_log10.cl b/libclc/clc/lib/generic/math/clc_native_log10.cl new file mode 100644 index 000000000000..624018e4481b --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_log10.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_log10 +#define __CLC_FUNCTION(x) __builtin_elementwise_log10 +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/generic/math/clc_native_log2.cl b/libclc/clc/lib/generic/math/clc_native_log2.cl new file mode 100644 index 000000000000..2c8c18e61ca5 --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_log2.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_log2 +#define __CLC_FUNCTION(x) __builtin_elementwise_log2 +#define __CLC_BODY + +#include diff --git a/libclc/amdgpu/lib/math/native_exp.cl b/libclc/clc/lib/generic/math/clc_native_rsqrt.cl similarity index 79% rename from libclc/amdgpu/lib/math/native_exp.cl rename to libclc/clc/lib/generic/math/clc_native_rsqrt.cl index e62b79d4ec9f..d5e6fcdae491 100644 --- a/libclc/amdgpu/lib/math/native_exp.cl +++ b/libclc/clc/lib/generic/math/clc_native_rsqrt.cl @@ -6,8 +6,9 @@ // //===----------------------------------------------------------------------===// -#include +#include +#include -#define __CLC_BODY +#define __CLC_BODY #define __FLOAT_ONLY #include diff --git a/libclc/generic/lib/math/native_rsqrt.inc b/libclc/clc/lib/generic/math/clc_native_rsqrt.inc similarity index 76% rename from libclc/generic/lib/math/native_rsqrt.inc rename to libclc/clc/lib/generic/math/clc_native_rsqrt.inc index 058209bcb8a1..7a3b0b2af272 100644 --- a/libclc/generic/lib/math/native_rsqrt.inc +++ b/libclc/clc/lib/generic/math/clc_native_rsqrt.inc @@ -6,6 +6,6 @@ // //===----------------------------------------------------------------------===// -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_rsqrt(__CLC_GENTYPE val) { - return 1.0f / native_sqrt(val); +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_rsqrt(__CLC_GENTYPE val) { + return 1.0f / __clc_native_sqrt(val); } diff --git a/libclc/clc/lib/generic/math/clc_native_sin.cl b/libclc/clc/lib/generic/math/clc_native_sin.cl new file mode 100644 index 000000000000..22b988bf4375 --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_sin.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_sin +#define __CLC_FUNCTION(x) __builtin_elementwise_sin +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/generic/math/clc_native_sqrt.cl b/libclc/clc/lib/generic/math/clc_native_sqrt.cl new file mode 100644 index 000000000000..ed022ef1fee1 --- /dev/null +++ b/libclc/clc/lib/generic/math/clc_native_sqrt.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_sqrt +#define __CLC_FUNCTION(x) __builtin_elementwise_sqrt +#define __CLC_BODY + +#include diff --git a/libclc/clc/lib/r600/SOURCES b/libclc/clc/lib/r600/SOURCES index 8f66107e0454..8d5caf167aa4 100644 --- a/libclc/clc/lib/r600/SOURCES +++ b/libclc/clc/lib/r600/SOURCES @@ -1 +1,2 @@ +math/clc_native_rsqrt.cl math/clc_rsqrt_override.cl diff --git a/libclc/r600/lib/math/native_rsqrt.cl b/libclc/clc/lib/r600/math/clc_native_rsqrt.cl similarity index 65% rename from libclc/r600/lib/math/native_rsqrt.cl rename to libclc/clc/lib/r600/math/clc_native_rsqrt.cl index e916147c3e05..ee09814eb1e7 100644 --- a/libclc/r600/lib/math/native_rsqrt.cl +++ b/libclc/clc/lib/r600/math/clc_native_rsqrt.cl @@ -6,12 +6,11 @@ // //===----------------------------------------------------------------------===// -#include #include +#include -_CLC_OVERLOAD _CLC_DEF float native_rsqrt(float x) -{ - return __builtin_r600_recipsqrt_ieeef(x); +_CLC_OVERLOAD _CLC_DEF float __clc_native_rsqrt(float x) { + return __builtin_r600_recipsqrt_ieeef(x); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, native_rsqrt, float); +_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_native_rsqrt, float); diff --git a/libclc/generic/lib/math/native_cos.cl b/libclc/generic/lib/math/native_cos.cl index d81a067f1104..0f4d46c1d73e 100644 --- a/libclc/generic/lib/math/native_cos.cl +++ b/libclc/generic/lib/math/native_cos.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC cos - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_cos +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_exp.cl b/libclc/generic/lib/math/native_exp.cl index 8f4531343415..c1d08ec2c7e1 100644 --- a/libclc/generic/lib/math/native_exp.cl +++ b/libclc/generic/lib/math/native_exp.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC exp - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_exp +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_exp2.cl b/libclc/generic/lib/math/native_exp2.cl index ecde4a6761d2..ceb570733b97 100644 --- a/libclc/generic/lib/math/native_exp2.cl +++ b/libclc/generic/lib/math/native_exp2.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC exp2 - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_exp2 +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_log.cl b/libclc/generic/lib/math/native_log.cl index 5731e09d21c9..adc2ff495f8b 100644 --- a/libclc/generic/lib/math/native_log.cl +++ b/libclc/generic/lib/math/native_log.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC log - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_log +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_log10.cl b/libclc/generic/lib/math/native_log10.cl index eab7a6f14d03..f63292124f3b 100644 --- a/libclc/generic/lib/math/native_log10.cl +++ b/libclc/generic/lib/math/native_log10.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC log10 - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_log10 +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_log2.cl b/libclc/generic/lib/math/native_log2.cl index 0db4be0b5e08..6b079872b1e0 100644 --- a/libclc/generic/lib/math/native_log2.cl +++ b/libclc/generic/lib/math/native_log2.cl @@ -7,8 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC log2 -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_log2 +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_rsqrt.cl b/libclc/generic/lib/math/native_rsqrt.cl index 14430c04fb72..cb49b2d1d670 100644 --- a/libclc/generic/lib/math/native_rsqrt.cl +++ b/libclc/generic/lib/math/native_rsqrt.cl @@ -7,7 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_rsqrt +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_sin.cl b/libclc/generic/lib/math/native_sin.cl index 0e2ced09fa2d..50265b393627 100644 --- a/libclc/generic/lib/math/native_sin.cl +++ b/libclc/generic/lib/math/native_sin.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC sin - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_sin +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_sqrt.cl b/libclc/generic/lib/math/native_sqrt.cl index 1b668e5976ef..4cd022e8bbeb 100644 --- a/libclc/generic/lib/math/native_sqrt.cl +++ b/libclc/generic/lib/math/native_sqrt.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include +#include -#define __CLC_NATIVE_INTRINSIC sqrt - -#define __CLC_BODY #define __FLOAT_ONLY +#define FUNCTION native_sqrt +#define __CLC_BODY + #include diff --git a/libclc/generic/lib/math/native_unary_intrinsic.inc b/libclc/generic/lib/math/native_unary_intrinsic.inc deleted file mode 100644 index c118ec095692..000000000000 --- a/libclc/generic/lib/math/native_unary_intrinsic.inc +++ /dev/null @@ -1,26 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 - -#ifdef __CLC_SCALAR -#define __CLC_FUNCTION __CLC_XCONCAT(__clc_native_, __CLC_NATIVE_INTRINSIC) -#define __CLC_INTRINSIC "llvm." __CLC_XSTR(__CLC_NATIVE_INTRINSIC) - -#undef cl_khr_fp64 -#include - -#endif - -#define __CLC_FUNCTION __CLC_XCONCAT(native_, __CLC_NATIVE_INTRINSIC) - -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __CLC_FUNCTION(__CLC_GENTYPE val) { - return __CLC_XCONCAT(__clc_native_, __CLC_NATIVE_INTRINSIC)(val); -} - -#undef __CLC_FUNCTION diff --git a/libclc/r600/lib/SOURCES b/libclc/r600/lib/SOURCES index cad45785dc48..4342ac38201c 100644 --- a/libclc/r600/lib/SOURCES +++ b/libclc/r600/lib/SOURCES @@ -1,6 +1,5 @@ math/fmax.cl math/fmin.cl -math/native_rsqrt.cl synchronization/barrier.cl workitem/get_global_offset.cl workitem/get_group_id.cl