Author: Steffen Larsen Date: 2026-09-17T09:00:05+02:00 New Revision: f557f80378b08024b79efbde73fbe3a03036abf7
URL: https://github.com/llvm/llvm-project/commit/f557f80378b08024b79efbde73fbe3a03036abf7 DIFF: https://github.com/llvm/llvm-project/commit/f557f80378b08024b79efbde73fbe3a03036abf7.diff LOG: [CIR][AMDGPU][OpenCL] Classify __hip_atomic_*/__opencl_atomic_* operations (#220856) load, store, exchange, compare_exchange_strong/weak, and every fetch op were missing from the operand-classification switches, so they fell to the default and reported an unimplemented atomic. They take the same operands as their __c11_atomic_*/__atomic_* counterparts, so they belong with them. This is the ground work for implementing more AMDGPU atomics. --------- Signed-off-by: Steffen Holst Larsen <[email protected]> Added: clang/test/CIR/CodeGenHIP/atomic-classify.hip Modified: clang/lib/CIR/CodeGen/CIRGenAtomic.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index 8bbcba6b576e6..c4c3b455bf11c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -670,11 +670,15 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, llvm_unreachable("already handled!"); case AtomicExpr::AO__c11_atomic_compare_exchange_strong: + case AtomicExpr::AO__hip_atomic_compare_exchange_strong: + case AtomicExpr::AO__opencl_atomic_compare_exchange_strong: emitAtomicCmpXchgFailureSet(cgf, expr, /*isWeak=*/false, dest, ptr, val1, val2, failureOrderExpr, size, order, scope); return; case AtomicExpr::AO__c11_atomic_compare_exchange_weak: + case AtomicExpr::AO__hip_atomic_compare_exchange_weak: + case AtomicExpr::AO__opencl_atomic_compare_exchange_weak: emitAtomicCmpXchgFailureSet(cgf, expr, /*isWeak=*/true, dest, ptr, val1, val2, failureOrderExpr, size, order, scope); return; @@ -699,7 +703,9 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__atomic_load_n: case AtomicExpr::AO__atomic_load: case AtomicExpr::AO__scoped_atomic_load_n: - case AtomicExpr::AO__scoped_atomic_load: { + case AtomicExpr::AO__scoped_atomic_load: + case AtomicExpr::AO__hip_atomic_load: + case AtomicExpr::AO__opencl_atomic_load: { cir::LoadOp load = builder.createLoad(loc, ptr, /*isVolatile=*/expr->isVolatile()); @@ -714,7 +720,9 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__atomic_store_n: case AtomicExpr::AO__atomic_store: case AtomicExpr::AO__scoped_atomic_store: - case AtomicExpr::AO__scoped_atomic_store_n: { + case AtomicExpr::AO__scoped_atomic_store_n: + case AtomicExpr::AO__hip_atomic_store: + case AtomicExpr::AO__opencl_atomic_store: { cir::LoadOp loadVal1 = builder.createLoad(loc, val1); assert(!cir::MissingFeatures::atomicSyncScopeID()); @@ -730,6 +738,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange_n: case AtomicExpr::AO__scoped_atomic_exchange: + case AtomicExpr::AO__hip_atomic_exchange: + case AtomicExpr::AO__opencl_atomic_exchange: opName = cir::AtomicXchgOp::getOperationName(); break; @@ -740,6 +750,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_add: case AtomicExpr::AO__atomic_fetch_add: case AtomicExpr::AO__scoped_atomic_fetch_add: + case AtomicExpr::AO__hip_atomic_fetch_add: + case AtomicExpr::AO__opencl_atomic_fetch_add: handleFetchOp(cir::AtomicFetchKind::Add); break; @@ -750,6 +762,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_sub: case AtomicExpr::AO__atomic_fetch_sub: case AtomicExpr::AO__scoped_atomic_fetch_sub: + case AtomicExpr::AO__hip_atomic_fetch_sub: + case AtomicExpr::AO__opencl_atomic_fetch_sub: handleFetchOp(cir::AtomicFetchKind::Sub); break; @@ -760,6 +774,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_min: case AtomicExpr::AO__atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_fetch_min: + case AtomicExpr::AO__hip_atomic_fetch_min: + case AtomicExpr::AO__opencl_atomic_fetch_min: handleFetchOp(cir::AtomicFetchKind::Min); break; @@ -784,6 +800,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_max: case AtomicExpr::AO__atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_fetch_max: + case AtomicExpr::AO__hip_atomic_fetch_max: + case AtomicExpr::AO__opencl_atomic_fetch_max: handleFetchOp(cir::AtomicFetchKind::Max); break; @@ -808,6 +826,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_and: case AtomicExpr::AO__atomic_fetch_and: case AtomicExpr::AO__scoped_atomic_fetch_and: + case AtomicExpr::AO__hip_atomic_fetch_and: + case AtomicExpr::AO__opencl_atomic_fetch_and: handleFetchOp(cir::AtomicFetchKind::And); break; @@ -818,6 +838,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_or: case AtomicExpr::AO__atomic_fetch_or: case AtomicExpr::AO__scoped_atomic_fetch_or: + case AtomicExpr::AO__hip_atomic_fetch_or: + case AtomicExpr::AO__opencl_atomic_fetch_or: handleFetchOp(cir::AtomicFetchKind::Or); break; @@ -828,6 +850,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__c11_atomic_fetch_xor: case AtomicExpr::AO__atomic_fetch_xor: case AtomicExpr::AO__scoped_atomic_fetch_xor: + case AtomicExpr::AO__hip_atomic_fetch_xor: + case AtomicExpr::AO__opencl_atomic_fetch_xor: handleFetchOp(cir::AtomicFetchKind::Xor); break; @@ -869,43 +893,6 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, break; case AtomicExpr::AO__opencl_atomic_init: - - case AtomicExpr::AO__hip_atomic_compare_exchange_strong: - case AtomicExpr::AO__opencl_atomic_compare_exchange_strong: - - case AtomicExpr::AO__opencl_atomic_compare_exchange_weak: - case AtomicExpr::AO__hip_atomic_compare_exchange_weak: - - case AtomicExpr::AO__opencl_atomic_load: - case AtomicExpr::AO__hip_atomic_load: - - case AtomicExpr::AO__opencl_atomic_store: - case AtomicExpr::AO__hip_atomic_store: - - case AtomicExpr::AO__hip_atomic_exchange: - case AtomicExpr::AO__opencl_atomic_exchange: - - case AtomicExpr::AO__hip_atomic_fetch_add: - case AtomicExpr::AO__opencl_atomic_fetch_add: - - case AtomicExpr::AO__hip_atomic_fetch_sub: - case AtomicExpr::AO__opencl_atomic_fetch_sub: - - case AtomicExpr::AO__hip_atomic_fetch_min: - case AtomicExpr::AO__opencl_atomic_fetch_min: - - case AtomicExpr::AO__hip_atomic_fetch_max: - case AtomicExpr::AO__opencl_atomic_fetch_max: - - case AtomicExpr::AO__hip_atomic_fetch_and: - case AtomicExpr::AO__opencl_atomic_fetch_and: - - case AtomicExpr::AO__hip_atomic_fetch_or: - case AtomicExpr::AO__opencl_atomic_fetch_or: - - case AtomicExpr::AO__hip_atomic_fetch_xor: - case AtomicExpr::AO__opencl_atomic_fetch_xor: - cgf.cgm.errorNYI(expr->getSourceRange(), "emitAtomicOp: expr op NYI"); return; } @@ -1435,6 +1422,8 @@ RValue CIRGenFunction::emitAtomicExpr(AtomicExpr *e) { case AtomicExpr::AO__atomic_load_n: case AtomicExpr::AO__scoped_atomic_load_n: case AtomicExpr::AO__c11_atomic_load: + case AtomicExpr::AO__opencl_atomic_load: + case AtomicExpr::AO__hip_atomic_load: case AtomicExpr::AO__atomic_test_and_set: case AtomicExpr::AO__atomic_clear: break; @@ -1461,6 +1450,10 @@ RValue CIRGenFunction::emitAtomicExpr(AtomicExpr *e) { case AtomicExpr::AO__c11_atomic_compare_exchange_strong: case AtomicExpr::AO__scoped_atomic_compare_exchange: case AtomicExpr::AO__scoped_atomic_compare_exchange_n: + case AtomicExpr::AO__hip_atomic_compare_exchange_strong: + case AtomicExpr::AO__hip_atomic_compare_exchange_weak: + case AtomicExpr::AO__opencl_atomic_compare_exchange_strong: + case AtomicExpr::AO__opencl_atomic_compare_exchange_weak: val1 = emitPointerWithAlignment(e->getVal1()); if (e->getOp() == AtomicExpr::AO__atomic_compare_exchange || e->getOp() == AtomicExpr::AO__scoped_atomic_compare_exchange) @@ -1562,6 +1555,24 @@ RValue CIRGenFunction::emitAtomicExpr(AtomicExpr *e) { case AtomicExpr::AO__atomic_fetch_fmaximum: case AtomicExpr::AO__atomic_fetch_fminimum_num: case AtomicExpr::AO__atomic_fetch_fmaximum_num: + case AtomicExpr::AO__hip_atomic_exchange: + case AtomicExpr::AO__hip_atomic_store: + case AtomicExpr::AO__hip_atomic_fetch_add: + case AtomicExpr::AO__hip_atomic_fetch_sub: + case AtomicExpr::AO__hip_atomic_fetch_min: + case AtomicExpr::AO__hip_atomic_fetch_max: + case AtomicExpr::AO__hip_atomic_fetch_and: + case AtomicExpr::AO__hip_atomic_fetch_or: + case AtomicExpr::AO__hip_atomic_fetch_xor: + case AtomicExpr::AO__opencl_atomic_exchange: + case AtomicExpr::AO__opencl_atomic_store: + case AtomicExpr::AO__opencl_atomic_fetch_add: + case AtomicExpr::AO__opencl_atomic_fetch_sub: + case AtomicExpr::AO__opencl_atomic_fetch_min: + case AtomicExpr::AO__opencl_atomic_fetch_max: + case AtomicExpr::AO__opencl_atomic_fetch_and: + case AtomicExpr::AO__opencl_atomic_fetch_or: + case AtomicExpr::AO__opencl_atomic_fetch_xor: val1 = emitValToTemp(*this, e->getVal1()); break; } diff --git a/clang/test/CIR/CodeGenHIP/atomic-classify.hip b/clang/test/CIR/CodeGenHIP/atomic-classify.hip new file mode 100644 index 0000000000000..ad3ad771e6fbe --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/atomic-classify.hip @@ -0,0 +1,207 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fcuda-is-device \ +// RUN: -emit-cir %s -o - | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fcuda-is-device \ +// RUN: -emit-llvm %s -o - | FileCheck --check-prefix=LLVM %s + +#define __device__ __attribute__((device)) + +// CIR-LABEL: @_Z8hip_loadPi +// CIR: cir.load {{.*}} atomic(relaxed) +// LLVM-LABEL: @_Z8hip_loadPi +// LLVM: load atomic i32, ptr {{.*}} monotonic, align 4 +__device__ int hip_load(int *p) { + return __hip_atomic_load(p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z11opencl_loadPVU7_Atomici +// CIR: cir.load volatile {{.*}} atomic(relaxed) +// LLVM-LABEL: @_Z11opencl_loadPVU7_Atomici +// LLVM: load atomic volatile i32, ptr {{.*}} monotonic, align 4 +__device__ int opencl_load(_Atomic(int) volatile *p) { + return __opencl_atomic_load(p, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z9hip_storePii +// CIR: cir.store {{.*}} atomic(relaxed) +// LLVM-LABEL: @_Z9hip_storePii +// LLVM: store atomic i32 {{.*}} monotonic, align 4 +__device__ void hip_store(int *p, int v) { + __hip_atomic_store(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z12opencl_storePVU7_Atomicii +// CIR: cir.store volatile {{.*}} atomic(relaxed) +// LLVM-LABEL: @_Z12opencl_storePVU7_Atomicii +// LLVM: store atomic volatile i32 {{.*}} monotonic, align 4 +__device__ void opencl_store(_Atomic(int) volatile *p, int v) { + __opencl_atomic_store(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z12hip_exchangePii +// CIR: cir.atomic.xchg relaxed +// LLVM-LABEL: @_Z12hip_exchangePii +// LLVM: atomicrmw xchg ptr {{.*}} monotonic +__device__ int hip_exchange(int *p, int v) { + return __hip_atomic_exchange(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z15opencl_exchangePVU7_Atomicii +// CIR: cir.atomic.xchg relaxed +// LLVM-LABEL: @_Z15opencl_exchangePVU7_Atomicii +// LLVM: atomicrmw volatile xchg ptr {{.*}} monotonic +__device__ int opencl_exchange(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_exchange(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z11hip_cmpxchgPiii +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) +// LLVM-LABEL: @_Z11hip_cmpxchgPiii +// LLVM: cmpxchg ptr {{.*}} monotonic monotonic +__device__ bool hip_cmpxchg(int *p, int expected, int v) { + return __hip_atomic_compare_exchange_strong(p, &expected, v, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16hip_cmpxchg_weakPiii +// CIR: cir.atomic.cmpxchg weak success(relaxed) failure(relaxed) +// LLVM-LABEL: @_Z16hip_cmpxchg_weakPiii +// LLVM: cmpxchg weak ptr {{.*}} monotonic monotonic +__device__ bool hip_cmpxchg_weak(int *p, int expected, int v) { + return __hip_atomic_compare_exchange_weak(p, &expected, v, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z14opencl_cmpxchgPVU7_AtomiciPii +// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) +// LLVM-LABEL: @_Z14opencl_cmpxchgPVU7_AtomiciPii +// LLVM: cmpxchg volatile ptr {{.*}} monotonic monotonic +__device__ bool opencl_cmpxchg(_Atomic(int) volatile *p, int *expected, int v) { + return __opencl_atomic_compare_exchange_strong(p, expected, v, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, + __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z19opencl_cmpxchg_weakPVU7_AtomiciPii +// CIR: cir.atomic.cmpxchg weak success(relaxed) failure(relaxed) +// LLVM-LABEL: @_Z19opencl_cmpxchg_weakPVU7_AtomiciPii +// LLVM: cmpxchg weak volatile ptr {{.*}} monotonic monotonic +__device__ bool opencl_cmpxchg_weak(_Atomic(int) volatile *p, int *expected, int v) { + return __opencl_atomic_compare_exchange_weak(p, expected, v, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, + __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z13hip_fetch_addPii +// CIR: cir.atomic.fetch add relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z13hip_fetch_addPii +// LLVM: atomicrmw add ptr {{.*}} monotonic +__device__ int hip_fetch_add(int *p, int v) { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16opencl_fetch_addPVU7_Atomicii +// CIR: cir.atomic.fetch add relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z16opencl_fetch_addPVU7_Atomicii +// LLVM: atomicrmw volatile add ptr {{.*}} monotonic +__device__ int opencl_fetch_add(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z13hip_fetch_subPii +// CIR: cir.atomic.fetch sub relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z13hip_fetch_subPii +// LLVM: atomicrmw sub ptr {{.*}} monotonic +__device__ int hip_fetch_sub(int *p, int v) { + return __hip_atomic_fetch_sub(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16opencl_fetch_subPVU7_Atomicii +// CIR: cir.atomic.fetch sub relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z16opencl_fetch_subPVU7_Atomicii +// LLVM: atomicrmw volatile sub ptr {{.*}} monotonic +__device__ int opencl_fetch_sub(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_sub(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z13hip_fetch_minPii +// CIR: cir.atomic.fetch min relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z13hip_fetch_minPii +// LLVM: atomicrmw min ptr {{.*}} monotonic +__device__ int hip_fetch_min(int *p, int v) { + return __hip_atomic_fetch_min(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16opencl_fetch_minPVU7_Atomicii +// CIR: cir.atomic.fetch min relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z16opencl_fetch_minPVU7_Atomicii +// LLVM: atomicrmw volatile min ptr {{.*}} monotonic +__device__ int opencl_fetch_min(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_min(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z13hip_fetch_maxPii +// CIR: cir.atomic.fetch max relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z13hip_fetch_maxPii +// LLVM: atomicrmw max ptr {{.*}} monotonic +__device__ int hip_fetch_max(int *p, int v) { + return __hip_atomic_fetch_max(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16opencl_fetch_maxPVU7_Atomicii +// CIR: cir.atomic.fetch max relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z16opencl_fetch_maxPVU7_Atomicii +// LLVM: atomicrmw volatile max ptr {{.*}} monotonic +__device__ int opencl_fetch_max(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_max(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z13hip_fetch_andPii +// CIR: cir.atomic.fetch and relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z13hip_fetch_andPii +// LLVM: atomicrmw and ptr {{.*}} monotonic +__device__ int hip_fetch_and(int *p, int v) { + return __hip_atomic_fetch_and(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16opencl_fetch_andPVU7_Atomicii +// CIR: cir.atomic.fetch and relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z16opencl_fetch_andPVU7_Atomicii +// LLVM: atomicrmw volatile and ptr {{.*}} monotonic +__device__ int opencl_fetch_and(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_and(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z12hip_fetch_orPii +// CIR: cir.atomic.fetch or relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z12hip_fetch_orPii +// LLVM: atomicrmw or ptr {{.*}} monotonic +__device__ int hip_fetch_or(int *p, int v) { + return __hip_atomic_fetch_or(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z15opencl_fetch_orPVU7_Atomicii +// CIR: cir.atomic.fetch or relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z15opencl_fetch_orPVU7_Atomicii +// LLVM: atomicrmw volatile or ptr {{.*}} monotonic +__device__ int opencl_fetch_or(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_or(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} + +// CIR-LABEL: @_Z13hip_fetch_xorPii +// CIR: cir.atomic.fetch xor relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z13hip_fetch_xorPii +// LLVM: atomicrmw xor ptr {{.*}} monotonic +__device__ int hip_fetch_xor(int *p, int v) { + return __hip_atomic_fetch_xor(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); +} + +// CIR-LABEL: @_Z16opencl_fetch_xorPVU7_Atomicii +// CIR: cir.atomic.fetch xor relaxed {{.*}} fetch_first +// LLVM-LABEL: @_Z16opencl_fetch_xorPVU7_Atomicii +// LLVM: atomicrmw volatile xor ptr {{.*}} monotonic +__device__ int opencl_fetch_xor(_Atomic(int) volatile *p, int v) { + return __opencl_atomic_fetch_xor(p, v, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_DEVICE); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
