Skip to content

Commit b4db951

Browse files
wenju-hewenju-he
committed
use __CLC_SCALAR, set __CLC_MIN_VECSIZE
1 parent 4e6aa2a commit b4db951

File tree

7 files changed

+40
-22
lines changed

7 files changed

+40
-22
lines changed

‎libclc/clc/include/clc/shared/binary_def_scalarize.inc

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,11 @@
88

99
#include <clc/utils.h>
1010

11-
#if __CLC_VECSIZE_OR_1 == 1
11+
#ifdef __CLC_SCALAR
12+
13+
#ifndef __CLC_MIN_VECSIZE
14+
#define __CLC_MIN_VECSIZE 2
15+
#endif
1216

1317
#ifndef __IMPL_FUNCTION
1418
#define __IMPL_FUNCTION FUNCTION
@@ -30,13 +34,14 @@
3034
#define __CLC_ARG2_TYPE __CLC_GENTYPE
3135
#endif
3236

33-
#ifdef __CLC_HAS_SCALAR
37+
#if __CLC_MIN_VECSIZE == 1
3438
_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE FUNCTION(__CLC_ARG1_TYPE x,
3539
__CLC_ARG2_TYPE y) {
3640
return __IMPL_FUNCTION(x, y);
3741
}
38-
#endif
42+
#endif // __CLC_MIN_VECSIZE == 1
3943

