Skip to content

[libspirv] Replace typo _CLC_DECL with _CLC_DEF for function definition #20093

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Sep 16, 2025
+85 −96
Merged

[libspirv] Replace typo _CLC_DECL with _CLC_DEF for function definition #20093

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4bb7d39
[libspirv] Replace typo _CLC_DECL with _CLC_DEF for function definition
wenju-he Sep 16, 2025
2605572
undo changes to clc
wenju-he Sep 16, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_add.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ AMDGPU_ATOMIC(__spirv_AtomicIAdd, long, __hip_atomic_fetch_add)
AMDGPU_ATOMIC(__spirv_AtomicIAdd, unsigned long, __hip_atomic_fetch_add)

#define AMDGPU_ATOMIC_FP32_ADD_IMPL(AS, CHECK, NEW_BUILTIN) \
_CLC_OVERLOAD _CLC_DECL float __spirv_AtomicFAddEXT( \
_CLC_OVERLOAD _CLC_DEF float __spirv_AtomicFAddEXT( \
AS float *p, int scope, int semantics, float val) { \
if (CHECK) \
return NEW_BUILTIN(p, val); \
Expand All @@ -40,7 +40,7 @@ AMDGPU_ATOMIC_FP32_ADD_IMPL(, AMDGPU_ARCH_BETWEEN(9400, 10000),
__builtin_amdgcn_flat_atomic_fadd_f32)

#define AMDGPU_ATOMIC_FP64_ADD_IMPL(AS, CHECK, NEW_BUILTIN) \
_CLC_OVERLOAD _CLC_DECL double __spirv_AtomicFAddEXT( \
_CLC_OVERLOAD _CLC_DEF double __spirv_AtomicFAddEXT( \
AS double *p, int scope, int semantics, double val) { \
if (CHECK) \
return NEW_BUILTIN(p, val); \
Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,8 +60,8 @@ extern int __oclc_amdgpu_reflect(__constant char *);
}

#define AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, AS, BUILTIN) \
_CLC_OVERLOAD _CLC_DECL TYPE FUNC_NAME(AS TYPE *p, int scope, int semantics, \
TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE FUNC_NAME(AS TYPE *p, int scope, int semantics, \
TYPE val) { \
int atomic_scope = 0, memory_order = 0; \
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
return BUILTIN(p, val, memory_order, atomic_scope); \
Expand Down
11 changes: 5 additions & 6 deletions libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,16 +20,15 @@ _CLC_DEF static bool __clc_amdgcn_is_global(generic void *ptr) {
}

#define GenericCastToPtrExplicit_To(ADDRSPACE, NAME) \
_CLC_DECL _CLC_OVERLOAD \
ADDRSPACE void *__spirv_GenericCastToPtrExplicit_To##NAME( \
generic void *ptr, int unused) { \
_CLC_OVERLOAD _CLC_DEF ADDRSPACE void * \
__spirv_GenericCastToPtrExplicit_To##NAME(generic void *ptr, int unused) { \
if (__clc_amdgcn_is_##ADDRSPACE(ptr)) \
return (ADDRSPACE void *)ptr; \
return 0; \
} \
_CLC_DECL _CLC_OVERLOAD \
ADDRSPACE const void *__spirv_GenericCastToPtrExplicit_To##NAME( \
generic const void *ptr, int unused) { \
_CLC_OVERLOAD _CLC_DEF ADDRSPACE const void * \
__spirv_GenericCastToPtrExplicit_To##NAME(generic const void *ptr, \
int unused) { \
return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr, \
unused); \
}
Expand Down
35 changes: 17 additions & 18 deletions libclc/libspirv/lib/generic/atomic/atomic_dec.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,46 +8,45 @@

#include <libspirv/spirv.h>

_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIDecrement(local int *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIDecrement(local int *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (int)1);
}

_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIDecrement(global int *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIDecrement(global int *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (int)1);
}

_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIDecrement(local uint *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIDecrement(local uint *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (uint)1);
}

_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIDecrement(global uint *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIDecrement(global uint *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (uint)1);
}

#ifdef cl_khr_int64_base_atomics
_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIDecrement(local long *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIDecrement(local long *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (long)1);
}

_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIDecrement(global long *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIDecrement(global long *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (long)1);
}

_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIDecrement(local ulong *p,
int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIDecrement(local ulong *p, int scope,
int semantics) {
return __sync_fetch_and_sub(p, (ulong)1);
}

_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIDecrement(global ulong *p,
int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIDecrement(global ulong *p,
int scope,
int semantics) {
return __sync_fetch_and_sub(p, (ulong)1);
}
#endif
35 changes: 17 additions & 18 deletions libclc/libspirv/lib/generic/atomic/atomic_inc.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,46 +8,45 @@

#include <libspirv/spirv.h>

_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIIncrement(local int *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIIncrement(local int *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (int)1);
}

_CLC_OVERLOAD _CLC_DECL int __spirv_AtomicIIncrement(global int *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF int __spirv_AtomicIIncrement(global int *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (int)1);
}

_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIIncrement(local uint *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIIncrement(local uint *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (uint)1);
}

_CLC_OVERLOAD _CLC_DECL uint __spirv_AtomicIIncrement(global uint *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF uint __spirv_AtomicIIncrement(global uint *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (uint)1);
}

#ifdef cl_khr_int64_base_atomics
_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIIncrement(local long *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIIncrement(local long *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (long)1);
}

_CLC_OVERLOAD _CLC_DECL long __spirv_AtomicIIncrement(global long *p, int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF long __spirv_AtomicIIncrement(global long *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (long)1);
}

_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIIncrement(local ulong *p,
int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIIncrement(local ulong *p, int scope,
int semantics) {
return __sync_fetch_and_add(p, (ulong)1);
}

_CLC_OVERLOAD _CLC_DECL ulong __spirv_AtomicIIncrement(global ulong *p,
int scope,
int semantics) {
_CLC_OVERLOAD _CLC_DEF ulong __spirv_AtomicIIncrement(global ulong *p,
int scope,
int semantics) {
return __sync_fetch_and_add(p, (ulong)1);
}
#endif
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/generic/atomic/atomic_load.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,8 +16,8 @@
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, acquire) \
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \
_CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicLoad(AS TYPE *p, int scope, \
int semantics) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicLoad(AS TYPE *p, int scope, \
int semantics) { \
if (semantics & Acquire) { \
return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_acquire(p); \
} \
Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/generic/atomic/atomic_max.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,8 +9,8 @@
#include <libspirv/spirv.h>

#define IMPL(TYPE, AS, NAME, PREFIX, SUFFIX) \
_CLC_OVERLOAD _CLC_DECL TYPE NAME(AS TYPE *p, int scope, int semantics, \
TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE NAME(AS TYPE *p, int scope, int semantics, \
TYPE val) { \
return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \
}

Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/generic/atomic/atomic_min.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,8 +9,8 @@
#include <libspirv/spirv.h>

#define IMPL(TYPE, AS, NAME, PREFIX, SUFFIX) \
_CLC_OVERLOAD _CLC_DECL TYPE NAME(AS TYPE *p, int scope, int semantics, \
TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE NAME(AS TYPE *p, int scope, int semantics, \
TYPE val) { \
return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \
}

Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/generic/atomic/atomic_or.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,8 +9,8 @@
#include <libspirv/spirv.h>

#define IMPL(TYPE, AS, FN_NAME) \
_CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicOr(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicOr(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
return FN_NAME(p, val); \
}

Expand Down
12 changes: 6 additions & 6 deletions libclc/libspirv/lib/generic/atomic/atomic_store.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,13 +8,13 @@

#include <libspirv/spirv.h>

_CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(global float *p, int scope,
int semantics, float val) {
_CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore(global float *p, int scope,
int semantics, float val) {
__spirv_AtomicStore((global uint *)p, scope, semantics, __clc_as_uint(val));
}

_CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(local float *p, int scope,
int semantics, float val) {
_CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore(local float *p, int scope,
int semantics, float val) {
__spirv_AtomicStore((local uint *)p, scope, semantics, __clc_as_uint(val));
}

Expand All @@ -26,8 +26,8 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(local float *p, int scope,
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, release) \
FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \
_CLC_OVERLOAD _CLC_DECL void __spirv_AtomicStore(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
_CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
if (semantics == Release) { \
__clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_release(p, val); \
} else if (semantics == SequentiallyConsistent) { \
Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/generic/atomic/atomic_sub.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,8 +9,8 @@
#include <libspirv/spirv.h>

#define IMPL(TYPE, AS, FN_NAME) \
_CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicISub(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicISub(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
return FN_NAME(p, val); \
}

Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/generic/atomic/atomic_xor.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,8 +9,8 @@
#include <libspirv/spirv.h>

#define IMPL(TYPE, AS, FN_NAME) \
_CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicXor(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicXor(AS TYPE *p, int scope, \
int semantics, TYPE val) { \
return FN_NAME(p, val); \
}

Expand Down
5 changes: 2 additions & 3 deletions libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_add.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,9 +19,8 @@ __CLC_NVVM_ATOMIC(float, float, f, add, __spirv_AtomicFAddEXT)
#ifdef cl_khr_int64_base_atomics

#define __CLC_NVVM_ATOMIC_ADD_DOUBLE_IMPL(ADDR_SPACE, ADDR_SPACE_NV) \
__attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL double \
__spirv_AtomicFAddEXT(ADDR_SPACE double *pointer, int scope, int semantics, \
double value) { \
_CLC_OVERLOAD _CLC_DEF double __spirv_AtomicFAddEXT( \
ADDR_SPACE double *pointer, int scope, int semantics, double value) { \
/* Semantics mask may include memory order, storage class and other info \
Memory order is stored in the lowest 5 bits */ \
unsigned int order = semantics & 0x1F; \
Expand Down
7 changes: 3 additions & 4 deletions libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -76,10 +76,9 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);
// Type V, Type C);
#define __CLC_NVVM_ATOMIC_CAS_IMPL(TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, \
OP_MANGLED, ADDR_SPACE, ADDR_SPACE_NV) \
__attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL TYPE \
__spirv_Atomic##OP_MANGLED(ADDR_SPACE TYPE *pointer, int scope, \
int semantics1, int semantics2, TYPE cmp, \
TYPE value) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_Atomic##OP_MANGLED( \
ADDR_SPACE TYPE *pointer, int scope, int semantics1, int semantics2, \
TYPE cmp, TYPE value) { \
/* Semantics mask may include memory order, storage class and other info \
Memory order is stored in the lowest 5 bits */ \
unsigned int order = semantics1 & 0x1F; \
Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,8 +85,8 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);

#define __CLC_NVVM_ATOMIC_IMPL(TYPE, TYPE_NV, TYPE_MANGLED_NV, OP, NAME, \
ADDR_SPACE, ADDR_SPACE_NV) \
__attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL TYPE NAME( \
ADDR_SPACE TYPE *pointer, int scope, int semantics, TYPE value) { \
_CLC_OVERLOAD _CLC_DEF TYPE NAME(ADDR_SPACE TYPE *pointer, int scope, \
int semantics, TYPE value) { \
/* Semantics mask may include memory order, storage class and other info \
Memory order is stored in the lowest 5 bits */ \
unsigned int order = semantics & 0x1F; \
Expand Down
5 changes: 2 additions & 3 deletions libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_inc_dec_helpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,9 +13,8 @@
#include <libspirv/spirv_types.h>

#define __CLC_NVVM_ATOMIC_INCDEC_IMPL(TYPE, OP_MANGLED, VAL, ADDR_SPACE) \
__attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL TYPE \
__spirv_Atomic##OP_MANGLED(ADDR_SPACE TYPE *pointer, int scope, \
int semantics) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_Atomic##OP_MANGLED( \
ADDR_SPACE TYPE *pointer, int scope, int semantics) { \
return __spirv_AtomicIAdd(pointer, scope, semantics, VAL); \
}

Expand Down
4 changes: 2 additions & 2 deletions libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_load.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,8 +38,8 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int);

#define __CLC_NVVM_ATOMIC_LOAD_IMPL(TYPE, TYPE_NV, TYPE_MANGLED_NV, \
ADDR_SPACE, ADDR_SPACE_NV) \
__attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL TYPE \
__spirv_AtomicLoad(ADDR_SPACE TYPE *pointer, int scope, int semantics) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_AtomicLoad(ADDR_SPACE TYPE *pointer, \
int scope, int semantics) { \
/* Semantics mask may include memory order, storage class and other info \
Memory order is stored in the lowest 5 bits */ \
unsigned int order = semantics & 0x1F; \
Expand Down
5 changes: 2 additions & 3 deletions libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_max.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,9 +20,8 @@ __CLC_NVVM_ATOMIC(unsigned long, unsigned long, ul, max, __spirv_AtomicUMax)
#undef __CLC_NVVM_ATOMIC_IMPL

#define __CLC_NVVM_ATOMIC_MAX_IMPL(TYPE, TYPE_INT, OP_MANGLED, ADDR_SPACE) \
__attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL TYPE \
__spirv_Atomic##OP_MANGLED(ADDR_SPACE TYPE *pointer, int scope, \
int semantics, TYPE val) { \
_CLC_OVERLOAD _CLC_DEF TYPE __spirv_Atomic##OP_MANGLED( \
ADDR_SPACE TYPE *pointer, int scope, int semantics, TYPE val) { \
int load_order; \
switch (semantics) { \
case SequentiallyConsistent: \
Expand Down
Loading