[clang] f557f80 - [CIR][AMDGPU][OpenCL] Classify __hip_atomic_*/__opencl_atomic_* operations (#220856)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 17 00:00:10 PDT 2026
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 <sholstla at amd.com>
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);
+}
More information about the cfe-commits
mailing list