[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

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

libclc/amdgpu/lib/SOURCES

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

libclc/amdgpu/lib/math/native_log.inc

@@ -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);
-}

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_COS_H__
+#define __CLC_MATH_CLC_NATIVE_COS_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_cos
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_COS_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_EXP_H__
+#define __CLC_MATH_CLC_NATIVE_EXP_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_exp
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_EXP_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_EXP2_H__
+#define __CLC_MATH_CLC_NATIVE_EXP2_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_exp2
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_EXP2_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_LOG_H__
+#define __CLC_MATH_CLC_NATIVE_LOG_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_LOG10_H__
+#define __CLC_MATH_CLC_NATIVE_LOG10_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log10
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG10_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_LOG2_H__
+#define __CLC_MATH_CLC_NATIVE_LOG2_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log2
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG2_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_RSQRT_H__
+#define __CLC_MATH_CLC_NATIVE_RSQRT_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_rsqrt
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_RSQRT_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_SIN_H__
+#define __CLC_MATH_CLC_NATIVE_SIN_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_sin
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_SIN_H__

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

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MATH_CLC_NATIVE_SQRT_H__
+#define __CLC_MATH_CLC_NATIVE_SQRT_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_sqrt
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_SQRT_H__

libclc/clc/include/clc/math/unary_intrin.inc

@@ -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

libclc/clc/lib/amdgpu/SOURCES

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

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

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

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

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

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

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

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

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

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

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

libclc/clc/lib/generic/SOURCES

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

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>

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>

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>

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>

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>

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>

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>

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);
 }

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>

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>

libclc/clc/lib/r600/SOURCES

@@ -1 +1,2 @@
+math/clc_native_rsqrt.cl
 math/clc_rsqrt_override.cl

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);

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>

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>

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>

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>

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>

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>

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>

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>

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>

libclc/generic/lib/math/native_unary_intrinsic.inc

@@ -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

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