diff --git a/libclc/amdgcn-amdhsa/libspirv/SOURCES b/libclc/amdgcn-amdhsa/libspirv/SOURCES index 1a6e3dfa7d3e7..2a0650facae28 100644 --- a/libclc/amdgcn-amdhsa/libspirv/SOURCES +++ b/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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/acos.cl b/libclc/amdgcn-amdhsa/libspirv/math/acos.cl new file mode 100644 index 0000000000000..a18feb52d332e --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/acosh.cl b/libclc/amdgcn-amdhsa/libspirv/math/acosh.cl new file mode 100644 index 0000000000000..7eb8c824ba58c --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/asin.cl b/libclc/amdgcn-amdhsa/libspirv/math/asin.cl new file mode 100644 index 0000000000000..365f47819b204 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/asinh.cl b/libclc/amdgcn-amdhsa/libspirv/math/asinh.cl new file mode 100644 index 0000000000000..8efcbc04bac71 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/atan2.cl b/libclc/amdgcn-amdhsa/libspirv/math/atan2.cl new file mode 100644 index 0000000000000..2a64cf8a80680 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/atanh.cl b/libclc/amdgcn-amdhsa/libspirv/math/atanh.cl new file mode 100644 index 0000000000000..91f34d2056361 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/ceil.cl b/libclc/amdgcn-amdhsa/libspirv/math/ceil.cl new file mode 100644 index 0000000000000..b9e36c26f8986 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/copysign.cl b/libclc/amdgcn-amdhsa/libspirv/math/copysign.cl new file mode 100644 index 0000000000000..473e86f9943a6 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/cos.cl b/libclc/amdgcn-amdhsa/libspirv/math/cos.cl index 2a9fd49724ce0..b93ad3dcc38d6 100644 --- a/libclc/amdgcn-amdhsa/libspirv/math/cos.cl +++ b/libclc/amdgcn-amdhsa/libspirv/math/cos.cl @@ -6,9 +6,15 @@ // //===----------------------------------------------------------------------===// +#include #include -_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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/cosh.cl b/libclc/amdgcn-amdhsa/libspirv/math/cosh.cl new file mode 100644 index 0000000000000..1ad0f0427c269 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/cospi.cl b/libclc/amdgcn-amdhsa/libspirv/math/cospi.cl new file mode 100644 index 0000000000000..e20fe100152d5 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/erf.cl b/libclc/amdgcn-amdhsa/libspirv/math/erf.cl new file mode 100644 index 0000000000000..a0b5686ed134c --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/erfc.cl b/libclc/amdgcn-amdhsa/libspirv/math/erfc.cl new file mode 100644 index 0000000000000..e3a97a8ad6efe --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/exp.cl b/libclc/amdgcn-amdhsa/libspirv/math/exp.cl new file mode 100644 index 0000000000000..9c324175cfa3c --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/exp10.cl b/libclc/amdgcn-amdhsa/libspirv/math/exp10.cl new file mode 100644 index 0000000000000..22c22d6b2c1d9 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/exp2.cl b/libclc/amdgcn-amdhsa/libspirv/math/exp2.cl new file mode 100644 index 0000000000000..2f173e32c6308 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/expm1.cl b/libclc/amdgcn-amdhsa/libspirv/math/expm1.cl new file mode 100644 index 0000000000000..5d6d5b240dff8 --- /dev/null +++ b/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 +#include + +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 diff --git a/libclc/amdgcn-amdhsa/libspirv/math/fabs.cl b/libclc/amdgcn-amdhsa/libspirv/math/fabs.cl new file mode 100644 index 0000000000000..193c4cadea4d9 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/fabs.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 +#include + +double __ocml_fabs_f64(double); +float __ocml_fabs_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_fabs +#define __CLC_BUILTIN __ocml_fabs +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/fdim.cl b/libclc/amdgcn-amdhsa/libspirv/math/fdim.cl new file mode 100644 index 0000000000000..11fbec26b65bb --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/fdim.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 +#include + +double __ocml_fdim_f64(double,double); +float __ocml_fdim_f32(float,float); + +#define __CLC_FUNCTION __spirv_ocl_fdim +#define __CLC_BUILTIN __ocml_fdim +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/floor.cl b/libclc/amdgcn-amdhsa/libspirv/math/floor.cl new file mode 100644 index 0000000000000..4d21cc7435898 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/floor.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 +#include + +double __ocml_floor_f64(double); +float __ocml_floor_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_floor +#define __CLC_BUILTIN __ocml_floor +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/fmax.cl b/libclc/amdgcn-amdhsa/libspirv/math/fmax.cl new file mode 100644 index 0000000000000..08e3878885be3 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/fmax.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 +#include + +double __ocml_fmax_f64(double,double); +float __ocml_fmax_f32(float,float); + +#define __CLC_FUNCTION __spirv_ocl_fmax +#define __CLC_BUILTIN __ocml_fmax +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/fmin.cl b/libclc/amdgcn-amdhsa/libspirv/math/fmin.cl new file mode 100644 index 0000000000000..6f60e65f190d0 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/fmin.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 +#include + +double __ocml_fmin_f64(double,double); +float __ocml_fmin_f32(float,float); + +#define __CLC_FUNCTION __spirv_ocl_fmin +#define __CLC_BUILTIN __ocml_fmin +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/fmod.cl b/libclc/amdgcn-amdhsa/libspirv/math/fmod.cl new file mode 100644 index 0000000000000..60ea7dac0278e --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/fmod.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 +#include + +double __ocml_fmod_f64(double,double); +float __ocml_fmod_f32(float,float); + +#define __CLC_FUNCTION __spirv_ocl_fmod +#define __CLC_BUILTIN __ocml_fmod +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/frexp.cl b/libclc/amdgcn-amdhsa/libspirv/math/frexp.cl new file mode 100644 index 0000000000000..eb0256822acf1 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/frexp.cl @@ -0,0 +1,86 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +double __ocml_frexp_f64(double, int *); +float __ocml_frexp_f32(float, int *); + +_CLC_OVERLOAD _CLC_DEF float __clc_spirv_ocl_frexp(float x, private int *ep) { + return __ocml_frexp_f32(x, ep); +} + +#ifdef cl_khr_fp64 + +#pragma OPENCL EXTENSION cl_khr_fp64 : enable + +_CLC_OVERLOAD _CLC_DEF double __clc_spirv_ocl_frexp(double x, private int *ep) { + return __ocml_frexp_f64(x, ep); +} + +#endif + +#ifdef cl_khr_fp16 + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +_CLC_OVERLOAD _CLC_DEF half __clc_spirv_ocl_frexp(half x, private int *ep) { + float t = x; + return __ocml_frexp_f32(t, ep); +} + +#endif + +#define __CLC_ADDRESS_SPACE private +#define __CLC_GENTYPE float +#include +#undef __CLC_GENTYPE +#ifdef cl_khr_fp64 +#define __CLC_GENTYPE double +#include +#undef __CLC_GENTYPE +#endif +#ifdef cl_khr_fp16 +#define __CLC_GENTYPE half +#include +#undef __CLC_GENTYPE +#endif +#undef __CLC_ADDRESS_SPACE + +#define __CLC_ADDRESS_SPACE global +#define __CLC_GENTYPE float +#include +#undef __CLC_GENTYPE +#ifdef cl_khr_fp64 +#define __CLC_GENTYPE double +#include +#undef __CLC_GENTYPE +#endif +#ifdef cl_khr_fp16 +#define __CLC_GENTYPE half +#include +#undef __CLC_GENTYPE +#endif +#undef __CLC_ADDRESS_SPACE + +#define __CLC_ADDRESS_SPACE local +#define __CLC_GENTYPE float +#include +#undef __CLC_GENTYPE +#ifdef cl_khr_fp64 +#define __CLC_GENTYPE double +#include +#undef __CLC_GENTYPE +#endif +#ifdef cl_khr_fp16 +#define __CLC_GENTYPE half +#include +#undef __CLC_GENTYPE +#endif +#undef __CLC_ADDRESS_SPACE diff --git a/libclc/amdgcn-amdhsa/libspirv/math/frexp.inc b/libclc/amdgcn-amdhsa/libspirv/math/frexp.inc new file mode 100644 index 0000000000000..7c8409040772a --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/frexp.inc @@ -0,0 +1,64 @@ +//===----------------------------------------------------------------------===// +// +// 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 +__spirv_ocl_frexp(__CLC_GENTYPE x, __CLC_ADDRESS_SPACE int *ep) { + int t_ep; + __CLC_GENTYPE res = __clc_spirv_ocl_frexp(x, &t_ep); + *ep = t_ep; + return res; +} + +#define __CLC_GENTYPE_VEC __CLC_XCONCAT(__CLC_GENTYPE, 2) + +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE_VEC +__spirv_ocl_frexp(__CLC_GENTYPE_VEC x, __CLC_ADDRESS_SPACE int2 *ep) { + int ep_x; + int ep_y; + __CLC_GENTYPE_VEC res = (__CLC_GENTYPE_VEC)(__spirv_ocl_frexp(x.x, &ep_x), + __spirv_ocl_frexp(x.y, &ep_y)); + *ep = (int2)(ep_x, ep_y); + return res; +} + +#undef __CLC_GENTYPE_VEC +#define __CLC_GENTYPE_VEC __CLC_XCONCAT(__CLC_GENTYPE, 3) + +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE_VEC +__spirv_ocl_frexp(__CLC_GENTYPE_VEC x, __CLC_ADDRESS_SPACE int3 *ep) { + int ep_x; + int ep_y; + int ep_z; + __CLC_GENTYPE_VEC res = (__CLC_GENTYPE_VEC)(__spirv_ocl_frexp(x.x, &ep_x), + __spirv_ocl_frexp(x.y, &ep_y), + __spirv_ocl_frexp(x.z, &ep_z)); + *ep = (int3)(ep_x, ep_y, ep_z); + return res; +} + +#define __CLC_VECTORIZE_FREXP_IMPL(GENTYPE, VEC_LEN, HALF_VEC_LEN) \ + _CLC_OVERLOAD _CLC_DEF GENTYPE __spirv_ocl_frexp( \ + GENTYPE x, __CLC_ADDRESS_SPACE int##VEC_LEN *ep) { \ + int##HALF_VEC_LEN ep_lo; \ + int##HALF_VEC_LEN ep_hi; \ + GENTYPE res = (GENTYPE)(__spirv_ocl_frexp(x.lo, &ep_lo), \ + __spirv_ocl_frexp(x.hi, &ep_hi)); \ + *ep = (int##VEC_LEN)(ep_lo, ep_hi); \ + return res; \ + } + +#define __CLC_VECTORIZE_FREXP(VEC_LEN, HALF_VEC_LEN) \ + __CLC_VECTORIZE_FREXP_IMPL(__CLC_XCONCAT(__CLC_GENTYPE, VEC_LEN), VEC_LEN, \ + HALF_VEC_LEN) + +__CLC_VECTORIZE_FREXP(4, 2) +__CLC_VECTORIZE_FREXP(8, 4) +__CLC_VECTORIZE_FREXP(16, 8) + +#undef __CLC_GENTYPE_VEC +#undef __CLC_VECTORIZE_FREXP diff --git a/libclc/amdgcn-amdhsa/libspirv/math/hypot.cl b/libclc/amdgcn-amdhsa/libspirv/math/hypot.cl new file mode 100644 index 0000000000000..53701fd7ca264 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/hypot.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 +#include + +double __ocml_hypot_f64(double,double); +float __ocml_hypot_f32(float,float); + +#define __CLC_FUNCTION __spirv_ocl_hypot +#define __CLC_BUILTIN __ocml_hypot +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/ldexp.cl b/libclc/amdgcn-amdhsa/libspirv/math/ldexp.cl new file mode 100644 index 0000000000000..2365ac6cf600b --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/ldexp.cl @@ -0,0 +1,35 @@ +//===----------------------------------------------------------------------===// +// +// 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 +#include +#include + +double __ocml_ldexp_f64(double, int); +float __ocml_ldexp_f32(float, int); + +#define __CLC_FUNCTION __spirv_ocl_ldexp +#define __CLC_BUILTIN __ocml_ldexp +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) + +_CLC_DEFINE_BINARY_BUILTIN(float, __spirv_ocl_ldexp, __ocml_ldexp_f32, float, int) +_CLC_DEFINE_BINARY_BUILTIN(float, __spirv_ocl_ldexp, __ocml_ldexp_f32, float, uint) + +#ifdef cl_khr_fp64 +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +_CLC_DEFINE_BINARY_BUILTIN(double, __spirv_ocl_ldexp, __ocml_ldexp_f64, double, int) +_CLC_DEFINE_BINARY_BUILTIN(double, __spirv_ocl_ldexp, __ocml_ldexp_f64, double, uint) +#endif + +#ifdef cl_khr_fp16 +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +_CLC_DEFINE_BINARY_BUILTIN(half, __spirv_ocl_ldexp, __ocml_ldexp_f32, half, int) +_CLC_DEFINE_BINARY_BUILTIN(half, __spirv_ocl_ldexp, __ocml_ldexp_f32, half, uint) +#endif diff --git a/libclc/amdgcn-amdhsa/libspirv/math/lgamma.cl b/libclc/amdgcn-amdhsa/libspirv/math/lgamma.cl new file mode 100644 index 0000000000000..003d7df0a11f2 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/lgamma.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 +#include + +double __ocml_lgamma_f64(double); +float __ocml_lgamma_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_lgamma +#define __CLC_BUILTIN __ocml_lgamma +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/log.cl b/libclc/amdgcn-amdhsa/libspirv/math/log.cl new file mode 100644 index 0000000000000..b2a7709bf8c35 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/log.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 +#include + +double __ocml_log_f64(double); +float __ocml_log_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_log +#define __CLC_BUILTIN __ocml_log +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/log10.cl b/libclc/amdgcn-amdhsa/libspirv/math/log10.cl new file mode 100644 index 0000000000000..f3fd6467062c5 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/log10.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 +#include + +double __ocml_log10_f64(double); +float __ocml_log10_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_log10 +#define __CLC_BUILTIN __ocml_log10 +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/log1p.cl b/libclc/amdgcn-amdhsa/libspirv/math/log1p.cl new file mode 100644 index 0000000000000..b42d1d93bca8b --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/log1p.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 +#include + +double __ocml_log1p_f64(double); +float __ocml_log1p_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_log1p +#define __CLC_BUILTIN __ocml_log1p +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/logb.cl b/libclc/amdgcn-amdhsa/libspirv/math/logb.cl new file mode 100644 index 0000000000000..32653df78c8a5 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/logb.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 +#include + +double __ocml_logb_f64(double); +float __ocml_logb_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_logb +#define __CLC_BUILTIN __ocml_logb +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/modf.cl b/libclc/amdgcn-amdhsa/libspirv/math/modf.cl new file mode 100644 index 0000000000000..daaf816418179 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/modf.cl @@ -0,0 +1,60 @@ +//===----------------------------------------------------------------------===// +// +// 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 + +double __ocml_modf_f64(double, double *); +float __ocml_modf_f32(float, float *); + +#define __CLC_MODF_IMPL(ADDRSPACE, BUILTIN, FP_TYPE, ARG_TYPE) \ + _CLC_OVERLOAD _CLC_DEF ARG_TYPE __spirv_ocl_modf(ARG_TYPE x, \ + ADDRSPACE ARG_TYPE *iptr) { \ + FP_TYPE stack_iptr; \ + ARG_TYPE ret = BUILTIN(x, &stack_iptr); \ + *iptr = stack_iptr; \ + return ret; \ + } + +#define __CLC_MODF(BUILTIN, FP_TYPE, ARG_TYPE) \ + __CLC_MODF_IMPL(private, BUILTIN, FP_TYPE, ARG_TYPE) \ + __CLC_MODF_IMPL(local, BUILTIN, FP_TYPE, ARG_TYPE) \ + __CLC_MODF_IMPL(global, BUILTIN, FP_TYPE, ARG_TYPE) + +__CLC_MODF(__ocml_modf_f32, float, float) + + +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __spirv_ocl_modf, float, + private, float) +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __spirv_ocl_modf, float, + local, float) +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __spirv_ocl_modf, float, + global, float) + +#ifdef cl_khr_fp64 +__CLC_MODF(__ocml_modf_f64, double, double) + +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __spirv_ocl_modf, double, + private, double) +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __spirv_ocl_modf, double, + local, double) +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __spirv_ocl_modf, double, + global, double) +#endif + +#ifdef cl_khr_fp16 +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +__CLC_MODF(__ocml_modf_f32, float, half) + +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __spirv_ocl_modf, half, + private, half) +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __spirv_ocl_modf, half, + local, half) +_CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __spirv_ocl_modf, half, + global, half) +#endif diff --git a/libclc/amdgcn-amdhsa/libspirv/math/nextafter.cl b/libclc/amdgcn-amdhsa/libspirv/math/nextafter.cl new file mode 100644 index 0000000000000..50a659d0d3035 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/nextafter.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 +#include + +double __ocml_nextafter_f64(double,double); +float __ocml_nextafter_f32(float,float); + +#define __CLC_FUNCTION __spirv_ocl_nextafter +#define __CLC_BUILTIN __ocml_nextafter +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/pow.cl b/libclc/amdgcn-amdhsa/libspirv/math/pow.cl new file mode 100644 index 0000000000000..e5db8a67cc416 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/pow.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 +#include + +double __ocml_pow_f64(double, double); +float __ocml_pow_f32(float, float); + +#define __CLC_FUNCTION __spirv_ocl_pow +#define __CLC_BUILTIN __ocml_pow +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/round.cl b/libclc/amdgcn-amdhsa/libspirv/math/round.cl new file mode 100644 index 0000000000000..77ef5c8c170a5 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/round.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 +#include + +double __ocml_round_f64(double); +float __ocml_round_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_round +#define __CLC_BUILTIN __ocml_round +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/rsqrt.cl b/libclc/amdgcn-amdhsa/libspirv/math/rsqrt.cl new file mode 100644 index 0000000000000..3d8189e98231f --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/rsqrt.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 +#include + +double __ocml_rsqrt_f64(double); +float __ocml_rsqrt_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_rsqrt +#define __CLC_BUILTIN __ocml_rsqrt +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/sin.cl b/libclc/amdgcn-amdhsa/libspirv/math/sin.cl index f976818f14caf..79e1dc75ab2e0 100644 --- a/libclc/amdgcn-amdhsa/libspirv/math/sin.cl +++ b/libclc/amdgcn-amdhsa/libspirv/math/sin.cl @@ -6,9 +6,14 @@ // //===----------------------------------------------------------------------===// +#include #include -_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t -__spirv_ocl_sin(__clc_fp32_t In) { - return __builtin_amdgcn_sinf(In); -} +double __ocml_sin_f64(double); +float __ocml_sin_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_sin +#define __CLC_BUILTIN __ocml_sin +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/sincos.cl b/libclc/amdgcn-amdhsa/libspirv/math/sincos.cl new file mode 100644 index 0000000000000..415f18eec4975 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/sincos.cl @@ -0,0 +1,60 @@ +//===----------------------------------------------------------------------===// +// +// 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 + + void __ocml_sincos_f64(double, double *, double *); + void __ocml_sincos_f32(float, float *, float * ); + + #define __CLC_SINCOS_IMPL(ADDRSPACE, BUILTIN, FP_TYPE, ARG_TYPE) \ + _CLC_OVERLOAD _CLC_DEF ARG_TYPE __spirv_ocl_sincos( \ + ARG_TYPE x, ADDRSPACE ARG_TYPE *cosval_ptr) { \ + FP_TYPE sinval; \ + FP_TYPE cosval; \ + BUILTIN(x, &sinval, &cosval); \ + *cosval_ptr = cosval; \ + return sinval; \ + } + + #define __CLC_SINCOS(BUILTIN, FP_TYPE, ARG_TYPE) \ + __CLC_SINCOS_IMPL(global, BUILTIN, FP_TYPE, ARG_TYPE) \ + __CLC_SINCOS_IMPL(local, BUILTIN, FP_TYPE, ARG_TYPE) \ + __CLC_SINCOS_IMPL(private, BUILTIN, FP_TYPE, ARG_TYPE) + + __CLC_SINCOS(__ocml_sincos_f32, float, float) + + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __spirv_ocl_sincos, float, + private, float) + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __spirv_ocl_sincos, float, + local, float) + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __spirv_ocl_sincos, float, + global, float) + + #ifdef cl_khr_fp64 + __CLC_SINCOS(__ocml_sincos_f64, double, double) + + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __spirv_ocl_sincos, + double, private, double) + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __spirv_ocl_sincos, + double, local, double) + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __spirv_ocl_sincos, + double, global, double) + #endif + + #ifdef cl_khr_fp16 + #pragma OPENCL EXTENSION cl_khr_fp16 : enable + __CLC_SINCOS(__ocml_sincos_f32, float, half) + + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __spirv_ocl_sincos, half, + private, half) + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __spirv_ocl_sincos, half, + local, half) + _CLC_V_V_VP_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __spirv_ocl_sincos, half, + global, half) + #endif diff --git a/libclc/amdgcn-amdhsa/libspirv/math/sinpi.cl b/libclc/amdgcn-amdhsa/libspirv/math/sinpi.cl new file mode 100644 index 0000000000000..918b43acb2c0f --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/sinpi.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 +#include + +double __ocml_sinpi_f64(double); +float __ocml_sinpi_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_sinpi +#define __CLC_BUILTIN __ocml_sinpi +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/sqrt.cl b/libclc/amdgcn-amdhsa/libspirv/math/sqrt.cl index d499e1de2a90a..8a91b898e360a 100644 --- a/libclc/amdgcn-amdhsa/libspirv/math/sqrt.cl +++ b/libclc/amdgcn-amdhsa/libspirv/math/sqrt.cl @@ -6,14 +6,14 @@ // //===----------------------------------------------------------------------===// +#include #include -_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t -__spirv_ocl_sqrt(__clc_fp32_t In) { - return __builtin_amdgcn_sqrtf(In); -} +double __ocml_sqrt_f64(double); +float __ocml_sqrt_f32(float); -_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp64_t -__spirv_ocl_sqrt(__clc_fp64_t In) { - return __builtin_amdgcn_sqrt(In); -} +#define __CLC_FUNCTION __spirv_ocl_sqrt +#define __CLC_BUILTIN __ocml_sqrt +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/tan.cl b/libclc/amdgcn-amdhsa/libspirv/math/tan.cl new file mode 100644 index 0000000000000..7324b43dff33b --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/tan.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 +#include + +double __ocml_tan_f64(double); +float __ocml_tan_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_tan +#define __CLC_BUILTIN __ocml_tan +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/tanh.cl b/libclc/amdgcn-amdhsa/libspirv/math/tanh.cl new file mode 100644 index 0000000000000..f9a9d87623cb3 --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/tanh.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 +#include + +double __ocml_tanh_f64(double); +float __ocml_tanh_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_tanh +#define __CLC_BUILTIN __ocml_tanh +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include diff --git a/libclc/amdgcn-amdhsa/libspirv/math/trunc.cl b/libclc/amdgcn-amdhsa/libspirv/math/trunc.cl new file mode 100644 index 0000000000000..06740ffaef28c --- /dev/null +++ b/libclc/amdgcn-amdhsa/libspirv/math/trunc.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 +#include + +double __ocml_trunc_f64(double); +float __ocml_trunc_f32(float); + +#define __CLC_FUNCTION __spirv_ocl_trunc +#define __CLC_BUILTIN __ocml_trunc +#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32) +#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64) +#include