Skip to content

Commit d2d1b58

Browse files
frasercrmckfrasercrmck
authored
[libclc] Move clcmacro.h to CLC library. NFC (#114845)
1 parent 8d023b7 commit d2d1b58

Some content is hidden

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

95 files changed

+220
-220
lines changed

‎libclc/amdgcn/lib/integer/popcount.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
#include <clc/clc.h>
2-
#include <utils.h>
2+
#include <clc/utils.h>
33
#include <integer/popcount.h>
44

55
#define __CLC_BODY "popcount.inc"

‎libclc/amdgcn/lib/math/fmax.cl‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,5 @@
11
#include <clc/clc.h>
2-
3-
#include "../../../generic/lib/clcmacro.h"
2+
#include <clc/clcmacro.h>
43

54
_CLC_DEF _CLC_OVERLOAD float fmax(float x, float y)
65
{

‎libclc/amdgcn/lib/math/fmin.cl‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,5 @@
11
#include <clc/clc.h>
2-
3-
#include "../../../generic/lib/clcmacro.h"
2+
#include <clc/clcmacro.h>
43

54
_CLC_DEF _CLC_OVERLOAD float fmin(float x, float y)
65
{

‎libclc/amdgcn/lib/math/ldexp.cl‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,8 +21,7 @@
2121
*/
2222

2323
#include <clc/clc.h>
24-
25-
#include "../../../generic/lib/clcmacro.h"
24+
#include <clc/clcmacro.h>
2625

2726
#ifdef __HAS_LDEXPF__
2827
#define BUILTINF __builtin_amdgcn_ldexpf

‎libclc/amdgpu/lib/math/half_native_unary.inc‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
#include <utils.h>
1+
#include <clc/utils.h>
22

33
#define __CLC_HALF_FUNC(x) __CLC_CONCAT(half_, x)
44
#define __CLC_NATIVE_FUNC(x) __CLC_CONCAT(native_, x)

‎libclc/amdgpu/lib/math/nextafter.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
#include <clc/clc.h>
2-
#include "../lib/clcmacro.h"
2+
#include <clc/clcmacro.h>
33
#include <math/clc_nextafter.h>
44

55
_CLC_DEFINE_BINARY_BUILTIN(float, nextafter, __clc_nextafter, float, float)

‎libclc/amdgpu/lib/math/sqrt.cl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,9 +20,9 @@
2020
* THE SOFTWARE.
2121
*/
2222

23-
#include <clc/clc.h>
24-
#include "../../../generic/lib/clcmacro.h"
2523
#include "math/clc_sqrt.h"
24+
#include <clc/clc.h>
25+
#include <clc/clcmacro.h>
2626

2727
_CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)
2828

‎libclc/generic/lib/clcmacro.h‎ renamed to ‎libclc/clc/include/clc/clcmacro.h‎

Lines changed: 116 additions & 98 deletions
Original file line numberDiff line numberDiff line change
@@ -1,92 +1,105 @@
1-
#include <clc/clc.h>
2-
#include <utils.h>
3-
4-
#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \
5-
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \
6-
return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y)); \
7-
} \
8-
\
9-
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
10-
return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
11-
} \
12-
\
13-
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
14-
return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
15-
} \
16-
\
17-
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
18-
return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
19-
} \
20-
\
21-
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
22-
return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
1+
#ifndef __CLC_CLCMACRO_H__
2+
#define __CLC_CLCMACRO_H__
3+
4+
#include <clc/internal/clc.h>
5+
#include <clc/utils.h>
6+
7+
#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \
8+
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \
9+
return (RET_TYPE##2)(FUNCTION(x.x), FUNCTION(x.y)); \
10+
} \
11+
\
12+
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
13+
return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
14+
} \
15+
\
16+
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
17+
return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
18+
} \
19+
\
20+
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
21+
return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
22+
} \
23+
\
24+
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
25+
return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
2326
}
2427

25-
#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
26-
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) { \
27-
return (RET_TYPE##2)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y)); \
28-
} \
29-
\
30-
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
31-
return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
32-
FUNCTION(x.z, y.z)); \
33-
} \
34-
\
35-
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
36-
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
37-
} \
38-
\
39-
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
40-
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
41-
} \
42-
\
43-
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
44-
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
28+
#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
29+
ARG2_TYPE) \
30+
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) { \
31+
return (RET_TYPE##2)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y)); \
32+
} \
33+
\
34+
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
35+
return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
36+
FUNCTION(x.z, y.z)); \
37+
} \
38+
\
39+
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
40+
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
41+
} \
42+
\
43+
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
44+
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
45+
} \
46+
\
47+
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
48+
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
4549
}
4650

