Skip to content

Commit 4bb7d39

Browse files
wenju-hewenju-he
committed
[libspirv] Replace typo _CLC_DECL with _CLC_DEF for function definition
Some built-ins miss alwaysinline attribute due to the typo.
1 parent e0ccfda commit 4bb7d39

26 files changed

+101
-112
lines changed

‎libclc/clc/include/clc/misc/shuffle2_def.inc‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -18,22 +18,22 @@
1818
// The return type is same base type as the input type, with the same vector
1919
// size as the mask. Elements in the mask must be the same size (number of bits)
2020
// as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask);
21-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
21+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
2222
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x,
2323
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) y, __CLC_U_GENTYPE mask) {
2424
return __IMPL_FUNCTION(FUNCTION)(x, y, mask);
2525
}
26-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
26+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
2727
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x,
2828
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) y, __CLC_U_GENTYPE mask) {
2929
return __IMPL_FUNCTION(FUNCTION)(x, y, mask);
3030
}
31-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
31+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
3232
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x,
3333
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) y, __CLC_U_GENTYPE mask) {
3434
return __IMPL_FUNCTION(FUNCTION)(x, y, mask);
3535
}
36-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
36+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
3737
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x,
3838
__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) y, __CLC_U_GENTYPE mask) {
3939
return __IMPL_FUNCTION(FUNCTION)(x, y, mask);

‎libclc/clc/include/clc/misc/shuffle_def.inc‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -18,19 +18,19 @@
1818
// The return type is same base type as the input type, with the same vector
1919
// size as the mask. Elements in the mask must be the same size (number of bits)
2020
// as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask);
21-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
21+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
2222
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_U_GENTYPE mask) {
2323
return __IMPL_FUNCTION(FUNCTION)(x, mask);
2424
}
25-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
25+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
2626
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_U_GENTYPE mask) {
2727
return __IMPL_FUNCTION(FUNCTION)(x, mask);
2828
}
29-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
29+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
3030
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_U_GENTYPE mask) {
3131
return __IMPL_FUNCTION(FUNCTION)(x, mask);
3232
}
33-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
33+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
3434
FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x, __CLC_U_GENTYPE mask) {
3535
return __IMPL_FUNCTION(FUNCTION)(x, mask);
3636
}

‎libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@
2424
#ifdef __CLC_FPSIZE
2525

