https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/220856
>From dcae190072318ae83d7854dc83888f36893005c2 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 1 Sep 2026 09:02:12 -0500 Subject: [PATCH 1/3] [CIR][AMDGPU][OpenCL] Classify __hip_atomic_*/__opencl_atomic_* operations 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. --- clang/lib/CIR/CodeGen/CIRGenAtomic.cpp | 94 ++++---- clang/test/CIR/CodeGenHIP/atomic-classify.hip | 207 ++++++++++++++++++ 2 files changed, 255 insertions(+), 46 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/atomic-classify.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index 4a028e0be34f7..6d2392c85eb1e 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; @@ -700,6 +704,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__atomic_load: case AtomicExpr::AO__scoped_atomic_load_n: 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()); @@ -715,6 +721,8 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, case AtomicExpr::AO__atomic_store: case AtomicExpr::AO__scoped_atomic_store: 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; @@ -870,41 +894,17 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, 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; @@ -1219,10 +1219,6 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_compare_exchange_n: case AtomicExpr::AO__c11_atomic_compare_exchange_weak: case AtomicExpr::AO__c11_atomic_compare_exchange_strong: - case AtomicExpr::AO__hip_atomic_compare_exchange_weak: - case AtomicExpr::AO__hip_atomic_compare_exchange_strong: - case AtomicExpr::AO__opencl_atomic_compare_exchange_weak: - case AtomicExpr::AO__opencl_atomic_compare_exchange_strong: case AtomicExpr::AO__scoped_atomic_compare_exchange: case AtomicExpr::AO__scoped_atomic_compare_exchange_n: cgf.cgm.errorNYI( @@ -1234,8 +1230,6 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_exchange: case AtomicExpr::AO__atomic_exchange_n: case AtomicExpr::AO__c11_atomic_exchange: - case AtomicExpr::AO__hip_atomic_exchange: - case AtomicExpr::AO__opencl_atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange_n: cgf.cgm.errorNYI(loc, "emitLibCallForAtomicExpr: atomic exchange NYI"); @@ -1256,8 +1250,6 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, break; } - case AtomicExpr::AO__hip_atomic_store: - case AtomicExpr::AO__opencl_atomic_store: cgf.cgm.errorNYI(loc, "emitLibCallForAtomicExpr: atomic store for hip/opencl"); return RValue::get(nullptr); @@ -1294,36 +1286,26 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__scoped_atomic_add_fetch: case AtomicExpr::AO__atomic_fetch_add: case AtomicExpr::AO__c11_atomic_fetch_add: - case AtomicExpr::AO__hip_atomic_fetch_add: - case AtomicExpr::AO__opencl_atomic_fetch_add: case AtomicExpr::AO__scoped_atomic_fetch_add: case AtomicExpr::AO__atomic_and_fetch: case AtomicExpr::AO__scoped_atomic_and_fetch: case AtomicExpr::AO__atomic_fetch_and: case AtomicExpr::AO__c11_atomic_fetch_and: - case AtomicExpr::AO__hip_atomic_fetch_and: - case AtomicExpr::AO__opencl_atomic_fetch_and: case AtomicExpr::AO__scoped_atomic_fetch_and: case AtomicExpr::AO__atomic_or_fetch: case AtomicExpr::AO__scoped_atomic_or_fetch: case AtomicExpr::AO__atomic_fetch_or: case AtomicExpr::AO__c11_atomic_fetch_or: - case AtomicExpr::AO__hip_atomic_fetch_or: - case AtomicExpr::AO__opencl_atomic_fetch_or: case AtomicExpr::AO__scoped_atomic_fetch_or: case AtomicExpr::AO__atomic_sub_fetch: case AtomicExpr::AO__scoped_atomic_sub_fetch: case AtomicExpr::AO__atomic_fetch_sub: case AtomicExpr::AO__c11_atomic_fetch_sub: - case AtomicExpr::AO__hip_atomic_fetch_sub: - case AtomicExpr::AO__opencl_atomic_fetch_sub: case AtomicExpr::AO__scoped_atomic_fetch_sub: case AtomicExpr::AO__atomic_xor_fetch: case AtomicExpr::AO__scoped_atomic_xor_fetch: case AtomicExpr::AO__atomic_fetch_xor: case AtomicExpr::AO__c11_atomic_fetch_xor: - case AtomicExpr::AO__hip_atomic_fetch_xor: - case AtomicExpr::AO__opencl_atomic_fetch_xor: case AtomicExpr::AO__scoped_atomic_fetch_xor: case AtomicExpr::AO__atomic_nand_fetch: case AtomicExpr::AO__atomic_fetch_nand: @@ -1333,15 +1315,11 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_min_fetch: case AtomicExpr::AO__atomic_fetch_min: case AtomicExpr::AO__c11_atomic_fetch_min: - case AtomicExpr::AO__hip_atomic_fetch_min: - case AtomicExpr::AO__opencl_atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_min_fetch: case AtomicExpr::AO__atomic_max_fetch: case AtomicExpr::AO__atomic_fetch_max: case AtomicExpr::AO__c11_atomic_fetch_max: - case AtomicExpr::AO__hip_atomic_fetch_max: - case AtomicExpr::AO__opencl_atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_max_fetch: case AtomicExpr::AO__scoped_atomic_fetch_uinc: @@ -1433,6 +1411,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; @@ -1459,6 +1439,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) @@ -1560,6 +1544,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..3c19ba393d4c0 --- /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 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 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 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 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 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 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 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 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); +} >From d265334c38dc7c8add61c9a9daa3a14016ebbc8e Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Thu, 3 Sep 2026 06:19:02 -0500 Subject: [PATCH 2/3] Fix formatting Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenAtomic.cpp | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index 6d2392c85eb1e..c70dcd6d32e4c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -893,19 +893,6 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, break; case AtomicExpr::AO__opencl_atomic_init: - - - - - - - - - - - - - cgf.cgm.errorNYI(expr->getSourceRange(), "emitAtomicOp: expr op NYI"); return; } >From b39e7571950fb16e2e16bece8ac9c389f77d1ebe Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Thu, 3 Sep 2026 06:24:57 -0500 Subject: [PATCH 3/3] Add back switch cases Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenAtomic.cpp | 23 ++++++++++++++++++++++- 1 file changed, 22 insertions(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index c70dcd6d32e4c..9d28443d98238 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -1206,6 +1206,10 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_compare_exchange_n: case AtomicExpr::AO__c11_atomic_compare_exchange_weak: case AtomicExpr::AO__c11_atomic_compare_exchange_strong: + case AtomicExpr::AO__hip_atomic_compare_exchange_weak: + case AtomicExpr::AO__hip_atomic_compare_exchange_strong: + case AtomicExpr::AO__opencl_atomic_compare_exchange_weak: + case AtomicExpr::AO__opencl_atomic_compare_exchange_strong: case AtomicExpr::AO__scoped_atomic_compare_exchange: case AtomicExpr::AO__scoped_atomic_compare_exchange_n: cgf.cgm.errorNYI( @@ -1217,6 +1221,8 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_exchange: case AtomicExpr::AO__atomic_exchange_n: case AtomicExpr::AO__c11_atomic_exchange: + case AtomicExpr::AO__hip_atomic_exchange: + case AtomicExpr::AO__opencl_atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange_n: cgf.cgm.errorNYI(loc, "emitLibCallForAtomicExpr: atomic exchange NYI"); @@ -1236,7 +1242,8 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, cgf.getContext().VoidPtrTy); break; } - + case AtomicExpr::AO__hip_atomic_store: + case AtomicExpr::AO__opencl_atomic_store: cgf.cgm.errorNYI(loc, "emitLibCallForAtomicExpr: atomic store for hip/opencl"); return RValue::get(nullptr); @@ -1273,26 +1280,36 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__scoped_atomic_add_fetch: case AtomicExpr::AO__atomic_fetch_add: case AtomicExpr::AO__c11_atomic_fetch_add: + case AtomicExpr::AO__hip_atomic_fetch_add: + case AtomicExpr::AO__opencl_atomic_fetch_add: case AtomicExpr::AO__scoped_atomic_fetch_add: case AtomicExpr::AO__atomic_and_fetch: case AtomicExpr::AO__scoped_atomic_and_fetch: case AtomicExpr::AO__atomic_fetch_and: case AtomicExpr::AO__c11_atomic_fetch_and: + case AtomicExpr::AO__hip_atomic_fetch_and: + case AtomicExpr::AO__opencl_atomic_fetch_and: case AtomicExpr::AO__scoped_atomic_fetch_and: case AtomicExpr::AO__atomic_or_fetch: case AtomicExpr::AO__scoped_atomic_or_fetch: case AtomicExpr::AO__atomic_fetch_or: case AtomicExpr::AO__c11_atomic_fetch_or: + case AtomicExpr::AO__hip_atomic_fetch_or: + case AtomicExpr::AO__opencl_atomic_fetch_or: case AtomicExpr::AO__scoped_atomic_fetch_or: case AtomicExpr::AO__atomic_sub_fetch: case AtomicExpr::AO__scoped_atomic_sub_fetch: case AtomicExpr::AO__atomic_fetch_sub: case AtomicExpr::AO__c11_atomic_fetch_sub: + case AtomicExpr::AO__hip_atomic_fetch_sub: + case AtomicExpr::AO__opencl_atomic_fetch_sub: case AtomicExpr::AO__scoped_atomic_fetch_sub: case AtomicExpr::AO__atomic_xor_fetch: case AtomicExpr::AO__scoped_atomic_xor_fetch: case AtomicExpr::AO__atomic_fetch_xor: case AtomicExpr::AO__c11_atomic_fetch_xor: + case AtomicExpr::AO__hip_atomic_fetch_xor: + case AtomicExpr::AO__opencl_atomic_fetch_xor: case AtomicExpr::AO__scoped_atomic_fetch_xor: case AtomicExpr::AO__atomic_nand_fetch: case AtomicExpr::AO__atomic_fetch_nand: @@ -1302,11 +1319,15 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_min_fetch: case AtomicExpr::AO__atomic_fetch_min: case AtomicExpr::AO__c11_atomic_fetch_min: + case AtomicExpr::AO__hip_atomic_fetch_min: + case AtomicExpr::AO__opencl_atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_min_fetch: case AtomicExpr::AO__atomic_max_fetch: case AtomicExpr::AO__atomic_fetch_max: case AtomicExpr::AO__c11_atomic_fetch_max: + case AtomicExpr::AO__hip_atomic_fetch_max: + case AtomicExpr::AO__opencl_atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_max_fetch: case AtomicExpr::AO__scoped_atomic_fetch_uinc: _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
