From 4bb7d39a0488cb351983d704842df0f1170df87b Mon Sep 17 00:00:00 2001 From: Wenju He Date: Tue, 16 Sep 2025 07:49:36 +0200 Subject: [PATCH 1/2] [libspirv] Replace typo _CLC_DECL with _CLC_DEF for function definition Some built-ins miss alwaysinline attribute due to the typo. --- libclc/clc/include/clc/misc/shuffle2_def.inc | 8 ++--- libclc/clc/include/clc/misc/shuffle_def.inc | 8 ++--- .../atomic/clc_atomic_compare_exchange.inc | 4 +-- .../clc/lib/generic/atomic/clc_atomic_def.inc | 12 +++---- .../lib/amdgcn-amdhsa/atomic/atomic_add.cl | 4 +-- .../lib/amdgcn-amdhsa/atomic/atomic_helpers.h | 4 +-- .../conversion/GenericCastToPtrExplicit.cl | 11 +++--- .../libspirv/lib/generic/atomic/atomic_dec.cl | 35 +++++++++---------- .../libspirv/lib/generic/atomic/atomic_inc.cl | 35 +++++++++---------- .../lib/generic/atomic/atomic_load.cl | 4 +-- .../libspirv/lib/generic/atomic/atomic_max.cl | 4 +-- .../libspirv/lib/generic/atomic/atomic_min.cl | 4 +-- .../libspirv/lib/generic/atomic/atomic_or.cl | 4 +-- .../lib/generic/atomic/atomic_store.cl | 12 +++---- .../libspirv/lib/generic/atomic/atomic_sub.cl | 4 +-- .../libspirv/lib/generic/atomic/atomic_xor.cl | 4 +-- .../lib/ptx-nvidiacl/atomic/atomic_add.cl | 5 ++- .../lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl | 7 ++-- .../lib/ptx-nvidiacl/atomic/atomic_helpers.h | 4 +-- .../atomic/atomic_inc_dec_helpers.h | 5 ++- .../lib/ptx-nvidiacl/atomic/atomic_load.cl | 4 +-- .../lib/ptx-nvidiacl/atomic/atomic_max.cl | 5 ++- .../lib/ptx-nvidiacl/atomic/atomic_min.cl | 5 ++- .../lib/ptx-nvidiacl/atomic/atomic_store.cl | 5 ++- .../lib/ptx-nvidiacl/atomic/atomic_sub.cl | 5 ++- .../conversion/GenericCastToPtrExplicit.cl | 11 +++--- 26 files changed, 101 insertions(+), 112 deletions(-) diff --git a/libclc/clc/include/clc/misc/shuffle2_def.inc b/libclc/clc/include/clc/misc/shuffle2_def.inc index 0415b8c187291..960419992ebf2 100644 --- a/libclc/clc/include/clc/misc/shuffle2_def.inc +++ b/libclc/clc/include/clc/misc/shuffle2_def.inc @@ -18,22 +18,22 @@ // The return type is same base type as the input type, with the same vector // size as the mask. Elements in the mask must be the same size (number of bits) // as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask); -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); } -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); } -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); } -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); diff --git a/libclc/clc/include/clc/misc/shuffle_def.inc b/libclc/clc/include/clc/misc/shuffle_def.inc index d2a088498fc6c..29fcef15e34b8 100644 --- a/libclc/clc/include/clc/misc/shuffle_def.inc +++ b/libclc/clc/include/clc/misc/shuffle_def.inc @@ -18,19 +18,19 @@ // The return type is same base type as the input type, with the same vector // size as the mask. Elements in the mask must be the same size (number of bits) // as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask); -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } -_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } diff --git a/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc b/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc index 32ff9b45b769e..74284fd61024c 100644 --- a/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc +++ b/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc @@ -24,7 +24,7 @@ #ifdef __CLC_FPSIZE #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_atomic_compare_exchange( \ + _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_atomic_compare_exchange( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Comparator, \ __CLC_GENTYPE Value, int MemoryOrderEqual, int MemoryOrderUnequal, \ int MemoryScope) { \ @@ -38,7 +38,7 @@ #else #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_atomic_compare_exchange( \ + _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_atomic_compare_exchange( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Comparator, \ __CLC_GENTYPE Value, int MemoryOrderEqual, int MemoryOrderUnequal, \ int MemoryScope) { \ diff --git a/libclc/clc/lib/generic/atomic/clc_atomic_def.inc b/libclc/clc/lib/generic/atomic/clc_atomic_def.inc index 2c45f49f60848..635532c60db93 100644 --- a/libclc/clc/lib/generic/atomic/clc_atomic_def.inc +++ b/libclc/clc/lib/generic/atomic/clc_atomic_def.inc @@ -31,7 +31,7 @@ #ifdef __CLC_NO_VALUE_ARG #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \ + _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, int MemoryOrder, \ int MemoryScope) { \ return __CLC_AS_RETTYPE(__IMPL_FUNCTION( \ @@ -39,7 +39,7 @@ } #elif defined(__CLC_INC_DEC) #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \ + _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, int MemoryOrder, \ int MemoryScope) { \ return __CLC_AS_RETTYPE( \ @@ -48,15 +48,15 @@ } #elif defined(__CLC_RETURN_VOID) #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DECL void FUNCTION(volatile ADDRSPACE __CLC_GENTYPE *Ptr, \ - __CLC_GENTYPE Value, int MemoryOrder, \ - int MemoryScope) { \ + _CLC_OVERLOAD _CLC_DEF void FUNCTION(volatile ADDRSPACE __CLC_GENTYPE *Ptr, \ + __CLC_GENTYPE Value, int MemoryOrder, \ + int MemoryScope) { \ __IMPL_FUNCTION((ADDRSPACE __CLC_PTR_CASTTYPE *)Ptr, Value, MemoryOrder, \ MemoryScope); \ } #else #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \ + _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Value, \ int MemoryOrder, int MemoryScope) { \ return __CLC_AS_RETTYPE( \ diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_add.cl b/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_add.cl index be59bec58fa77..e0202b1aaeb5b 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_add.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_add.cl @@ -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); \ @@ -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); \ diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_helpers.h b/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_helpers.h index 5ca852ce5523c..811ba65e2fc8f 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_helpers.h +++ b/libclc/libspirv/lib/amdgcn-amdhsa/atomic/atomic_helpers.h @@ -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); \ diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl b/libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl index 374ed03e47851..636966f117cc9 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl @@ -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); \ } diff --git a/libclc/libspirv/lib/generic/atomic/atomic_dec.cl b/libclc/libspirv/lib/generic/atomic/atomic_dec.cl index 6363432bc4958..156e881836cc2 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_dec.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_dec.cl @@ -8,46 +8,45 @@ #include -_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 diff --git a/libclc/libspirv/lib/generic/atomic/atomic_inc.cl b/libclc/libspirv/lib/generic/atomic/atomic_inc.cl index 5f6e6afd35f06..4d5f1473a136e 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_inc.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_inc.cl @@ -8,46 +8,45 @@ #include -_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 diff --git a/libclc/libspirv/lib/generic/atomic/atomic_load.cl b/libclc/libspirv/lib/generic/atomic/atomic_load.cl index 7e6b341a9c483..9e485f6d4f36f 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_load.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_load.cl @@ -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); \ } \ diff --git a/libclc/libspirv/lib/generic/atomic/atomic_max.cl b/libclc/libspirv/lib/generic/atomic/atomic_max.cl index 90f0f340bb3d2..e6665cbcb2fba 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_max.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_max.cl @@ -9,8 +9,8 @@ #include #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); \ } diff --git a/libclc/libspirv/lib/generic/atomic/atomic_min.cl b/libclc/libspirv/lib/generic/atomic/atomic_min.cl index d3df63524c916..d2f8d6360721b 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_min.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_min.cl @@ -9,8 +9,8 @@ #include #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); \ } diff --git a/libclc/libspirv/lib/generic/atomic/atomic_or.cl b/libclc/libspirv/lib/generic/atomic/atomic_or.cl index d2e78a062d78b..43259a7828210 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_or.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_or.cl @@ -9,8 +9,8 @@ #include #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); \ } diff --git a/libclc/libspirv/lib/generic/atomic/atomic_store.cl b/libclc/libspirv/lib/generic/atomic/atomic_store.cl index 00c732da06b3e..adec4a1b83416 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_store.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_store.cl @@ -8,13 +8,13 @@ #include -_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)); } @@ -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) { \ diff --git a/libclc/libspirv/lib/generic/atomic/atomic_sub.cl b/libclc/libspirv/lib/generic/atomic/atomic_sub.cl index 42cdd416aa119..bafa60d4ac3f8 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_sub.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_sub.cl @@ -9,8 +9,8 @@ #include #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); \ } diff --git a/libclc/libspirv/lib/generic/atomic/atomic_xor.cl b/libclc/libspirv/lib/generic/atomic/atomic_xor.cl index 860dcb189e1ff..291eba060c6e2 100644 --- a/libclc/libspirv/lib/generic/atomic/atomic_xor.cl +++ b/libclc/libspirv/lib/generic/atomic/atomic_xor.cl @@ -9,8 +9,8 @@ #include #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); \ } diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_add.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_add.cl index 2ecc9762bdf46..998d97c29d2b7 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_add.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_add.cl @@ -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; \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl index ab8f06a4acbe5..9a3408b229154 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_cmpxchg.cl @@ -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; \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_helpers.h b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_helpers.h index b9ef9c0d846c2..2b2ec209eaf05 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_helpers.h +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_helpers.h @@ -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; \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_inc_dec_helpers.h b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_inc_dec_helpers.h index 478d895c90ed7..eac740d1c693d 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_inc_dec_helpers.h +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_inc_dec_helpers.h @@ -13,9 +13,8 @@ #include #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); \ } diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_load.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_load.cl index da23ef23d3635..690cefc2eeb5a 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_load.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_load.cl @@ -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; \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_max.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_max.cl index f67a257965bb1..dcc381f491dce 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_max.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_max.cl @@ -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: \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_min.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_min.cl index 60fe3d0329c65..a9d1bfe327750 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_min.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_min.cl @@ -20,9 +20,8 @@ __CLC_NVVM_ATOMIC(ulong, ulong, ul, min, __spirv_AtomicUMin) #undef __CLC_NVVM_ATOMIC_IMPL #define __CLC_NVVM_ATOMIC_MIN_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: \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_store.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_store.cl index 3382366a38ade..61894ffe8facc 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_store.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_store.cl @@ -38,9 +38,8 @@ _CLC_OVERLOAD _CLC_DECL void __spirv_MemoryBarrier(int, int); #define __CLC_NVVM_ATOMIC_STORE_IMPL(TYPE, TYPE_NV, TYPE_MANGLED_NV, \ ADDR_SPACE, ADDR_SPACE_NV) \ - __attribute__((always_inline)) _CLC_OVERLOAD _CLC_DECL void \ - __spirv_AtomicStore(ADDR_SPACE TYPE *pointer, int scope, int semantics, \ - TYPE value) { \ + _CLC_OVERLOAD _CLC_DEF void __spirv_AtomicStore( \ + 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; \ diff --git a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_sub.cl b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_sub.cl index 1aa10828bee28..3db07435298f0 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_sub.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/atomic/atomic_sub.cl @@ -12,9 +12,8 @@ #define __CLC_NVVM_ATOMIC_SUB_IMPL(TYPE, OP_MANGLED, ADDR_SPACE) \ _CLC_OVERLOAD _CLC_DECL TYPE __spirv_AtomicIAdd(ADDR_SPACE TYPE *, int, int, \ TYPE); \ - __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) { \ return __spirv_AtomicIAdd(pointer, scope, semantics, -val); \ } diff --git a/libclc/libspirv/lib/ptx-nvidiacl/conversion/GenericCastToPtrExplicit.cl b/libclc/libspirv/lib/ptx-nvidiacl/conversion/GenericCastToPtrExplicit.cl index 3dda6d1f41fe4..b92c1d515b32b 100644 --- a/libclc/libspirv/lib/ptx-nvidiacl/conversion/GenericCastToPtrExplicit.cl +++ b/libclc/libspirv/lib/ptx-nvidiacl/conversion/GenericCastToPtrExplicit.cl @@ -19,16 +19,15 @@ _CLC_DEF static bool __clc_nvvm_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_nvvm_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); \ } From 26055725b71f687adc33f8df2b58ecd7998dfff5 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Tue, 16 Sep 2025 08:24:24 +0200 Subject: [PATCH 2/2] undo changes to clc --- libclc/clc/include/clc/misc/shuffle2_def.inc | 8 ++++---- libclc/clc/include/clc/misc/shuffle_def.inc | 8 ++++---- .../generic/atomic/clc_atomic_compare_exchange.inc | 4 ++-- libclc/clc/lib/generic/atomic/clc_atomic_def.inc | 12 ++++++------ 4 files changed, 16 insertions(+), 16 deletions(-) diff --git a/libclc/clc/include/clc/misc/shuffle2_def.inc b/libclc/clc/include/clc/misc/shuffle2_def.inc index 960419992ebf2..0415b8c187291 100644 --- a/libclc/clc/include/clc/misc/shuffle2_def.inc +++ b/libclc/clc/include/clc/misc/shuffle2_def.inc @@ -18,22 +18,22 @@ // The return type is same base type as the input type, with the same vector // size as the mask. Elements in the mask must be the same size (number of bits) // as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask); -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); } -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); } -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); } -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x, __CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) y, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, y, mask); diff --git a/libclc/clc/include/clc/misc/shuffle_def.inc b/libclc/clc/include/clc/misc/shuffle_def.inc index 29fcef15e34b8..d2a088498fc6c 100644 --- a/libclc/clc/include/clc/misc/shuffle_def.inc +++ b/libclc/clc/include/clc/misc/shuffle_def.inc @@ -18,19 +18,19 @@ // The return type is same base type as the input type, with the same vector // size as the mask. Elements in the mask must be the same size (number of bits) // as the input value., e.g. char8 ret = shuffle(char2 x, uchar8 mask); -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 2) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 4) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 8) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } -_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE +_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION(__CLC_XCONCAT(__CLC_SCALAR_GENTYPE, 16) x, __CLC_U_GENTYPE mask) { return __IMPL_FUNCTION(FUNCTION)(x, mask); } diff --git a/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc b/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc index 74284fd61024c..32ff9b45b769e 100644 --- a/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc +++ b/libclc/clc/lib/generic/atomic/clc_atomic_compare_exchange.inc @@ -24,7 +24,7 @@ #ifdef __CLC_FPSIZE #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_atomic_compare_exchange( \ + _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_atomic_compare_exchange( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Comparator, \ __CLC_GENTYPE Value, int MemoryOrderEqual, int MemoryOrderUnequal, \ int MemoryScope) { \ @@ -38,7 +38,7 @@ #else #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_atomic_compare_exchange( \ + _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_atomic_compare_exchange( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Comparator, \ __CLC_GENTYPE Value, int MemoryOrderEqual, int MemoryOrderUnequal, \ int MemoryScope) { \ diff --git a/libclc/clc/lib/generic/atomic/clc_atomic_def.inc b/libclc/clc/lib/generic/atomic/clc_atomic_def.inc index 635532c60db93..2c45f49f60848 100644 --- a/libclc/clc/lib/generic/atomic/clc_atomic_def.inc +++ b/libclc/clc/lib/generic/atomic/clc_atomic_def.inc @@ -31,7 +31,7 @@ #ifdef __CLC_NO_VALUE_ARG #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \ + _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, int MemoryOrder, \ int MemoryScope) { \ return __CLC_AS_RETTYPE(__IMPL_FUNCTION( \ @@ -39,7 +39,7 @@ } #elif defined(__CLC_INC_DEC) #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \ + _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, int MemoryOrder, \ int MemoryScope) { \ return __CLC_AS_RETTYPE( \ @@ -48,15 +48,15 @@ } #elif defined(__CLC_RETURN_VOID) #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DEF void FUNCTION(volatile ADDRSPACE __CLC_GENTYPE *Ptr, \ - __CLC_GENTYPE Value, int MemoryOrder, \ - int MemoryScope) { \ + _CLC_OVERLOAD _CLC_DECL void FUNCTION(volatile ADDRSPACE __CLC_GENTYPE *Ptr, \ + __CLC_GENTYPE Value, int MemoryOrder, \ + int MemoryScope) { \ __IMPL_FUNCTION((ADDRSPACE __CLC_PTR_CASTTYPE *)Ptr, Value, MemoryOrder, \ MemoryScope); \ } #else #define __CLC_DEFINE_ATOMIC(ADDRSPACE) \ - _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE FUNCTION( \ + _CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE FUNCTION( \ volatile ADDRSPACE __CLC_GENTYPE *Ptr, __CLC_GENTYPE Value, \ int MemoryOrder, int MemoryScope) { \ return __CLC_AS_RETTYPE( \