Skip to content

Commit 6ec3aed

Browse files
wenju-hewenju-he
authored
[libspirv] Replace typo _CLC_DECL with _CLC_DEF for function definition (#20093)
Some built-ins miss alwaysinline attribute due to the typo.
1 parent e0ccfda commit 6ec3aed

22 files changed

+85
-96
lines changed

‎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
} \

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@
99
#include <libspirv/spirv.h>
1010

1111
#define IMPL(TYPE, AS, NAME, PREFIX, SUFFIX) \
12-
_CLC_OVERLOAD _CLC_DECL TYPE NAME(AS TYPE *p, int scope, int semantics, \
13-
TYPE val) { \
12+
_CLC_OVERLOAD _CLC_DEF TYPE NAME(AS TYPE *p, int scope, int semantics, \
13+
TYPE val) { \
1414
return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \
1515
}
1616

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@
99
#include <libspirv/spirv.h>
1010

1111
#define IMPL(TYPE, AS, NAME, PREFIX, SUFFIX) \
12-
_CLC_OVERLOAD _CLC_DECL TYPE NAME(AS TYPE *p, int scope, int semantics, \
13-
TYPE val) { \
12+
_CLC_OVERLOAD _CLC_DEF TYPE NAME(AS TYPE *p, int scope, int semantics, \
13+
TYPE val) { \
1414
return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \
1515
}
1616

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@
99
#include <libspirv/spirv.h>
1010

1111
#define IMPL(TYPE, AS, FN_NAME) \
12-
_CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicOr(AS TYPE *p, int scope, \
13-
int semantics, TYPE val) { \
12+
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicOr(AS TYPE *p, int scope, \
13+
int semantics, TYPE val) { \
1414
return FN_NAME(p, val); \
1515
}
1616

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

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -8,13 +8,13 @@
88

99
#include <libspirv/spirv.h>
1010

11-
_CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(global float *p, int scope,
12-
int semantics, float val) {
11+
_CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore(global float *p, int scope,
12+
int semantics, float val) {
1313
__spirv_AtomicStore((global uint *)p, scope, semantics, __clc_as_uint(val));
1414
}
1515

16-
_CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(local float *p, int scope,
17-
int semantics, float val) {
16+
_CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore(local float *p, int scope,
17+
int semantics, float val) {
1818
__spirv_AtomicStore((local uint *)p, scope, semantics, __clc_as_uint(val));
1919
}
2020

@@ -26,8 +26,8 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(local float *p, int scope,
2626
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \
2727
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, release) \
2828
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \
29-
_CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(AS TYPE *p, int scope, \
30-
int semantics, TYPE val) { \
29+
_CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore(AS TYPE *p, int scope, \
30+
int semantics, TYPE val) { \
3131
if (semantics == Release) { \
3232
__clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_release(p, val); \
3333
} else if (semantics == SequentiallyConsistent) { \

0 commit comments

Comments
 (0)