2626
#define __CLC_DEFINE_ATOMIC(ADDRSPACE) \
27-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_atomic_compare_exchange( \
27+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_atomic_compare_exchange( \
2828
volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Comparator, \
2929
__CLC_GENTYPE Value, int MemoryOrderEqual, int MemoryOrderUnequal, \
3030
int MemoryScope) { \
@@ -38,7 +38,7 @@
3838
#else
3939

4040
#define __CLC_DEFINE_ATOMIC(ADDRSPACE) \
41-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_atomic_compare_exchange( \
41+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_atomic_compare_exchange( \
4242
volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Comparator, \
4343
__CLC_GENTYPE Value, int MemoryOrderEqual, int MemoryOrderUnequal, \
4444
int MemoryScope) { \

‎libclc/clc/lib/generic/atomic/clc_atomic_def.inc‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -31,15 +31,15 @@
3131

3232
#ifdef __CLC_NO_VALUE_ARG
3333
#define __CLC_DEFINE_ATOMIC(ADDRSPACE) \
34-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \
34+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \
3535
volatile ADDRSPACE __CLC_GENTYPE *Ptr, int MemoryOrder, \
3636
int MemoryScope) { \
3737
return __CLC_AS_RETTYPE(__IMPL_FUNCTION( \
3838
(ADDRSPACE __CLC_PTR_CASTTYPE *)Ptr, MemoryOrder, MemoryScope)); \
3939
}
4040
#elif defined(__CLC_INC_DEC)
4141
#define __CLC_DEFINE_ATOMIC(ADDRSPACE) \
42-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \
42+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \
4343
volatile ADDRSPACE __CLC_GENTYPE *Ptr, int MemoryOrder, \
4444
int MemoryScope) { \
4545
return __CLC_AS_RETTYPE( \
@@ -48,15 +48,15 @@
4848
}
4949
#elif defined(__CLC_RETURN_VOID)
5050
#define __CLC_DEFINE_ATOMIC(ADDRSPACE) \
51-
_CLC_OVERLOAD _CLC_DECL void FUNCTION(volatile ADDRSPACE __CLC_GENTYPE *Ptr, \
52-
__CLC_GENTYPE Value, int MemoryOrder, \
53-
int MemoryScope) { \
51+
_CLC_OVERLOAD _CLC_DEF void FUNCTION(volatile ADDRSPACE __CLC_GENTYPE *Ptr, \
52+
__CLC_GENTYPE Value, int MemoryOrder, \
53+
int MemoryScope) { \
5454
__IMPL_FUNCTION((ADDRSPACE __CLC_PTR_CASTTYPE *)Ptr, Value, MemoryOrder, \
5555
MemoryScope); \
5656
}
5757
#else
5858
#define __CLC_DEFINE_ATOMIC(ADDRSPACE) \
59-
_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \
59+
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \
6060
volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Value, \
6161
int MemoryOrder, int MemoryScope) { \
6262
return __CLC_AS_RETTYPE( \

‎libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_add.cl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@ AMDGPU_ATOMIC(__spirv_AtomicIAdd, long, __hip_atomic_fetch_add)
1818
AMDGPU_ATOMIC(__spirv_AtomicIAdd, unsigned long, __hip_atomic_fetch_add)
1919

2020
#define AMDGPU_ATOMIC_FP32_ADD_IMPL(AS, CHECK, NEW_BUILTIN) \
21-
_CLC_OVERLOAD _CLC_DECL float __spirv_AtomicFAddEXT( \
21+
_CLC_OVERLOAD _CLC_DEF float __spirv_AtomicFAddEXT( \
2222
AS float *p, int scope, int semantics, float val) { \
2323
if (CHECK) \
2424
return NEW_BUILTIN(p, val); \
@@ -40,7 +40,7 @@ AMDGPU_ATOMIC_FP32_ADD_IMPL(, AMDGPU_ARCH_BETWEEN(9400, 10000),
4040
__builtin_amdgcn_flat_atomic_fadd_f32)
4141

4242
#define AMDGPU_ATOMIC_FP64_ADD_IMPL(AS, CHECK, NEW_BUILTIN) \
43-
_CLC_OVERLOAD _CLC_DECL double __spirv_AtomicFAddEXT( \
43+
_CLC_OVERLOAD _CLC_DEF double __spirv_AtomicFAddEXT( \
4444
AS double *p, int scope, int semantics, double val) { \
4545
if (CHECK) \
4646
return NEW_BUILTIN(p, val); \

‎libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_helpers.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -60,8 +60,8 @@ extern int __oclc_amdgpu_reflect(__constant char *);
6060
}
6161

6262
#define AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, AS, BUILTIN) \
63-
_CLC_OVERLOAD _CLC_DECL TYPE FUNC_NAME(AS TYPE *p, int scope, int semantics, \
64-
TYPE val) { \
63+
_CLC_OVERLOAD _CLC_DEF TYPE FUNC_NAME(AS TYPE *p, int scope, int semantics, \
64+
TYPE val) { \
6565
int atomic_scope = 0, memory_order = 0; \
6666
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
6767
return BUILTIN(p, val, memory_order, atomic_scope); \

‎libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -20,16 +20,15 @@ _CLC_DEF static bool __clc_amdgcn_is_global(generic void *ptr) {
2020
}
2121

2222
#define GenericCastToPtrExplicit_To(ADDRSPACE, NAME) \
23-
_CLC_DECL _CLC_OVERLOAD \
24-
ADDRSPACE void *__spirv_GenericCastToPtrExplicit_To##NAME( \
25-
generic void *ptr, int unused) { \
23+
_CLC_OVERLOAD _CLC_DEF ADDRSPACE void * \
24+
__spirv_GenericCastToPtrExplicit_To##NAME(generic void *ptr, int unused) { \
2625
if (__clc_amdgcn_is_##ADDRSPACE(ptr)) \
2726
return (ADDRSPACE void *)ptr; \
2827
return 0; \
2928
} \
30-
_CLC_DECL _CLC_OVERLOAD \
31-
ADDRSPACE const void *__spirv_GenericCastToPtrExplicit_To##NAME( \
32-
generic const void *ptr, int unused) { \
29+
_CLC_OVERLOAD _CLC_DEF ADDRSPACE const void * \
30+
__spirv_GenericCastToPtrExplicit_To##NAME(generic const void *ptr, \
31+
int unused) { \
3332
return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr, \
3433
unused); \
3534
}

‎libclc/libspirv/lib/generic/atomic/atomic_dec.cl‎

Lines changed: 17 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -8,46 +8,45 @@
88

99
#include <libspirv/spirv.h>
1010

11-
_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIDecrement(local int *p, int scope,
12-
int semantics) {
11+
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIDecrement(local int *p, int scope,
12+
int semantics) {
1313
return __sync_fetch_and_sub(p, (int)1);
1414
}
1515

16-
_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIDecrement(global int *p, int scope,
17-
int semantics) {
16+
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIDecrement(global int *p, int scope,
17+
int semantics) {
1818
return __sync_fetch_and_sub(p, (int)1);
1919
}
2020

21-
_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIDecrement(local uint *p, int scope,
22-
int semantics) {
21+
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIDecrement(local uint *p, int scope,
22+
int semantics) {
2323
return __sync_fetch_and_sub(p, (uint)1);
2424
}
2525

26-
_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIDecrement(global uint *p, int scope,
27-
int semantics) {
26+
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIDecrement(global uint *p, int scope,
27+
int semantics) {
2828
return __sync_fetch_and_sub(p, (uint)1);
2929
}
3030

3131
#ifdef cl_khr_int64_base_atomics
32-
_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIDecrement(local long *p, int scope,
33-
int semantics) {
32+
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIDecrement(local long *p, int scope,
33+
int semantics) {
3434
return __sync_fetch_and_sub(p, (long)1);
3535
}
3636

37-
_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIDecrement(global long *p, int scope,
38-
int semantics) {
37+
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIDecrement(global long *p, int scope,
38+
int semantics) {
3939
return __sync_fetch_and_sub(p, (long)1);
4040
}
4141

42-
_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIDecrement(local ulong *p,
43-
int scope,
44-
int semantics) {
42+
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIDecrement(local ulong *p, int scope,
43+
int semantics) {
4544
return __sync_fetch_and_sub(p, (ulong)1);
4645
}
4746

48-
_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIDecrement(global ulong *p,
49-
int scope,
50-
int semantics) {
47+
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIDecrement(global ulong *p,
48+
int scope,
49+
int semantics) {
5150
return __sync_fetch_and_sub(p, (ulong)1);
5251
}
5352
#endif

‎libclc/libspirv/lib/generic/atomic/atomic_inc.cl‎

Lines changed: 17 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -8,46 +8,45 @@
88

99
#include <libspirv/spirv.h>
1010

11-
_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIIncrement(local int *p, int scope,
12-
int semantics) {
11+
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIIncrement(local int *p, int scope,
12+
int semantics) {
1313
return __sync_fetch_and_add(p, (int)1);
1414
}
1515

16-
_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIIncrement(global int *p, int scope,
17-
int semantics) {
16+
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIIncrement(global int *p, int scope,
17+
int semantics) {
1818
return __sync_fetch_and_add(p, (int)1);
1919
}
2020

21-
_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIIncrement(local uint *p, int scope,
22-
int semantics) {
21+
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIIncrement(local uint *p, int scope,
22+
int semantics) {
2323
return __sync_fetch_and_add(p, (uint)1);
2424
}
2525

26-
_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIIncrement(global uint *p, int scope,
27-
int semantics) {
26+
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIIncrement(global uint *p, int scope,
27+
int semantics) {
2828
return __sync_fetch_and_add(p, (uint)1);
2929
}
3030

3131
#ifdef cl_khr_int64_base_atomics
32-
_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIIncrement(local long *p, int scope,
33-
int semantics) {
32+
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIIncrement(local long *p, int scope,
33+
int semantics) {
3434
return __sync_fetch_and_add(p, (long)1);
3535
}
3636

37-
_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIIncrement(global long *p, int scope,
38-
int semantics) {
37+
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIIncrement(global long *p, int scope,
38+
int semantics) {
3939
return __sync_fetch_and_add(p, (long)1);
4040
}
4141

42-
_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIIncrement(local ulong *p,
43-
int scope,
44-
int semantics) {
42+
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIIncrement(local ulong *p, int scope,
43+
int semantics) {
4544
return __sync_fetch_and_add(p, (ulong)1);
4645
}
4746

48-
_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIIncrement(global ulong *p,
49-
int scope,
50-
int semantics) {
47+
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIIncrement(global ulong *p,
48+
int scope,
49+
int semantics) {
5150
return __sync_fetch_and_add(p, (ulong)1);
5251
}
5352
#endif

‎libclc/libspirv/lib/generic/atomic/atomic_load.cl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,8 @@
1616
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \
1717
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, acquire) \
1818
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \
19-
_CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicLoad(AS TYPE *p, int scope, \
20-
int semantics) { \
19+
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicLoad(AS TYPE *p, int scope, \
20+
int semantics) { \
2121
if (semantics & Acquire) { \
2222
return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_acquire(p); \
2323
} \

0 commit comments

Comments
 (0)