Skip to content

Commit 789ec8b

Browse files
GYDmedwinGYDmedwin
authored
[libclc] Add new function implementations in math. (#4864)
In the math package, there are only atan.cl, cbrt.cl, cos.cl, sin.cl, and sqrt.cl files initially. The other files in the math package I submitted this time are all newly implemented. We are using the SYCL-CTS test set for testing, one of which is the math_builtin_api test. Only when these files are newly added can it be compiled and passed.
1 parent a203fe7 commit 789ec8b

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

45 files changed

+1063
-18
lines changed

‎libclc/amdgcn-amdhsa/libspirv/SOURCES

Lines changed: 45 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,54 @@
11
atomic/loadstore_helpers.ll
22
cl_khr_int64_extended_atomics/minmax_helpers.ll
33
synchronization/barrier.cl
4+
math/acos.cl
5+
math/acosh.cl
6+
math/asin.cl
7+
math/asinh.cl
8+
math/atan.cl
9+
math/atan2.cl
10+
math/atanh.cl
11+
math/cbrt.cl
12+
math/ceil.cl
13+
math/copysign.cl
414
math/cos.cl
15+
math/cosh.cl
16+
math/cospi.cl
17+
math/erf.cl
18+
math/erfc.cl
19+
math/exp.cl
20+
math/exp10.cl
21+
math/exp2.cl
22+
math/expm1.cl
23+
math/fabs.cl
24+
math/fdim.cl
25+
math/floor.cl
26+
math/fma.cl
27+
math/fmax.cl
28+
math/fmin.cl
29+
math/fmod.cl
30+
math/frexp.cl
31+
math/hypot.cl
32+
math/ilogb.cl
33+
math/ldexp.cl
34+
math/lgamma.cl
35+
math/log.cl
36+
math/log10.cl
37+
math/log1p.cl
38+
math/logb.cl
39+
math/modf.cl
40+
math/nextafter.cl
41+
math/pow.cl
42+
math/round.cl
43+
math/rsqrt.cl
544
math/sin.cl
45+
math/sincos.cl
46+
math/sinh.cl
47+
math/sinpi.cl
648
math/sqrt.cl
7-
math/atan.cl
8-
math/cbrt.cl
49+
math/tan.cl
50+
math/tanh.cl
51+
math/trunc.cl
952
workitem/get_global_size.cl
1053
workitem/get_local_size.cl
1154
workitem/get_num_groups.cl

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_acos_f64(double);
13+
float __ocml_acos_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_acos
16+
#define __CLC_BUILTIN __ocml_acos
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_acosh_f64(double);
13+
float __ocml_acosh_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_acosh
16+
#define __CLC_BUILTIN __ocml_acosh
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_asin_f64(double);
13+
float __ocml_asin_f32(float);
14+
15+
16+
#define __CLC_FUNCTION __spirv_ocl_asin
17+
#define __CLC_BUILTIN __ocml_asin
18+
19+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
20+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
21+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_asinh_f64(double);
13+
float __ocml_asinh_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_asinh
16+
#define __CLC_BUILTIN __ocml_asinh
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_atan2_f64(double,double);
13+
float __ocml_atan2_f32(float,float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_atan2
16+
#define __CLC_BUILTIN __ocml_atan2
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/binary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_atanh_f64(double);
13+
float __ocml_atanh_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_atanh
16+
#define __CLC_BUILTIN __ocml_atanh
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_ceil_f64(double);
13+
float __ocml_ceil_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_ceil
16+
#define __CLC_BUILTIN __ocml_ceil
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_copysign_f64(double, double);
13+
float __ocml_copysign_f32(float, float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_copysign
16+
#define __CLC_BUILTIN __ocml_copysign
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/binary_builtin.inc>

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

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6,9 +6,15 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9+
#include <clcmacro.h>
910
#include <spirv/spirv.h>
1011

11-
_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
12-
__spirv_ocl_cos(__clc_fp32_t In) {
13-
return __builtin_amdgcn_cosf(In);
14-
}
12+
double __ocml_cos_f64(double);
13+
14+
float __ocml_cos_f32(float);
15+
16+
#define __CLC_FUNCTION __spirv_ocl_cos
17+
#define __CLC_BUILTIN __ocml_cos
18+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
19+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
20+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_cosh_f64(double);
13+
float __ocml_cosh_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_cosh
16+
#define __CLC_BUILTIN __ocml_cosh
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_cospi_f64(double);
13+
float __ocml_cospi_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_cospi
16+
#define __CLC_BUILTIN __ocml_cospi
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_erf_f64(double);
13+
float __ocml_erf_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_erf
16+
#define __CLC_BUILTIN __ocml_erf
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_erfc_f64(double);
13+
float __ocml_erfc_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_erfc
16+
#define __CLC_BUILTIN __ocml_erfc
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_exp_f64(double);
13+
float __ocml_exp_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_exp
16+
#define __CLC_BUILTIN __ocml_exp
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_exp10_f64(double);
13+
float __ocml_exp10_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_exp10
16+
#define __CLC_BUILTIN __ocml_exp10
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_exp2_f64(double);
13+
float __ocml_exp2_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_exp2
16+
#define __CLC_BUILTIN __ocml_exp2
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clcmacro.h>
10+
#include <spirv/spirv.h>
11+
12+
double __ocml_expm1_f64(double);
13+
float __ocml_expm1_f32(float);
14+
15+
#define __CLC_FUNCTION __spirv_ocl_expm1
16+
#define __CLC_BUILTIN __ocml_expm1
17+
#define __CLC_BUILTIN_F __CLC_XCONCAT(__CLC_BUILTIN, _f32)
18+
#define __CLC_BUILTIN_D __CLC_XCONCAT(__CLC_BUILTIN, _f64)
19+
#include <math/unary_builtin.inc>

0 commit comments

Comments
 (0)