44+
#if __CLC_MIN_VECSIZE <= 2
4045
#define __CLC_RET_TYPE2 __CLC_XCONCAT(__CLC_RET_TYPE, 2)
4146
#define __CLC_ARG1_TYPE2 __CLC_XCONCAT(__CLC_ARG1_TYPE, 2)
4247
#define __CLC_ARG2_TYPE2 __CLC_XCONCAT(__CLC_ARG2_TYPE, 2)
@@ -48,6 +53,7 @@ _CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE2 FUNCTION(__CLC_ARG1_TYPE2 x,
4853
#undef __CLC_RET_TYPE2
4954
#undef __CLC_ARG1_TYPE2
5055
#undef __CLC_ARG2_TYPE2
56+
#endif // __CLC_MIN_VECSIZE <= 2
5157

5258
#define __CLC_RET_TYPE3 __CLC_XCONCAT(__CLC_RET_TYPE, 3)
5359
#define __CLC_ARG1_TYPE3 __CLC_XCONCAT(__CLC_ARG1_TYPE, 3)
@@ -118,4 +124,4 @@ _CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE16 FUNCTION(__CLC_ARG1_TYPE16 x,
118124
#undef __CLC_ARG1_TYPE16
119125
#undef __CLC_ARG2_TYPE16
120126

121-
#endif // __CLC_VECSIZE_OR_1 == 1
127+
#endif // __CLC_SCALAR

‎libclc/clc/include/clc/shared/ternary_def_scalarize.inc

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,11 @@
88

99
#include <clc/utils.h>
1010

11-
#if __CLC_VECSIZE_OR_1 == 1
11+
#ifdef __CLC_SCALAR
12+
13+
#ifndef __CLC_MIN_VECSIZE
14+
#define __CLC_MIN_VECSIZE 2
15+
#endif
1216

1317
#ifndef __IMPL_FUNCTION
1418
#define __IMPL_FUNCTION FUNCTION
@@ -34,14 +38,15 @@
3438
#define __CLC_ARG3_TYPE __CLC_GENTYPE
3539
#endif
3640

37-
#ifdef __CLC_HAS_SCALAR
41+
#if __CLC_MIN_VECSIZE == 1
3842
_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE FUNCTION(__CLC_ARG1_TYPE x,
3943
__CLC_ARG2_TYPE y,
4044
__CLC_ARG3_TYPE z) {
4145
return __IMPL_FUNCTION(x, y, z);
4246
}
43-
#endif
47+
#endif // __CLC_MIN_VECSIZE == 1
4448

49+
#if __CLC_MIN_VECSIZE <= 2
4550
#define __CLC_RET_TYPE2 __CLC_XCONCAT(__CLC_RET_TYPE, 2)
4651
#define __CLC_ARG1_TYPE2 __CLC_XCONCAT(__CLC_ARG1_TYPE, 2)
4752
#define __CLC_ARG2_TYPE2 __CLC_XCONCAT(__CLC_ARG2_TYPE, 2)
@@ -56,6 +61,7 @@ _CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE2 FUNCTION(__CLC_ARG1_TYPE2 x,
5661
#undef __CLC_ARG1_TYPE2
5762
#undef __CLC_ARG2_TYPE2
5863
#undef __CLC_ARG3_TYPE2
64+
#endif // __CLC_MIN_VECSIZE <= 2
5965

6066
#define __CLC_RET_TYPE3 __CLC_XCONCAT(__CLC_RET_TYPE, 3)
6167
#define __CLC_ARG1_TYPE3 __CLC_XCONCAT(__CLC_ARG1_TYPE, 3)
@@ -140,4 +146,4 @@ _CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE16 FUNCTION(__CLC_ARG1_TYPE16 x,
140146
#undef __CLC_ARG2_TYPE16
141147
#undef __CLC_ARG3_TYPE16
142148

143-
#endif // __CLC_VECSIZE_OR_1 == 1
149+
#endif // __CLC_SCALAR

‎libclc/clc/include/clc/shared/unary_def_scalarize.inc

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,11 @@
88

99
#include <clc/utils.h>
1010

11-
#if __CLC_VECSIZE_OR_1 == 1
11+
#ifdef __CLC_SCALAR
12+
13+
#ifndef __CLC_MIN_VECSIZE
14+
#define __CLC_MIN_VECSIZE 2
15+
#endif
1216

1317
#ifndef __IMPL_FUNCTION
1418
#define __IMPL_FUNCTION FUNCTION
@@ -26,19 +30,21 @@
2630
#define __CLC_ARG2_TYPE __CLC_GENTYPE
2731
#endif
2832

29-
#ifdef __CLC_HAS_SCALAR
33+
#if __CLC_MIN_VECSIZE == 1
3034
_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE FUNCTION(__CLC_ARG1_TYPE x) {
3135
return __IMPL_FUNCTION(x);
3236
}
33-
#endif
37+
#endif // __CLC_MIN_VECSIZE == 1
3438

39+
#if __CLC_MIN_VECSIZE <= 2
3540
#define __CLC_RET_TYPE2 __CLC_XCONCAT(__CLC_RET_TYPE, 2)
3641
#define __CLC_ARG1_TYPE2 __CLC_XCONCAT(__CLC_ARG1_TYPE, 2)
3742
_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE2 FUNCTION(__CLC_ARG1_TYPE2 x) {
3843
return (__CLC_RET_TYPE2)(__IMPL_FUNCTION(x.s0), __IMPL_FUNCTION(x.s1));
3944
}
4045
#undef __CLC_RET_TYPE2
4146
#undef __CLC_ARG1_TYPE2
47+
#endif // __CLC_MIN_VECSIZE <= 2
4248

4349
#define __CLC_RET_TYPE3 __CLC_XCONCAT(__CLC_RET_TYPE, 3)
4450
#define __CLC_ARG1_TYPE3 __CLC_XCONCAT(__CLC_ARG1_TYPE, 3)
@@ -84,4 +90,4 @@ _CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE16 FUNCTION(__CLC_ARG1_TYPE16 x) {
8490
#undef __CLC_RET_TYPE16
8591
#undef __CLC_ARG1_TYPE16
8692

87-
#endif // __CLC_VECSIZE_OR_1 == 1
93+
#endif // __CLC_SCALAR

‎libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@
1212

1313
#define FUNCTION __clc_ldexp
1414
#define __CLC_ARG2_TYPE int
15-
#define __CLC_HAS_SCALAR
15+
#define __CLC_MIN_VECSIZE 1
1616

1717
#ifdef __HAS_LDEXPF__
1818
// This defines all the ldexp(floatN, intN) variants.

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@
1010
#include <clc/internal/clc.h>
1111

1212
#define __FLOAT_ONLY
13-
#define __CLC_HAS_SCALAR
13+
#define __CLC_MIN_VECSIZE 1
1414
#define FUNCTION __clc_native_exp2
1515
#define __IMPL_FUNCTION __builtin_amdgcn_exp2f
1616
#define __CLC_BODY <clc/shared/unary_def_scalarize.inc>

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11,12 +11,12 @@
1111
#include <clc/relational/clc_isnan.h>
1212

1313
#define __FLOAT_ONLY
14-
#define __CLC_HAS_SCALAR
14+
#define __CLC_MIN_VECSIZE 1
1515
#define FUNCTION __clc_fmax
1616
#define __IMPL_FUNCTION __builtin_fmaxf
1717
#define __CLC_BODY <clc/shared/binary_def_scalarize.inc>
1818
#include <clc/math/gentype.inc>
19-
#undef __CLC_HAS_SCALAR
19+
#undef __CLC_MIN_VECSIZE
2020
#undef FUNCTION
2121
#undef __IMPL_FUNCTION
2222

@@ -25,12 +25,12 @@
2525
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
2626

2727
#define __DOUBLE_ONLY
28-
#define __CLC_HAS_SCALAR
28+
#define __CLC_MIN_VECSIZE 1
2929
#define FUNCTION __clc_fmax
3030
#define __IMPL_FUNCTION __builtin_fmax
3131
#define __CLC_BODY <clc/shared/binary_def_scalarize.inc>
3232
#include <clc/math/gentype.inc>
33-
#undef __CLC_HAS_SCALAR
33+
#undef __CLC_MIN_VECSIZE
3434
#undef FUNCTION
3535
#undef __IMPL_FUNCTION
3636

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11,12 +11,12 @@
1111
#include <clc/relational/clc_isnan.h>
1212

1313
#define __FLOAT_ONLY
14-
#define __CLC_HAS_SCALAR
14+
#define __CLC_MIN_VECSIZE 1
1515
#define FUNCTION __clc_fmin
1616
#define __IMPL_FUNCTION __builtin_fminf
1717
#define __CLC_BODY <clc/shared/binary_def_scalarize.inc>
1818
#include <clc/math/gentype.inc>
19-
#undef __CLC_HAS_SCALAR
19+
#undef __CLC_MIN_VECSIZE
2020
#undef FUNCTION
2121
#undef __IMPL_FUNCTION
2222

@@ -25,12 +25,12 @@
2525
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
2626

2727
#define __DOUBLE_ONLY
28-
#define __CLC_HAS_SCALAR
28+
#define __CLC_MIN_VECSIZE 1
2929
#define FUNCTION __clc_fmin
3030
#define __IMPL_FUNCTION __builtin_fmin
3131
#define __CLC_BODY <clc/shared/binary_def_scalarize.inc>
3232
#include <clc/math/gentype.inc>
33-
#undef __CLC_HAS_SCALAR
33+
#undef __CLC_MIN_VECSIZE
3434
#undef FUNCTION
3535
#undef __IMPL_FUNCTION
3636

0 commit comments

Comments
 (0)