[libclc] Add new function implementations in math.

libclc/amdgcn-amdhsa/libspirv/SOURCES

@@ -1,11 +1,54 @@
 atomic/loadstore_helpers.ll
 cl_khr_int64_extended_atomics/minmax_helpers.ll
 synchronization/barrier.cl
+math/acos.cl
+math/acosh.cl
+math/asin.cl
+math/asinh.cl
+math/atan.cl
+math/atan2.cl
+math/atanh.cl
+math/cbrt.cl
+math/ceil.cl
+math/copysign.cl
 math/cos.cl
+math/cosh.cl
+math/cospi.cl
+math/erf.cl
+math/erfc.cl
+math/exp.cl
+math/exp10.cl
+math/exp2.cl
+math/expm1.cl
+math/fabs.cl
+math/fdim.cl
+math/floor.cl
+math/fma.cl
+math/fmax.cl
+math/fmin.cl
+math/fmod.cl
+math/frexp.cl
+math/hypot.cl
+math/ilogb.cl
+math/ldexp.cl
+math/lgamma.cl
+math/log.cl
+math/log10.cl
+math/log1p.cl
+math/logb.cl
+math/modf.cl
+math/nextafter.cl
+math/pow.cl
+math/round.cl
+math/rsqrt.cl
 math/sin.cl
+math/sincos.cl
+math/sinh.cl
+math/sinpi.cl
 math/sqrt.cl
-math/atan.cl
-math/cbrt.cl
+math/tan.cl
+math/tanh.cl
+math/trunc.cl
 workitem/get_global_size.cl
 workitem/get_local_size.cl
 workitem/get_num_groups.cl

libclc/amdgcn-amdhsa/libspirv/math/acos.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_acos_f64(double);
+float __ocml_acos_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_acos
+#define __CLC_BUILTIN __ocml_acos
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/acosh.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_acosh_f64(double);
+float __ocml_acosh_f32(float); 
+
+#define __CLC_FUNCTION __spirv_ocl_acosh
+#define __CLC_BUILTIN __ocml_acosh
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/asin.cl

@@ -0,0 +1,21 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_asin_f64(double);
+float __ocml_asin_f32(float);
+
+
+#define __CLC_FUNCTION __spirv_ocl_asin
+#define __CLC_BUILTIN __ocml_asin
+
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/asinh.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_asinh_f64(double);
+float __ocml_asinh_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_asinh
+#define __CLC_BUILTIN __ocml_asinh
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/atan2.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_atan2_f64(double,double);
+float __ocml_atan2_f32(float,float);
+
+#define __CLC_FUNCTION __spirv_ocl_atan2
+#define __CLC_BUILTIN __ocml_atan2
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/binary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/atanh.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_atanh_f64(double);
+float __ocml_atanh_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_atanh
+#define __CLC_BUILTIN __ocml_atanh
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/ceil.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_ceil_f64(double);
+float __ocml_ceil_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_ceil
+#define __CLC_BUILTIN __ocml_ceil
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/copysign.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_copysign_f64(double, double);
+float __ocml_copysign_f32(float, float);
+
+#define __CLC_FUNCTION __spirv_ocl_copysign
+#define __CLC_BUILTIN __ocml_copysign
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/binary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/cos.cl

@@ -6,9 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clcmacro.h>
 #include <spirv/spirv.h>
 
-_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
-__spirv_ocl_cos(__clc_fp32_t In) {
-  return __builtin_amdgcn_cosf(In);
-}
+double __ocml_cos_f64(double);
+
+float __ocml_cos_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_cos
+#define __CLC_BUILTIN __ocml_cos
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/cosh.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_cosh_f64(double);
+float __ocml_cosh_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_cosh
+#define __CLC_BUILTIN __ocml_cosh
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/cospi.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_cospi_f64(double);
+float __ocml_cospi_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_cospi
+#define __CLC_BUILTIN __ocml_cospi
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/erf.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_erf_f64(double);
+float __ocml_erf_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_erf
+#define __CLC_BUILTIN __ocml_erf
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/erfc.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_erfc_f64(double);
+float __ocml_erfc_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_erfc
+#define __CLC_BUILTIN __ocml_erfc
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/exp.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_exp_f64(double);
+float __ocml_exp_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_exp
+#define __CLC_BUILTIN __ocml_exp
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/exp10.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_exp10_f64(double);
+float __ocml_exp10_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_exp10
+#define __CLC_BUILTIN __ocml_exp10
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/exp2.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_exp2_f64(double);
+float __ocml_exp2_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_exp2
+#define __CLC_BUILTIN __ocml_exp2
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>

libclc/amdgcn-amdhsa/libspirv/math/expm1.cl

@@ -0,0 +1,19 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clcmacro.h>
+#include <spirv/spirv.h>
+
+double __ocml_expm1_f64(double);
+float __ocml_expm1_f32(float);
+
+#define __CLC_FUNCTION __spirv_ocl_expm1
+#define __CLC_BUILTIN __ocml_expm1
+#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
+#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
+#include <math/unary_builtin.inc>