47-
#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
48-
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
49-
return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
50-
} \
51-
\
52-
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
53-
return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
54-
FUNCTION(x, y.z)); \
55-
} \
56-
\
57-
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
58-
return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
59-
} \
60-
\
61-
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
62-
return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
63-
} \
64-
\
65-
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
66-
return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
67-
} \
68-
\
69-
70-
#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE, ARG3_TYPE) \
71-
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y, ARG3_TYPE##2 z) { \
72-
return (RET_TYPE##2)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
73-
} \
74-
\
75-
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, ARG3_TYPE##3 z) { \
76-
return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
77-
FUNCTION(x.z, y.z, z.z)); \
78-
} \
79-
\
80-
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, ARG3_TYPE##4 z) { \
81-
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
82-
} \
83-
\
84-
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, ARG3_TYPE##8 z) { \
85-
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
86-
} \
87-
\
88-
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, ARG3_TYPE##16 z) { \
89-
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
51+
#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
52+
ARG2_TYPE) \
53+
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
54+
return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
55+
} \
56+
\
57+
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
58+
return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
59+
FUNCTION(x, y.z)); \
60+
} \
61+
\
62+
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
63+
return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
64+
} \
65+
\
66+
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
67+
return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
68+
} \
69+
\
70+
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
71+
return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
72+
}
73+
74+
#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
75+
ARG2_TYPE, ARG3_TYPE) \
76+
DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y, \
77+
ARG3_TYPE##2 z) { \
78+
return (RET_TYPE##2)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
79+
} \
80+
\
81+
DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, \
82+
ARG3_TYPE##3 z) { \
83+
return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
84+
FUNCTION(x.z, y.z, z.z)); \
85+
} \
86+
\
87+
DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, \
88+
ARG3_TYPE##4 z) { \
89+
return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), \
90+
FUNCTION(x.hi, y.hi, z.hi)); \
91+
} \
92+
\
93+
DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, \
94+
ARG3_TYPE##8 z) { \
95+
return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), \
96+
FUNCTION(x.hi, y.hi, z.hi)); \
97+
} \
98+
\
99+
DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, \
100+
ARG3_TYPE##16 z) { \
101+
return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), \
102+
FUNCTION(x.hi, y.hi, z.hi)); \
90103
}
91104

92105
#define _CLC_V_S_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
@@ -161,21 +174,24 @@
161174
ARG2_TYPE, 8) *)((ADDR_SPACE ARG2_TYPE *)y + 8))); \
162175
}
163176

164-
#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
165-
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) { \
166-
return BUILTIN(x, y); \
167-
} \
168-
_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
177+
#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, \
178+
ARG2_TYPE) \
179+
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) { \
180+
return BUILTIN(x, y); \
181+
} \
182+
_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, \
183+
ARG2_TYPE)
169184

170-
#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
171-
_CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
172-
_CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
185+
#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG( \
186+
RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
187+
_CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, \
188+
ARG2_TYPE) \
189+
_CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, \
190+
FUNCTION, ARG1_TYPE, ARG2_TYPE)
173191

174-
#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
175-
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { \
176-
return BUILTIN(x); \
177-
} \
178-
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
192+
#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
193+
_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { return BUILTIN(x); } \
194+
_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
179195

180196
#ifdef cl_khr_fp16
181197

@@ -199,3 +215,5 @@ _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
199215
#define _CLC_DEFINE_BINARY_BUILTIN_FP16(FUNCTION)
200216

201217
#endif
218+
219+
#endif // __CLC_CLCMACRO_H__

‎libclc/clspv/lib/math/fma.cl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,9 +24,9 @@
2424
// (__clc_sw_fma), but avoids the use of ulong in favor of uint2. The logic has
2525
// been updated as appropriate.
2626

27-
#include <clc/clc.h>
28-
#include "../../../generic/lib/clcmacro.h"
2927
#include "../../../generic/lib/math/math.h"
28+
#include <clc/clc.h>
29+
#include <clc/clcmacro.h>
3030

3131
struct fp {
3232
uint2 mantissa;

‎libclc/generic/include/math/clc_sqrt.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,6 @@
1+
#include <clc/clcfunc.h>
2+
#include <clc/clctypes.h>
3+
14
#define __CLC_FUNCTION __clc_sqrt
25
#define __CLC_BODY <clc/math/unary_decl.inc>
36
#include <clc/math/gentype.inc>

0 commit comments

Comments
 (0)