Author: Fraser Cormack Date: 2025-04-01T09:20:54+01:00 New Revision: ad48fffb53003333456b327a2a3f22bd81a77a00
URL: https://github.com/llvm/llvm-project/commit/ad48fffb53003333456b327a2a3f22bd81a77a00 DIFF: https://github.com/llvm/llvm-project/commit/ad48fffb53003333456b327a2a3f22bd81a77a00.diff LOG: [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. Added: libclc/clc/include/clc/math/clc_native_cos.h libclc/clc/include/clc/math/clc_native_exp.h libclc/clc/include/clc/math/clc_native_exp2.h libclc/clc/include/clc/math/clc_native_log.h libclc/clc/include/clc/math/clc_native_log10.h libclc/clc/include/clc/math/clc_native_log2.h libclc/clc/include/clc/math/clc_native_rsqrt.h libclc/clc/include/clc/math/clc_native_sin.h libclc/clc/include/clc/math/clc_native_sqrt.h libclc/clc/lib/amdgpu/math/clc_native_exp.cl libclc/clc/lib/amdgpu/math/clc_native_exp.inc libclc/clc/lib/amdgpu/math/clc_native_exp2.cl libclc/clc/lib/amdgpu/math/clc_native_log10.cl libclc/clc/lib/amdgpu/math/clc_native_log10.inc libclc/clc/lib/generic/math/clc_native_cos.cl libclc/clc/lib/generic/math/clc_native_exp.cl libclc/clc/lib/generic/math/clc_native_exp2.cl libclc/clc/lib/generic/math/clc_native_log.cl libclc/clc/lib/generic/math/clc_native_log10.cl libclc/clc/lib/generic/math/clc_native_log2.cl libclc/clc/lib/generic/math/clc_native_rsqrt.cl libclc/clc/lib/generic/math/clc_native_rsqrt.inc libclc/clc/lib/generic/math/clc_native_sin.cl libclc/clc/lib/generic/math/clc_native_sqrt.cl libclc/clc/lib/r600/math/clc_native_rsqrt.cl Modified: libclc/CMakeLists.txt libclc/amdgpu/lib/SOURCES libclc/clc/lib/amdgpu/SOURCES libclc/clc/lib/generic/SOURCES libclc/clc/lib/r600/SOURCES libclc/generic/lib/math/native_cos.cl libclc/generic/lib/math/native_exp.cl libclc/generic/lib/math/native_exp2.cl libclc/generic/lib/math/native_log.cl libclc/generic/lib/math/native_log10.cl libclc/generic/lib/math/native_log2.cl libclc/generic/lib/math/native_rsqrt.cl libclc/generic/lib/math/native_sin.cl libclc/generic/lib/math/native_sqrt.cl libclc/r600/lib/SOURCES Removed: libclc/amdgpu/lib/math/native_exp.cl libclc/amdgpu/lib/math/native_exp.inc libclc/amdgpu/lib/math/native_exp2.cl libclc/amdgpu/lib/math/native_log.cl libclc/amdgpu/lib/math/native_log.inc libclc/amdgpu/lib/math/native_log10.cl libclc/amdgpu/lib/math/native_log10.inc libclc/clc/include/clc/math/unary_intrin.inc libclc/generic/lib/math/native_rsqrt.inc libclc/generic/lib/math/native_unary_intrinsic.inc libclc/r600/lib/math/native_rsqrt.cl ################################################################################ diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 59a70a200c95c..efe7f5804e8fb 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 ed5e45a37c18d..ab5da40711aa4 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 820e4929f02cf..0000000000000 --- 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 0000000000000..8a580f13ad2aa --- /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 <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__ 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 0000000000000..48a1be616ea3e --- /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 <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__ 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 0000000000000..bc0b32d6212b5 --- /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 <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__ 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 0000000000000..ea0362503f670 --- /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 <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__ 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 0000000000000..c5cceeeba2952 --- /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 <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__ 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 0000000000000..25375970cedc0 --- /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 <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__ 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 0000000000000..59fd2134107db --- /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 <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__ 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 0000000000000..878e19882ae6b --- /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 <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__ 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 0000000000000..afff77c89b22f --- /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 <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__ 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 8028470114b8e..0000000000000 --- 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 <clc/clcfunc.h> -#include <clc/clctypes.h> - -_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 fd64a862021e8..31e07b608c4cc 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 7cbe1f98988d5..591ecb0ac00b5 100644 --- a/libclc/amdgpu/lib/math/native_log10.cl +++ b/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> 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 d7dbd888d2988..cf2d48c4054ed 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 39ae914b19634..76b1850fce574 100644 --- a/libclc/amdgpu/lib/math/native_exp2.cl +++ b/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) 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 ce3258bb5b37b..0668a635d24d7 100644 --- a/libclc/amdgpu/lib/math/native_log.cl +++ b/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> 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 2aef43a8ed1a5..c91d698609793 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 6a1ba95362220..c31963c59e950 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 0000000000000..de56fdec48d24 --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_cos +#define __CLC_FUNCTION(x) __builtin_elementwise_cos +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 0000000000000..400270a6163a4 --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_exp +#define __CLC_FUNCTION(x) __builtin_elementwise_exp +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 0000000000000..427d901fcdb19 --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_exp2 +#define __CLC_FUNCTION(x) __builtin_elementwise_exp2 +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 0000000000000..85f188b654282 --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_log +#define __CLC_FUNCTION(x) __builtin_elementwise_log +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 0000000000000..624018e4481bf --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_log10 +#define __CLC_FUNCTION(x) __builtin_elementwise_log10 +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 0000000000000..2c8c18e61ca5d --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_log2 +#define __CLC_FUNCTION(x) __builtin_elementwise_log2 +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 e62b79d4ec9fa..d5e6fcdae491f 100644 --- a/libclc/amdgpu/lib/math/native_exp.cl +++ b/libclc/clc/lib/generic/math/clc_native_rsqrt.cl @@ -6,8 +6,9 @@ // //===----------------------------------------------------------------------===// -#include <clc/clc.h> +#include <clc/internal/clc.h> +#include <clc/math/clc_native_sqrt.h> -#define __CLC_BODY <native_exp.inc> +#define __CLC_BODY <clc_native_rsqrt.inc> #define __FLOAT_ONLY #include <clc/math/gentype.inc> 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 058209bcb8a15..7a3b0b2af2721 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 0000000000000..22b988bf4375f --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_sin +#define __CLC_FUNCTION(x) __builtin_elementwise_sin +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> 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 0000000000000..ed022ef1fee1f --- /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 <clc/internal/clc.h> + +#define __FLOAT_ONLY +#define FUNCTION __clc_native_sqrt +#define __CLC_FUNCTION(x) __builtin_elementwise_sqrt +#define __CLC_BODY <clc/shared/unary_def.inc> + +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/r600/SOURCES b/libclc/clc/lib/r600/SOURCES index 8f66107e0454e..8d5caf167aa4e 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 e916147c3e057..ee09814eb1e76 100644 --- a/libclc/r600/lib/math/native_rsqrt.cl +++ b/libclc/clc/lib/r600/math/clc_native_rsqrt.cl @@ -6,12 +6,11 @@ // //===----------------------------------------------------------------------===// -#include <clc/clc.h> #include <clc/clcmacro.h> +#include <clc/internal/clc.h> -_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 d81a067f11044..0f4d46c1d73e4 100644 --- a/libclc/generic/lib/math/native_cos.cl +++ b/libclc/generic/lib/math/native_cos.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_cos.h> -#define __CLC_NATIVE_INTRINSIC cos - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_cos +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_exp.cl b/libclc/generic/lib/math/native_exp.cl index 8f4531343415c..c1d08ec2c7e1f 100644 --- a/libclc/generic/lib/math/native_exp.cl +++ b/libclc/generic/lib/math/native_exp.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_exp.h> -#define __CLC_NATIVE_INTRINSIC exp - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_exp +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_exp2.cl b/libclc/generic/lib/math/native_exp2.cl index ecde4a6761d22..ceb570733b974 100644 --- a/libclc/generic/lib/math/native_exp2.cl +++ b/libclc/generic/lib/math/native_exp2.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_exp2.h> -#define __CLC_NATIVE_INTRINSIC exp2 - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_exp2 +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_log.cl b/libclc/generic/lib/math/native_log.cl index 5731e09d21c9c..adc2ff495f8b5 100644 --- a/libclc/generic/lib/math/native_log.cl +++ b/libclc/generic/lib/math/native_log.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_log.h> -#define __CLC_NATIVE_INTRINSIC log - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_log +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_log10.cl b/libclc/generic/lib/math/native_log10.cl index eab7a6f14d035..f63292124f3b7 100644 --- a/libclc/generic/lib/math/native_log10.cl +++ b/libclc/generic/lib/math/native_log10.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_log10.h> -#define __CLC_NATIVE_INTRINSIC log10 - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_log10 +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_log2.cl b/libclc/generic/lib/math/native_log2.cl index 0db4be0b5e083..6b079872b1e0a 100644 --- a/libclc/generic/lib/math/native_log2.cl +++ b/libclc/generic/lib/math/native_log2.cl @@ -7,8 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_log2.h> -#define __CLC_NATIVE_INTRINSIC log2 -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_log2 +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_rsqrt.cl b/libclc/generic/lib/math/native_rsqrt.cl index 14430c04fb72d..cb49b2d1d6706 100644 --- a/libclc/generic/lib/math/native_rsqrt.cl +++ b/libclc/generic/lib/math/native_rsqrt.cl @@ -7,7 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_rsqrt.h> -#define __CLC_BODY <native_rsqrt.inc> #define __FLOAT_ONLY +#define FUNCTION native_rsqrt +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_sin.cl b/libclc/generic/lib/math/native_sin.cl index 0e2ced09fa2dd..50265b3936272 100644 --- a/libclc/generic/lib/math/native_sin.cl +++ b/libclc/generic/lib/math/native_sin.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_sin.h> -#define __CLC_NATIVE_INTRINSIC sin - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_sin +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> diff --git a/libclc/generic/lib/math/native_sqrt.cl b/libclc/generic/lib/math/native_sqrt.cl index 1b668e5976ef7..4cd022e8bbeba 100644 --- a/libclc/generic/lib/math/native_sqrt.cl +++ b/libclc/generic/lib/math/native_sqrt.cl @@ -7,9 +7,10 @@ //===----------------------------------------------------------------------===// #include <clc/clc.h> +#include <clc/math/clc_native_sqrt.h> -#define __CLC_NATIVE_INTRINSIC sqrt - -#define __CLC_BODY <native_unary_intrinsic.inc> #define __FLOAT_ONLY +#define FUNCTION native_sqrt +#define __CLC_BODY <clc/shared/unary_def.inc> + #include <clc/math/gentype.inc> 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 c118ec095692f..0000000000000 --- 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 <clc/utils.h> - -#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 <clc/math/unary_intrin.inc> - -#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 cad45785dc483..4342ac38201c1 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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
