From ad48fffb53003333456b327a2a3f22bd81a77a00 Mon Sep 17 00:00:00 2001 From: Fraser Cormack Date: Tue, 1 Apr 2025 09:20:54 +0100 Subject: [PATCH] [libclc] Move several 'native' builtins to CLC library (#129679) This commit moves the 'native' builtins that use asm statements to generate LLVM intrinsics to the CLC library. In doing so it converts them to use the appropriate elementwise builtin to generate the same intrinsic; there are no codegen changes to any target except to AMDGPU targets where `native_log` is no longer custom implemented and instead used the clang elementwise builtin. This work forms part of #127196 and indeed with this commit there are no 'generic' builtins using/abusing asm statements - the remaining builtins are specific to the amdgpu and r600 targets. --- libclc/CMakeLists.txt | 21 ++++++++-- libclc/amdgpu/lib/SOURCES | 4 -- libclc/amdgpu/lib/math/native_log.inc | 11 ----- libclc/clc/include/clc/math/clc_native_cos.h | 22 ++++++++++ libclc/clc/include/clc/math/clc_native_exp.h | 22 ++++++++++ libclc/clc/include/clc/math/clc_native_exp2.h | 22 ++++++++++ libclc/clc/include/clc/math/clc_native_log.h | 22 ++++++++++ .../clc/include/clc/math/clc_native_log10.h | 22 ++++++++++ libclc/clc/include/clc/math/clc_native_log2.h | 22 ++++++++++ .../clc/include/clc/math/clc_native_rsqrt.h | 22 ++++++++++ libclc/clc/include/clc/math/clc_native_sin.h | 22 ++++++++++ libclc/clc/include/clc/math/clc_native_sqrt.h | 22 ++++++++++ libclc/clc/include/clc/math/unary_intrin.inc | 42 ------------------- libclc/clc/lib/amdgpu/SOURCES | 3 ++ .../lib/amdgpu/math/clc_native_exp.cl} | 6 ++- .../lib/amdgpu/math/clc_native_exp.inc} | 4 +- .../lib/amdgpu/math/clc_native_exp2.cl} | 6 +-- .../lib/amdgpu/math/clc_native_log10.cl} | 6 ++- .../lib/amdgpu/math/clc_native_log10.inc} | 4 +- libclc/clc/lib/generic/SOURCES | 9 ++++ libclc/clc/lib/generic/math/clc_native_cos.cl | 16 +++++++ libclc/clc/lib/generic/math/clc_native_exp.cl | 16 +++++++ .../clc/lib/generic/math/clc_native_exp2.cl | 16 +++++++ libclc/clc/lib/generic/math/clc_native_log.cl | 16 +++++++ .../clc/lib/generic/math/clc_native_log10.cl | 16 +++++++ .../clc/lib/generic/math/clc_native_log2.cl | 16 +++++++ .../lib/generic/math/clc_native_rsqrt.cl} | 5 ++- .../lib/generic/math/clc_native_rsqrt.inc} | 4 +- libclc/clc/lib/generic/math/clc_native_sin.cl | 16 +++++++ .../clc/lib/generic/math/clc_native_sqrt.cl | 16 +++++++ libclc/clc/lib/r600/SOURCES | 1 + .../lib/r600/math/clc_native_rsqrt.cl} | 9 ++-- libclc/generic/lib/math/native_cos.cl | 7 ++-- libclc/generic/lib/math/native_exp.cl | 7 ++-- libclc/generic/lib/math/native_exp2.cl | 7 ++-- libclc/generic/lib/math/native_log.cl | 7 ++-- libclc/generic/lib/math/native_log10.cl | 7 ++-- libclc/generic/lib/math/native_log2.cl | 6 ++- libclc/generic/lib/math/native_rsqrt.cl | 5 ++- libclc/generic/lib/math/native_sin.cl | 7 ++-- libclc/generic/lib/math/native_sqrt.cl | 7 ++-- .../lib/math/native_unary_intrinsic.inc | 26 ------------ libclc/r600/lib/SOURCES | 1 - 43 files changed, 416 insertions(+), 132 deletions(-) delete mode 100644 libclc/amdgpu/lib/math/native_log.inc create mode 100644 libclc/clc/include/clc/math/clc_native_cos.h create mode 100644 libclc/clc/include/clc/math/clc_native_exp.h create mode 100644 libclc/clc/include/clc/math/clc_native_exp2.h create mode 100644 libclc/clc/include/clc/math/clc_native_log.h create mode 100644 libclc/clc/include/clc/math/clc_native_log10.h create mode 100644 libclc/clc/include/clc/math/clc_native_log2.h create mode 100644 libclc/clc/include/clc/math/clc_native_rsqrt.h create mode 100644 libclc/clc/include/clc/math/clc_native_sin.h create mode 100644 libclc/clc/include/clc/math/clc_native_sqrt.h delete mode 100644 libclc/clc/include/clc/math/unary_intrin.inc rename libclc/{amdgpu/lib/math/native_log10.cl => clc/lib/amdgpu/math/clc_native_exp.cl} (74%) rename libclc/{amdgpu/lib/math/native_exp.inc => clc/lib/amdgpu/math/clc_native_exp.inc} (75%) rename libclc/{amdgpu/lib/math/native_exp2.cl => clc/lib/amdgpu/math/clc_native_exp2.cl} (72%) rename libclc/{amdgpu/lib/math/native_log.cl => clc/lib/amdgpu/math/clc_native_log10.cl} (74%) rename libclc/{amdgpu/lib/math/native_log10.inc => clc/lib/amdgpu/math/clc_native_log10.inc} (73%) create mode 100644 libclc/clc/lib/generic/math/clc_native_cos.cl create mode 100644 libclc/clc/lib/generic/math/clc_native_exp.cl create mode 100644 libclc/clc/lib/generic/math/clc_native_exp2.cl create mode 100644 libclc/clc/lib/generic/math/clc_native_log.cl create mode 100644 libclc/clc/lib/generic/math/clc_native_log10.cl create mode 100644 libclc/clc/lib/generic/math/clc_native_log2.cl rename libclc/{amdgpu/lib/math/native_exp.cl => clc/lib/generic/math/clc_native_rsqrt.cl} (79%) rename libclc/{generic/lib/math/native_rsqrt.inc => clc/lib/generic/math/clc_native_rsqrt.inc} (76%) create mode 100644 libclc/clc/lib/generic/math/clc_native_sin.cl create mode 100644 libclc/clc/lib/generic/math/clc_native_sqrt.cl rename libclc/{r600/lib/math/native_rsqrt.cl => clc/lib/r600/math/clc_native_rsqrt.cl} (65%) delete mode 100644 libclc/generic/lib/math/native_unary_intrinsic.inc 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