SWDEV-448288 - Remove references to deprecated llvm references.
Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72
Bu işleme şunda yer alıyor:
işlemeyi yapan:
Karthik Jayaprakash
ebeveyn
7f84df9f74
işleme
5b28e386f6
@@ -973,22 +973,13 @@ inline
|
||||
unsigned int atomicInc(unsigned int* address, unsigned int val)
|
||||
{
|
||||
#if defined(__gfx941__)
|
||||
__device__
|
||||
extern
|
||||
unsigned int __builtin_amdgcn_atomic_inc(
|
||||
unsigned int*,
|
||||
unsigned int,
|
||||
unsigned int,
|
||||
unsigned int,
|
||||
bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32");
|
||||
|
||||
return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
|
||||
address,
|
||||
val,
|
||||
[](unsigned int& x, unsigned int y) { x = (x >= y) ? 0 : (x + 1); },
|
||||
[=]() {
|
||||
return
|
||||
__builtin_amdgcn_atomic_inc(address, val, __ATOMIC_RELAXED, 1, false);
|
||||
__builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
|
||||
});
|
||||
#else
|
||||
return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
|
||||
@@ -1001,22 +992,13 @@ inline
|
||||
unsigned int atomicDec(unsigned int* address, unsigned int val)
|
||||
{
|
||||
#if defined(__gfx941__)
|
||||
__device__
|
||||
extern
|
||||
unsigned int __builtin_amdgcn_atomic_dec(
|
||||
unsigned int*,
|
||||
unsigned int,
|
||||
unsigned int,
|
||||
unsigned int,
|
||||
bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32");
|
||||
|
||||
return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
|
||||
address,
|
||||
val,
|
||||
[](unsigned int& x, unsigned int y) { x = (!x || x > y) ? y : (x - 1); },
|
||||
[=]() {
|
||||
return
|
||||
__builtin_amdgcn_atomic_dec(address, val, __ATOMIC_RELAXED, 1, false);
|
||||
__builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
|
||||
});
|
||||
#else
|
||||
return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
|
||||
|
||||
Yeni konuda referans
Bir kullanıcı engelle