Skip to content

Commit aabdf57

Browse files
frasercrmckfrasercrmck
committed
[libclc] Move several 'native' builtins to CLC library
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. 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.
1 parent 87602f6 commit aabdf57

36 files changed

+369
-110
lines changed

‎libclc/amdgpu/lib/SOURCES

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,3 @@
1-
math/native_exp.cl
2-
math/native_exp2.cl
3-
math/native_log.cl
4-
math/native_log10.cl
51
math/half_exp.cl
62
math/half_exp10.cl
73
math/half_exp2.cl

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_COS_H__
10+
#define __CLC_MATH_CLC_NATIVE_COS_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_cos
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_COS_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_EXP_H__
10+
#define __CLC_MATH_CLC_NATIVE_EXP_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_exp
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_EXP_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_EXP2_H__
10+
#define __CLC_MATH_CLC_NATIVE_EXP2_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_exp2
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_EXP2_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_LOG_H__
10+
#define __CLC_MATH_CLC_NATIVE_LOG_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_log
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_LOG_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_LOG10_H__
10+
#define __CLC_MATH_CLC_NATIVE_LOG10_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_log10
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_LOG10_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_LOG2_H__
10+
#define __CLC_MATH_CLC_NATIVE_LOG2_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_log2
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_LOG2_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_SIN_H__
10+
#define __CLC_MATH_CLC_NATIVE_SIN_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_sin
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_SIN_H__

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
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+
#ifndef __CLC_MATH_CLC_NATIVE_SQRT_H__
10+
#define __CLC_MATH_CLC_NATIVE_SQRT_H__
11+
12+
#define __FLOAT_ONLY
13+
#define __CLC_FUNCTION __clc_native_sqrt
14+
#define __CLC_BODY <clc/shared/unary_decl.inc>
15+
16+
#include <clc/math/gentype.inc>
17+
18+
#undef __CLC_BODY
19+
#undef __CLC_FUNCTION
20+
#undef __FLOAT_ONLY
21+
22+
#endif // __CLC_MATH_CLC_NATIVE_SQRT_H__

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

Lines changed: 0 additions & 42 deletions
This file was deleted.

‎libclc/clc/lib/amdgpu/SOURCES

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1 +1,5 @@
1+
math/clc_native_exp2.cl
2+
math/clc_native_exp.cl
3+
math/clc_native_log10.cl
4+
math/clc_native_log.cl
15
math/clc_sqrt_fp64.cl

‎libclc/amdgpu/lib/math/native_exp.cl renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_exp.cl

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

9-
#include <clc/clc.h>
9+
#include <clc/float/definitions.h>
10+
#include <clc/internal/clc.h>
11+
#include <clc/math/clc_native_exp2.h>
1012

11-
#define __CLC_BODY <native_exp.inc>
13+
#define __CLC_BODY <clc_native_exp.inc>
1214
#define __FLOAT_ONLY
1315
#include <clc/math/gentype.inc>

‎libclc/amdgpu/lib/math/native_exp.inc renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_exp.inc

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

9-
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_exp(__CLC_GENTYPE val) {
10-
return native_exp2(val * M_LOG2E_F);
9+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_exp(__CLC_GENTYPE val) {
10+
return __clc_native_exp2(val * M_LOG2E_F);
1111
}

‎libclc/amdgpu/lib/math/native_exp2.cl renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_exp2.cl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6,11 +6,11 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9-
#include <clc/clc.h>
109
#include <clc/clcmacro.h>
10+
#include <clc/internal/clc.h>
1111

12-
_CLC_OVERLOAD _CLC_DEF float native_exp2(float val) {
12+
_CLC_OVERLOAD _CLC_DEF float __clc_native_exp2(float val) {
1313
return __builtin_amdgcn_exp2f(val);
1414
}
1515

16-
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, native_exp2, float)
16+
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_native_exp2, float)

‎libclc/amdgpu/lib/math/native_log10.cl renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_log.cl

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

9-
#include <clc/clc.h>
9+
#include <clc/float/definitions.h>
10+
#include <clc/internal/clc.h>
11+
#include <clc/math/clc_native_log2.h>
1012

11-
#define __CLC_BODY <native_log10.inc>
13+
#define __CLC_BODY <clc_native_log.inc>
1214
#define __FLOAT_ONLY
1315
#include <clc/math/gentype.inc>

‎libclc/amdgpu/lib/math/native_log10.inc renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_log.inc

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

9-
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_log10(__CLC_GENTYPE val) {
10-
return native_log2(val) * (M_LN2_F / M_LN10_F);
9+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_log(__CLC_GENTYPE val) {
10+
return __clc_native_log2(val) * (1.0f / M_LOG2E_F);
1111
}

‎libclc/amdgpu/lib/math/native_log.cl renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_log10.cl

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

9-
#include <clc/clc.h>
9+
#include <clc/float/definitions.h>
10+
#include <clc/internal/clc.h>
11+
#include <clc/math/clc_native_log2.h>
1012

11-
#define __CLC_BODY <native_log.inc>
13+
#define __CLC_BODY <clc_native_log10.inc>
1214
#define __FLOAT_ONLY
1315
#include <clc/math/gentype.inc>

‎libclc/amdgpu/lib/math/native_log.inc renamed to ‎libclc/clc/lib/amdgpu/math/clc_native_log10.inc

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

9-
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_log(__CLC_GENTYPE val) {
10-
return native_log2(val) * (1.0f / M_LOG2E_F);
9+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_log10(__CLC_GENTYPE val) {
10+
return __clc_native_log2(val) * (M_LN2_F / M_LN10_F);
1111
}

‎libclc/clc/lib/generic/SOURCES

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,14 @@ math/clc_log2.cl
4545
math/clc_mad.cl
4646
math/clc_modf.cl
4747
math/clc_nan.cl
48+
math/clc_native_cos.cl
49+
math/clc_native_exp.cl
50+
math/clc_native_exp2.cl
51+
math/clc_native_log.cl
52+
math/clc_native_log10.cl
53+
math/clc_native_log2.cl
54+
math/clc_native_sin.cl
55+
math/clc_native_sqrt.cl
4856
math/clc_nextafter.cl
4957
math/clc_pow.cl
5058
math/clc_pown.cl

‎libclc/clc/lib/generic/math/clc_native_cos.cl

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
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 <clc/internal/clc.h>
10+
11+
#define __FLOAT_ONLY
12+
#define FUNCTION __clc_native_cos
13+
#define __CLC_FUNCTION(x) __builtin_elementwise_cos
14+
#define __CLC_BODY <clc/shared/unary_def.inc>
15+
16+
#include <clc/math/gentype.inc>

‎libclc/clc/lib/generic/math/clc_native_exp.cl

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
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 <clc/internal/clc.h>
10+
11+
#define __FLOAT_ONLY
12+
#define FUNCTION __clc_native_exp
13+
#define __CLC_FUNCTION(x) __builtin_elementwise_exp
14+
#define __CLC_BODY <clc/shared/unary_def.inc>
15+
16+
#include <clc/math/gentype.inc>

‎libclc/clc/lib/generic/math/clc_native_exp2.cl

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
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 <clc/internal/clc.h>
10+
11+
#define __FLOAT_ONLY
12+
#define FUNCTION __clc_native_exp2
13+
#define __CLC_FUNCTION(x) __builtin_elementwise_exp2
14+
#define __CLC_BODY <clc/shared/unary_def.inc>
15+
16+
#include <clc/math/gentype.inc>

0 commit comments

Comments
 (0)