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); \ }