This is an archive of the discontinued LLVM Phabricator instance.

Differential D157388

[AMDGPU] Support FMin/FMax in AMDGPUAtomicOptimizer.
ClosedPublic

Authored by pravinjagtap on Aug 8 2023, 5:40 AM.

Details

Reviewers
arsenm
b-sumner
foad
cdevadas
Group Reviewers
Restricted Project
Commits
rGedb9fab39022: [AMDGPU] Support FMin/FMax in AMDGPUAtomicOptimizer.

Diff Detail

Event Timeline

pravinjagtap created this revision.Aug 8 2023, 5:40 AM
Herald added a project: Restricted Project. · View Herald TranscriptAug 8 2023, 5:40 AM
Herald added subscribers: foad, kerbowa, hiraditya and 5 others. · View Herald Transcript
pravinjagtap requested review of this revision.Aug 8 2023, 5:40 AM
Herald added a project: Restricted Project. · View Herald TranscriptAug 8 2023, 5:40 AM
Herald added subscribers: llvm-commits, wdng. · View Herald Transcript
Harbormaster completed remote builds in B251068: Diff 548170.Aug 8 2023, 5:41 AM
pravinjagtap added a parent revision: D157265: [AMDGPU] Reorder atomic optimizer to avoid CAS loop..Aug 8 2023, 5:43 AM
pravinjagtap added reviewers: b-sumner, foad, cdevadas.
Herald added a subscriber: StephenFan. · View Herald TranscriptAug 8 2023, 5:44 AM
arsenm added inline comments.Aug 8 2023, 5:54 AM
llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp
320

I don't want to implicitly convert the intrinsics here. We should move towards getting rid of the intrinsics and autoupgrading them

405

This is incorrect, you should create minnum/maxnum

pravinjagtap added a comment.Aug 8 2023, 5:58 AM

For FMin and FMax cases, clang itself is emitting CAS loop for both

__device__ inline float atomicMax(float* addr, float val) and
__device__ inline float unsafeAtomicMax(float* addr, float val)

I am not sure how to potentially avoid this CAS loop before we reach atomic optimization pass.
CC: @b-sumner @arsenm

llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp
405

@arsenm you earlier suggested to use minnum/maxnum intrinsics for this. This also seems to give correct behavior. I am not sure what I am missing here

arsenm added inline comments.Aug 8 2023, 6:05 AM
llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp
405

Yes, it is wrong to use fcmp and select here. For example for fmax what you have returns the wrong result if LHS is a nan.

select (ugt nan, rhs), nan, rhs -> nan
maxnum(nan, rhs) -> rhs

arsenm added a comment.Aug 8 2023, 6:05 AM
In D157388#4569245, @pravinjagtap wrote:

For FMin and FMax cases, clang itself is emitting CAS loop for both

__device__ inline float atomicMax(float* addr, float val) and
__device__ inline float unsafeAtomicMax(float* addr, float val)

I am not sure how to potentially avoid this CAS loop before we reach atomic optimization pass.
CC: @b-sumner @arsenm

Clang should not be expanding any atomics itself

pravinjagtap added inline comments.Aug 8 2023, 6:08 AM
llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp
405

you should create minnum/maxnum

Are you referring to @llvm.amdgcn.fcmp.f32(float, float, i32) intrinsic here right ?

arsenm added inline comments.Aug 8 2023, 6:09 AM
llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp
405

No, IRBuilder.CreateMinNum and CreateMaxNum

pravinjagtap updated this revision to Diff 548193.Aug 8 2023, 6:33 AM

Switched to CreateMinNum/CreateMaxNum than fcmp and select

arsenm added a comment.Aug 8 2023, 6:46 AM

Also should have some end to end codegen tests, just this won't catch the interaction between the atomic expand and atomic optimizer

llvm/test/CodeGen/AMDGPU/global_atomics_iterative_scan_fp.ll
543

Should also test with different scopes, at least default system and agent

561

Should take this from the command line

Harbormaster completed remote builds in B251082: Diff 548193.Aug 8 2023, 8:54 AM
pravinjagtap updated this revision to Diff 548254.Aug 8 2023, 9:26 AM

Added few more test points

pravinjagtap added a comment.Aug 8 2023, 9:39 AM
In D157388#4569250, @arsenm wrote:
In D157388#4569245, @pravinjagtap wrote:

For FMin and FMax cases, clang itself is emitting CAS loop for both

__device__ inline float atomicMax(float* addr, float val) and
__device__ inline float unsafeAtomicMax(float* addr, float val)

I am not sure how to potentially avoid this CAS loop before we reach atomic optimization pass.
CC: @b-sumner @arsenm

Clang should not be expanding any atomics itself

All this is being implemented outside of llvm. The logic for creating CAS loop is inserted in hipamd/include/hip/amd_detail/amd_hip_atomic.h. How should we go about this? do I need to update hipamd repo ? CC: @b-sumner @arsenm

Harbormaster completed remote builds in B251135: Diff 548254.Aug 8 2023, 1:12 PM
pravinjagtap added a child revision: D157495: [Atomic-Expand] Run SimplifyCFG from Atomic-Expand on CAS loop blocks..Aug 9 2023, 4:54 AM
pravinjagtap updated this revision to Diff 551059.Aug 17 2023, 2:57 AM

Rebase & test-updates

Harbormaster completed remote builds in B253163: Diff 551059.Aug 17 2023, 2:58 AM
pravinjagtap added a reviewer: Restricted Project.Aug 17 2023, 3:47 AM
pravinjagtap removed a parent revision: D157265: [AMDGPU] Reorder atomic optimizer to avoid CAS loop..Aug 17 2023, 5:24 AM
pravinjagtap removed a child revision: D157495: [Atomic-Expand] Run SimplifyCFG from Atomic-Expand on CAS loop blocks..Aug 17 2023, 5:35 AM
pravinjagtap added a parent revision: D157265: [AMDGPU] Reorder atomic optimizer to avoid CAS loop..Aug 17 2023, 8:11 AM
pravinjagtap updated this revision to Diff 551145.Aug 17 2023, 8:21 AM

Rebased

pravinjagtap added a child revision: D157495: [Atomic-Expand] Run SimplifyCFG from Atomic-Expand on CAS loop blocks..Aug 17 2023, 8:21 AM
Harbormaster completed remote builds in B253226: Diff 551145.Aug 17 2023, 9:36 AM
pravinjagtap updated this revision to Diff 551428.Aug 18 2023, 1:42 AM

Addressed review comments and rebase

pravinjagtap removed a child revision: D157495: [Atomic-Expand] Run SimplifyCFG from Atomic-Expand on CAS loop blocks..Aug 18 2023, 1:45 AM
Harbormaster completed remote builds in B253432: Diff 551428.Aug 18 2023, 2:49 AM
pravinjagtap mentioned this in rGc931f2e6fd0c: [AMDGPU] Autogenerate & pre-commit tests for D156301 and D157388.Aug 18 2023, 6:51 AM
arsenm added inline comments.Aug 23 2023, 4:41 PM
llvm/test/CodeGen/AMDGPU/global_atomic_optimizer_fp_rtn.ll
809 ↗(On Diff #551428)

This broke the strictfp handling (I thought this was supposed to fail the verifier now?)

You probably need something like

Builder.setIsFPConstrained(
    RMWI->getFunction()->hasFnAttribute(Attribute::StrictFP));
pravinjagtap updated this revision to Diff 553027.Aug 24 2023, 1:04 AM

Rebased agaist updated dependancies.

Harbormaster completed remote builds in B254565: Diff 553027.Aug 24 2023, 1:44 AM
arsenm accepted this revision.Aug 29 2023, 3:09 PM
arsenm added inline comments.
llvm/test/CodeGen/AMDGPU/global_atomic_optimizer_fp_rtn.ll
631–633 ↗(On Diff #553027)

The canonical way to do this extract in the IR is trunc and trunc (lshr x, 32)

This revision is now accepted and ready to land.Aug 29 2023, 3:09 PM
pravinjagtap updated this revision to Diff 554644.Aug 30 2023, 2:31 AM

Addressed review comments

Harbormaster completed remote builds in B255742: Diff 554644.Aug 30 2023, 5:56 AM
This revision was landed with ongoing or failed builds.Aug 30 2023, 9:11 AM
Closed by commit rGedb9fab39022: [AMDGPU] Support FMin/FMax in AMDGPUAtomicOptimizer. (authored by pravinjagtap). · Explain Why
This revision was automatically updated to reflect the committed changes.
pravinjagtap added a commit: rGedb9fab39022: [AMDGPU] Support FMin/FMax in AMDGPUAtomicOptimizer..

Revision Contents

PathSize
llvm/
lib/
Target/
AMDGPU/
21 lines
test/
CodeGen/
AMDGPU/
276 lines

Diff 548170

llvm/lib/Target/AMDGPU/AMDGPUAtomicOptimizer.cpp

Show First 20 Lines • Show All 201 Lines • ▼ Show 20 Linesvoid AMDGPUAtomicOptimizerImpl::visitAtomicRMWInst(AtomicRMWInst &I) {
case AtomicRMWInst::Or: case AtomicRMWInst::Or:
case AtomicRMWInst::Xor: case AtomicRMWInst::Xor:
case AtomicRMWInst::Max: case AtomicRMWInst::Max:
case AtomicRMWInst::Min: case AtomicRMWInst::Min:
case AtomicRMWInst::UMax: case AtomicRMWInst::UMax:
case AtomicRMWInst::UMin: case AtomicRMWInst::UMin:
case AtomicRMWInst::FAdd: case AtomicRMWInst::FAdd:
case AtomicRMWInst::FSub: case AtomicRMWInst::FSub:
case AtomicRMWInst::FMax:
case AtomicRMWInst::FMin:
break; break;
} }
// Only 32-bit floating point atomic ops are supported. // Only 32-bit floating point atomic ops are supported.
if (AtomicRMWInst::isFPOperation(Op) && !I.getType()->isFloatTy()) { if (AtomicRMWInst::isFPOperation(Op) && !I.getType()->isFloatTy()) {
return; return;
} }
▲ Show 20 Lines • Show All 92 Lines • ▼ Show 20 Linesvoid AMDGPUAtomicOptimizerImpl::visitIntrinsicInst(IntrinsicInst &I) {
case Intrinsic::amdgcn_struct_ptr_buffer_atomic_umax: case Intrinsic::amdgcn_struct_ptr_buffer_atomic_umax:
case Intrinsic::amdgcn_raw_buffer_atomic_umax: case Intrinsic::amdgcn_raw_buffer_atomic_umax:
case Intrinsic::amdgcn_raw_ptr_buffer_atomic_umax: case Intrinsic::amdgcn_raw_ptr_buffer_atomic_umax:
Op = AtomicRMWInst::UMax; Op = AtomicRMWInst::UMax;
break; break;
case Intrinsic::amdgcn_global_atomic_fadd: case Intrinsic::amdgcn_global_atomic_fadd:
Op = AtomicRMWInst::FAdd; Op = AtomicRMWInst::FAdd;
break; break;
case Intrinsic::amdgcn_global_atomic_fmax:
arsenmUnsubmitted
Not Done
ReplyInline Actions

I don't want to implicitly convert the intrinsics here. We should move towards getting rid of the intrinsics and autoupgrading them

arsenm: I don't want to implicitly convert the intrinsics here. We should move towards getting rid of…
Op = AtomicRMWInst::FMax;
break;
case Intrinsic::amdgcn_global_atomic_fmin:
Op = AtomicRMWInst::FMin;
break;
} }
// Only 32-bit floating point atomic ops are supported. // Only 32-bit floating point atomic ops are supported.
if (AtomicRMWInst::isFPOperation(Op) && !I.getType()->isFloatTy()) { if (AtomicRMWInst::isFPOperation(Op) && !I.getType()->isFloatTy()) {
return; return;
} }
unsigned ValIdx = 0; unsigned ValIdx = 0;
// TODO: Operand order is not consistent for atomic fadd intrinsics // TODO: Operand order is not consistent for atomic fadd intrinsics
if (Op == AtomicRMWInst::FAdd) { if (Op == AtomicRMWInst::FAdd || Op == AtomicRMWInst::FMax ||
Op == AtomicRMWInst::FMin) {
ValIdx = 1; ValIdx = 1;
} }
const bool ValDivergent = UA->isDivergentUse(I.getOperandUse(ValIdx)); const bool ValDivergent = UA->isDivergentUse(I.getOperandUse(ValIdx));
// If the value operand is divergent, each lane is contributing a different // If the value operand is divergent, each lane is contributing a different
// value to the atomic calculation. We can only optimize divergent values if // value to the atomic calculation. We can only optimize divergent values if
// we have DPP available on our subtarget, and the atomic operation is 32 // we have DPP available on our subtarget, and the atomic operation is 32
▲ Show 20 Lines • Show All 50 Lines • ▼ Show 20 Lines case AtomicRMWInst::Min:
Pred = CmpInst::ICMP_SLT; Pred = CmpInst::ICMP_SLT;
break; break;
case AtomicRMWInst::UMax: case AtomicRMWInst::UMax:
Pred = CmpInst::ICMP_UGT; Pred = CmpInst::ICMP_UGT;
break; break;
case AtomicRMWInst::UMin: case AtomicRMWInst::UMin:
Pred = CmpInst::ICMP_ULT; Pred = CmpInst::ICMP_ULT;
break; break;
case AtomicRMWInst::FMax:
return B.CreateSelect(B.CreateFCmp(FCmpInst::FCMP_UGT, LHS, RHS), LHS, RHS);
pravinjagtapAuthorUnsubmitted
Done
ReplyInline Actions

@arsenm you earlier suggested to use minnum/maxnum intrinsics for this. This also seems to give correct behavior. I am not sure what I am missing here

pravinjagtap: @arsenm you earlier suggested to use minnum/maxnum intrinsics for this. This also seems to give…
arsenmUnsubmitted
Not Done
ReplyInline Actions

This is incorrect, you should create minnum/maxnum

arsenm: This is incorrect, you should create minnum/maxnum
arsenmUnsubmitted
Not Done
ReplyInline Actions

Yes, it is wrong to use fcmp and select here. For example for fmax what you have returns the wrong result if LHS is a nan.

select (ugt nan, rhs), nan, rhs -> nan
maxnum(nan, rhs) -> rhs

arsenm: Yes, it is wrong to use fcmp and select here. For example for fmax what you have returns the…
pravinjagtapAuthorUnsubmitted
Done
ReplyInline Actions

you should create minnum/maxnum

Are you referring to @llvm.amdgcn.fcmp.f32(float, float, i32) intrinsic here right ?

pravinjagtap: >you should create minnum/maxnum Are you referring to `@llvm.amdgcn.fcmp.f32(float, float…
arsenmUnsubmitted
Not Done
ReplyInline Actions

No, IRBuilder.CreateMinNum and CreateMaxNum

arsenm: No, IRBuilder.CreateMinNum and CreateMaxNum
case AtomicRMWInst::FMin:
return B.CreateSelect(B.CreateFCmp(FCmpInst::FCMP_ULT, LHS, RHS), LHS, RHS);
} }
Value *Cond = B.CreateICmp(Pred, LHS, RHS); Value *Cond = B.CreateICmp(Pred, LHS, RHS);
return B.CreateSelect(Cond, LHS, RHS); return B.CreateSelect(Cond, LHS, RHS);
} }
// Use the builder to create a reduction of V across the wavefront, with all // Use the builder to create a reduction of V across the wavefront, with all
// lanes active, returning the same result in all lanes. // lanes active, returning the same result in all lanes.
Value *AMDGPUAtomicOptimizerImpl::buildReduction( Value *AMDGPUAtomicOptimizerImpl::buildReduction(
▲ Show 20 Lines • Show All 304 Lines • ▼ Show 20 Linesstatic APFloat getIdentityValueForFAtomicOp(AtomicRMWInst::BinOp Op,
const fltSemantics &Semantics) { const fltSemantics &Semantics) {
switch (Op) { switch (Op) {
default: default:
llvm_unreachable("Unhandled atomic op"); llvm_unreachable("Unhandled atomic op");
case AtomicRMWInst::FAdd: case AtomicRMWInst::FAdd:
return APFloat::getZero(Semantics, false); return APFloat::getZero(Semantics, false);
case AtomicRMWInst::FSub: case AtomicRMWInst::FSub:
return APFloat::getZero(Semantics, true); return APFloat::getZero(Semantics, true);
case AtomicRMWInst::FMin:
return APFloat::getInf(Semantics, false);
case AtomicRMWInst::FMax:
return APFloat::getInf(Semantics, true);
} }
} }
static APInt getIdentityValueForAtomicOp(AtomicRMWInst::BinOp Op, static APInt getIdentityValueForAtomicOp(AtomicRMWInst::BinOp Op,
unsigned BitWidth) { unsigned BitWidth) {
switch (Op) { switch (Op) {
default: default:
llvm_unreachable("Unhandled atomic op"); llvm_unreachable("Unhandled atomic op");
▲ Show 20 Lines • Show All 178 Lines • ▼ Show 20 Lines case AtomicRMWInst::FSub: {
break; break;
} }
case AtomicRMWInst::And: case AtomicRMWInst::And:
case AtomicRMWInst::Or: case AtomicRMWInst::Or:
case AtomicRMWInst::Max: case AtomicRMWInst::Max:
case AtomicRMWInst::Min: case AtomicRMWInst::Min:
case AtomicRMWInst::UMax: case AtomicRMWInst::UMax:
case AtomicRMWInst::UMin: case AtomicRMWInst::UMin:
case AtomicRMWInst::FMin:
case AtomicRMWInst::FMax:
// These operations with a uniform value are idempotent: doing the atomic // These operations with a uniform value are idempotent: doing the atomic
// operation multiple times has the same effect as doing it once. // operation multiple times has the same effect as doing it once.
NewV = V; NewV = V;
break; break;
case AtomicRMWInst::Xor: case AtomicRMWInst::Xor:
// The new value we will be contributing to the atomic operation is the // The new value we will be contributing to the atomic operation is the
// old value times the parity of the number of active lanes. // old value times the parity of the number of active lanes.
▲ Show 20 Lines • Show All 173 Lines • Show Last 20 Lines

llvm/test/CodeGen/AMDGPU/global_atomics_iterative_scan_fp.ll

Show First 20 Lines • Show All 276 Lines • ▼ Show 20 Lines
; IR-DPP-NEXT: ret void ; IR-DPP-NEXT: ret void
; ;
%id.x = call i32 @llvm.amdgcn.workitem.id.x() %id.x = call i32 @llvm.amdgcn.workitem.id.x()
%divValue = bitcast i32 %id.x to float %divValue = bitcast i32 %id.x to float
%result = atomicrmw fsub ptr addrspace(1) %ptr, float %divValue seq_cst %result = atomicrmw fsub ptr addrspace(1) %ptr, float %divValue seq_cst
ret void ret void
} }
define amdgpu_kernel void @global_atomic_fmin_uni_value(ptr addrspace(1) %ptr) #0 {
; IR-ITERATIVE-LABEL: @global_atomic_fmin_uni_value(
; IR-ITERATIVE-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-ITERATIVE-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-ITERATIVE-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-ITERATIVE-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-ITERATIVE-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-ITERATIVE-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-ITERATIVE-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-ITERATIVE-NEXT: br i1 [[TMP7]], label [[TMP8:%.*]], label [[TMP10:%.*]]
; IR-ITERATIVE: 8:
; IR-ITERATIVE-NEXT: [[TMP9:%.*]] = atomicrmw fmin ptr addrspace(1) [[PTR:%.*]], float 4.000000e+00 seq_cst, align 4
; IR-ITERATIVE-NEXT: br label [[TMP10]]
; IR-ITERATIVE: 10:
; IR-ITERATIVE-NEXT: ret void
;
; IR-DPP-LABEL: @global_atomic_fmin_uni_value(
; IR-DPP-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-DPP-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-DPP-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-DPP-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-DPP-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-DPP-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-DPP-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-DPP-NEXT: br i1 [[TMP7]], label [[TMP8:%.*]], label [[TMP10:%.*]]
; IR-DPP: 8:
; IR-DPP-NEXT: [[TMP9:%.*]] = atomicrmw fmin ptr addrspace(1) [[PTR:%.*]], float 4.000000e+00 seq_cst, align 4
; IR-DPP-NEXT: br label [[TMP10]]
; IR-DPP: 10:
; IR-DPP-NEXT: ret void
;
%result = atomicrmw fmin ptr addrspace(1) %ptr, float 4.0 seq_cst
ret void
}
define amdgpu_kernel void @global_atomic_fmin_div_value(ptr addrspace(1) %ptr) #0 {
; IR-ITERATIVE-LABEL: @global_atomic_fmin_div_value(
; IR-ITERATIVE-NEXT: [[ID_X:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
; IR-ITERATIVE-NEXT: [[DIVVALUE:%.*]] = bitcast i32 [[ID_X]] to float
; IR-ITERATIVE-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-ITERATIVE-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-ITERATIVE-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-ITERATIVE-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-ITERATIVE-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-ITERATIVE-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-ITERATIVE-NEXT: [[TMP7:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-ITERATIVE-NEXT: br label [[COMPUTELOOP:%.*]]
; IR-ITERATIVE: 8:
; IR-ITERATIVE-NEXT: [[TMP9:%.*]] = atomicrmw fmin ptr addrspace(1) [[PTR:%.*]], float [[TMP17:%.*]] seq_cst, align 4
; IR-ITERATIVE-NEXT: br label [[TMP10:%.*]]
; IR-ITERATIVE: 10:
; IR-ITERATIVE-NEXT: ret void
; IR-ITERATIVE: ComputeLoop:
; IR-ITERATIVE-NEXT: [[ACCUMULATOR:%.*]] = phi float [ 0x7FF0000000000000, [[TMP0:%.*]] ], [ [[TMP17]], [[COMPUTELOOP]] ]
; IR-ITERATIVE-NEXT: [[ACTIVEBITS:%.*]] = phi i64 [ [[TMP7]], [[TMP0]] ], [ [[TMP20:%.*]], [[COMPUTELOOP]] ]
; IR-ITERATIVE-NEXT: [[TMP11:%.*]] = call i64 @llvm.cttz.i64(i64 [[ACTIVEBITS]], i1 true)
; IR-ITERATIVE-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
; IR-ITERATIVE-NEXT: [[TMP13:%.*]] = bitcast float [[DIVVALUE]] to i32
; IR-ITERATIVE-NEXT: [[TMP14:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[TMP13]], i32 [[TMP12]])
; IR-ITERATIVE-NEXT: [[TMP15:%.*]] = bitcast i32 [[TMP14]] to float
; IR-ITERATIVE-NEXT: [[TMP16:%.*]] = fcmp ult float [[ACCUMULATOR]], [[TMP15]]
; IR-ITERATIVE-NEXT: [[TMP17]] = select i1 [[TMP16]], float [[ACCUMULATOR]], float [[TMP15]]
; IR-ITERATIVE-NEXT: [[TMP18:%.*]] = shl i64 1, [[TMP11]]
; IR-ITERATIVE-NEXT: [[TMP19:%.*]] = xor i64 [[TMP18]], -1
; IR-ITERATIVE-NEXT: [[TMP20]] = and i64 [[ACTIVEBITS]], [[TMP19]]
; IR-ITERATIVE-NEXT: [[TMP21:%.*]] = icmp eq i64 [[TMP20]], 0
; IR-ITERATIVE-NEXT: br i1 [[TMP21]], label [[COMPUTEEND:%.*]], label [[COMPUTELOOP]]
; IR-ITERATIVE: ComputeEnd:
; IR-ITERATIVE-NEXT: [[TMP22:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-ITERATIVE-NEXT: br i1 [[TMP22]], label [[TMP8:%.*]], label [[TMP10]]
;
; IR-DPP-LABEL: @global_atomic_fmin_div_value(
; IR-DPP-NEXT: [[ID_X:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
; IR-DPP-NEXT: [[DIVVALUE:%.*]] = bitcast i32 [[ID_X]] to float
; IR-DPP-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-DPP-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-DPP-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-DPP-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-DPP-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-DPP-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-DPP-NEXT: [[TMP7:%.*]] = bitcast float [[DIVVALUE]] to i32
; IR-DPP-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.set.inactive.i32(i32 [[TMP7]], i32 2139095040)
; IR-DPP-NEXT: [[TMP9:%.*]] = bitcast i32 [[TMP8]] to float
; IR-DPP-NEXT: [[TMP10:%.*]] = bitcast i32 [[TMP7]] to float
; IR-DPP-NEXT: [[TMP11:%.*]] = bitcast float [[TMP9]] to i32
; IR-DPP-NEXT: [[TMP12:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 2139095040, i32 [[TMP11]], i32 273, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP13:%.*]] = bitcast i32 [[TMP12]] to float
; IR-DPP-NEXT: [[TMP14:%.*]] = bitcast i32 [[TMP11]] to float
; IR-DPP-NEXT: [[TMP15:%.*]] = fcmp ult float [[TMP14]], [[TMP13]]
; IR-DPP-NEXT: [[TMP16:%.*]] = select i1 [[TMP15]], float [[TMP14]], float [[TMP13]]
; IR-DPP-NEXT: [[TMP17:%.*]] = bitcast float [[TMP16]] to i32
; IR-DPP-NEXT: [[TMP18:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 2139095040, i32 [[TMP17]], i32 274, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP19:%.*]] = bitcast i32 [[TMP18]] to float
; IR-DPP-NEXT: [[TMP20:%.*]] = bitcast i32 [[TMP17]] to float
; IR-DPP-NEXT: [[TMP21:%.*]] = fcmp ult float [[TMP20]], [[TMP19]]
; IR-DPP-NEXT: [[TMP22:%.*]] = select i1 [[TMP21]], float [[TMP20]], float [[TMP19]]
; IR-DPP-NEXT: [[TMP23:%.*]] = bitcast float [[TMP22]] to i32
; IR-DPP-NEXT: [[TMP24:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 2139095040, i32 [[TMP23]], i32 276, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP25:%.*]] = bitcast i32 [[TMP24]] to float
; IR-DPP-NEXT: [[TMP26:%.*]] = bitcast i32 [[TMP23]] to float
; IR-DPP-NEXT: [[TMP27:%.*]] = fcmp ult float [[TMP26]], [[TMP25]]
; IR-DPP-NEXT: [[TMP28:%.*]] = select i1 [[TMP27]], float [[TMP26]], float [[TMP25]]
; IR-DPP-NEXT: [[TMP29:%.*]] = bitcast float [[TMP28]] to i32
; IR-DPP-NEXT: [[TMP30:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 2139095040, i32 [[TMP29]], i32 280, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP31:%.*]] = bitcast i32 [[TMP30]] to float
; IR-DPP-NEXT: [[TMP32:%.*]] = bitcast i32 [[TMP29]] to float
; IR-DPP-NEXT: [[TMP33:%.*]] = fcmp ult float [[TMP32]], [[TMP31]]
; IR-DPP-NEXT: [[TMP34:%.*]] = select i1 [[TMP33]], float [[TMP32]], float [[TMP31]]
; IR-DPP-NEXT: [[TMP35:%.*]] = bitcast float [[TMP34]] to i32
; IR-DPP-NEXT: [[TMP36:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 2139095040, i32 [[TMP35]], i32 322, i32 10, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP37:%.*]] = bitcast i32 [[TMP36]] to float
; IR-DPP-NEXT: [[TMP38:%.*]] = bitcast i32 [[TMP35]] to float
; IR-DPP-NEXT: [[TMP39:%.*]] = fcmp ult float [[TMP38]], [[TMP37]]
; IR-DPP-NEXT: [[TMP40:%.*]] = select i1 [[TMP39]], float [[TMP38]], float [[TMP37]]
; IR-DPP-NEXT: [[TMP41:%.*]] = bitcast float [[TMP40]] to i32
; IR-DPP-NEXT: [[TMP42:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 2139095040, i32 [[TMP41]], i32 323, i32 12, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP43:%.*]] = bitcast i32 [[TMP42]] to float
; IR-DPP-NEXT: [[TMP44:%.*]] = bitcast i32 [[TMP41]] to float
; IR-DPP-NEXT: [[TMP45:%.*]] = fcmp ult float [[TMP44]], [[TMP43]]
; IR-DPP-NEXT: [[TMP46:%.*]] = select i1 [[TMP45]], float [[TMP44]], float [[TMP43]]
; IR-DPP-NEXT: [[TMP47:%.*]] = bitcast float [[TMP46]] to i32
; IR-DPP-NEXT: [[TMP48:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[TMP47]], i32 63)
; IR-DPP-NEXT: [[TMP49:%.*]] = bitcast i32 [[TMP48]] to float
; IR-DPP-NEXT: [[TMP50:%.*]] = call float @llvm.amdgcn.strict.wwm.f32(float [[TMP49]])
; IR-DPP-NEXT: [[TMP51:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-DPP-NEXT: br i1 [[TMP51]], label [[TMP52:%.*]], label [[TMP54:%.*]]
; IR-DPP: 52:
; IR-DPP-NEXT: [[TMP53:%.*]] = atomicrmw fmin ptr addrspace(1) [[PTR:%.*]], float [[TMP50]] seq_cst, align 4
; IR-DPP-NEXT: br label [[TMP54]]
; IR-DPP: 54:
; IR-DPP-NEXT: ret void
;
%id.x = call i32 @llvm.amdgcn.workitem.id.x()
%divValue = bitcast i32 %id.x to float
%result = atomicrmw fmin ptr addrspace(1) %ptr, float %divValue seq_cst
ret void
}
define amdgpu_kernel void @global_atomic_fmax_uni_value(ptr addrspace(1) %ptr) #0 {
; IR-ITERATIVE-LABEL: @global_atomic_fmax_uni_value(
; IR-ITERATIVE-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-ITERATIVE-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-ITERATIVE-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-ITERATIVE-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-ITERATIVE-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-ITERATIVE-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-ITERATIVE-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-ITERATIVE-NEXT: br i1 [[TMP7]], label [[TMP8:%.*]], label [[TMP10:%.*]]
; IR-ITERATIVE: 8:
; IR-ITERATIVE-NEXT: [[TMP9:%.*]] = atomicrmw fmax ptr addrspace(1) [[PTR:%.*]], float 4.000000e+00 seq_cst, align 4
; IR-ITERATIVE-NEXT: br label [[TMP10]]
; IR-ITERATIVE: 10:
; IR-ITERATIVE-NEXT: ret void
;
; IR-DPP-LABEL: @global_atomic_fmax_uni_value(
; IR-DPP-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-DPP-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-DPP-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-DPP-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-DPP-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-DPP-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-DPP-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-DPP-NEXT: br i1 [[TMP7]], label [[TMP8:%.*]], label [[TMP10:%.*]]
; IR-DPP: 8:
; IR-DPP-NEXT: [[TMP9:%.*]] = atomicrmw fmax ptr addrspace(1) [[PTR:%.*]], float 4.000000e+00 seq_cst, align 4
; IR-DPP-NEXT: br label [[TMP10]]
; IR-DPP: 10:
; IR-DPP-NEXT: ret void
;
%result = atomicrmw fmax ptr addrspace(1) %ptr, float 4.0 seq_cst
ret void
}
define amdgpu_kernel void @global_atomic_fmax_div_value(ptr addrspace(1) %ptr) #0 {
; IR-ITERATIVE-LABEL: @global_atomic_fmax_div_value(
; IR-ITERATIVE-NEXT: [[ID_X:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
; IR-ITERATIVE-NEXT: [[DIVVALUE:%.*]] = bitcast i32 [[ID_X]] to float
; IR-ITERATIVE-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-ITERATIVE-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-ITERATIVE-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-ITERATIVE-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-ITERATIVE-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-ITERATIVE-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-ITERATIVE-NEXT: [[TMP7:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-ITERATIVE-NEXT: br label [[COMPUTELOOP:%.*]]
; IR-ITERATIVE: 8:
; IR-ITERATIVE-NEXT: [[TMP9:%.*]] = atomicrmw fmax ptr addrspace(1) [[PTR:%.*]], float [[TMP17:%.*]] seq_cst, align 4
; IR-ITERATIVE-NEXT: br label [[TMP10:%.*]]
; IR-ITERATIVE: 10:
; IR-ITERATIVE-NEXT: ret void
; IR-ITERATIVE: ComputeLoop:
; IR-ITERATIVE-NEXT: [[ACCUMULATOR:%.*]] = phi float [ 0xFFF0000000000000, [[TMP0:%.*]] ], [ [[TMP17]], [[COMPUTELOOP]] ]
; IR-ITERATIVE-NEXT: [[ACTIVEBITS:%.*]] = phi i64 [ [[TMP7]], [[TMP0]] ], [ [[TMP20:%.*]], [[COMPUTELOOP]] ]
; IR-ITERATIVE-NEXT: [[TMP11:%.*]] = call i64 @llvm.cttz.i64(i64 [[ACTIVEBITS]], i1 true)
; IR-ITERATIVE-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
; IR-ITERATIVE-NEXT: [[TMP13:%.*]] = bitcast float [[DIVVALUE]] to i32
; IR-ITERATIVE-NEXT: [[TMP14:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[TMP13]], i32 [[TMP12]])
; IR-ITERATIVE-NEXT: [[TMP15:%.*]] = bitcast i32 [[TMP14]] to float
; IR-ITERATIVE-NEXT: [[TMP16:%.*]] = fcmp ugt float [[ACCUMULATOR]], [[TMP15]]
; IR-ITERATIVE-NEXT: [[TMP17]] = select i1 [[TMP16]], float [[ACCUMULATOR]], float [[TMP15]]
; IR-ITERATIVE-NEXT: [[TMP18:%.*]] = shl i64 1, [[TMP11]]
; IR-ITERATIVE-NEXT: [[TMP19:%.*]] = xor i64 [[TMP18]], -1
; IR-ITERATIVE-NEXT: [[TMP20]] = and i64 [[ACTIVEBITS]], [[TMP19]]
; IR-ITERATIVE-NEXT: [[TMP21:%.*]] = icmp eq i64 [[TMP20]], 0
; IR-ITERATIVE-NEXT: br i1 [[TMP21]], label [[COMPUTEEND:%.*]], label [[COMPUTELOOP]]
; IR-ITERATIVE: ComputeEnd:
; IR-ITERATIVE-NEXT: [[TMP22:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-ITERATIVE-NEXT: br i1 [[TMP22]], label [[TMP8:%.*]], label [[TMP10]]
;
; IR-DPP-LABEL: @global_atomic_fmax_div_value(
; IR-DPP-NEXT: [[ID_X:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
; IR-DPP-NEXT: [[DIVVALUE:%.*]] = bitcast i32 [[ID_X]] to float
; IR-DPP-NEXT: [[TMP1:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
; IR-DPP-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to <2 x i32>
; IR-DPP-NEXT: [[TMP3:%.*]] = extractelement <2 x i32> [[TMP2]], i32 0
; IR-DPP-NEXT: [[TMP4:%.*]] = extractelement <2 x i32> [[TMP2]], i32 1
; IR-DPP-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.mbcnt.lo(i32 [[TMP3]], i32 0)
; IR-DPP-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.mbcnt.hi(i32 [[TMP4]], i32 [[TMP5]])
; IR-DPP-NEXT: [[TMP7:%.*]] = bitcast float [[DIVVALUE]] to i32
; IR-DPP-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.set.inactive.i32(i32 [[TMP7]], i32 -8388608)
; IR-DPP-NEXT: [[TMP9:%.*]] = bitcast i32 [[TMP8]] to float
; IR-DPP-NEXT: [[TMP10:%.*]] = bitcast i32 [[TMP7]] to float
; IR-DPP-NEXT: [[TMP11:%.*]] = bitcast float [[TMP9]] to i32
; IR-DPP-NEXT: [[TMP12:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 -8388608, i32 [[TMP11]], i32 273, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP13:%.*]] = bitcast i32 [[TMP12]] to float
; IR-DPP-NEXT: [[TMP14:%.*]] = bitcast i32 [[TMP11]] to float
; IR-DPP-NEXT: [[TMP15:%.*]] = fcmp ugt float [[TMP14]], [[TMP13]]
; IR-DPP-NEXT: [[TMP16:%.*]] = select i1 [[TMP15]], float [[TMP14]], float [[TMP13]]
; IR-DPP-NEXT: [[TMP17:%.*]] = bitcast float [[TMP16]] to i32
; IR-DPP-NEXT: [[TMP18:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 -8388608, i32 [[TMP17]], i32 274, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP19:%.*]] = bitcast i32 [[TMP18]] to float
; IR-DPP-NEXT: [[TMP20:%.*]] = bitcast i32 [[TMP17]] to float
; IR-DPP-NEXT: [[TMP21:%.*]] = fcmp ugt float [[TMP20]], [[TMP19]]
; IR-DPP-NEXT: [[TMP22:%.*]] = select i1 [[TMP21]], float [[TMP20]], float [[TMP19]]
; IR-DPP-NEXT: [[TMP23:%.*]] = bitcast float [[TMP22]] to i32
; IR-DPP-NEXT: [[TMP24:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 -8388608, i32 [[TMP23]], i32 276, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP25:%.*]] = bitcast i32 [[TMP24]] to float
; IR-DPP-NEXT: [[TMP26:%.*]] = bitcast i32 [[TMP23]] to float
; IR-DPP-NEXT: [[TMP27:%.*]] = fcmp ugt float [[TMP26]], [[TMP25]]
; IR-DPP-NEXT: [[TMP28:%.*]] = select i1 [[TMP27]], float [[TMP26]], float [[TMP25]]
; IR-DPP-NEXT: [[TMP29:%.*]] = bitcast float [[TMP28]] to i32
; IR-DPP-NEXT: [[TMP30:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 -8388608, i32 [[TMP29]], i32 280, i32 15, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP31:%.*]] = bitcast i32 [[TMP30]] to float
; IR-DPP-NEXT: [[TMP32:%.*]] = bitcast i32 [[TMP29]] to float
; IR-DPP-NEXT: [[TMP33:%.*]] = fcmp ugt float [[TMP32]], [[TMP31]]
; IR-DPP-NEXT: [[TMP34:%.*]] = select i1 [[TMP33]], float [[TMP32]], float [[TMP31]]
; IR-DPP-NEXT: [[TMP35:%.*]] = bitcast float [[TMP34]] to i32
; IR-DPP-NEXT: [[TMP36:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 -8388608, i32 [[TMP35]], i32 322, i32 10, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP37:%.*]] = bitcast i32 [[TMP36]] to float
; IR-DPP-NEXT: [[TMP38:%.*]] = bitcast i32 [[TMP35]] to float
; IR-DPP-NEXT: [[TMP39:%.*]] = fcmp ugt float [[TMP38]], [[TMP37]]
; IR-DPP-NEXT: [[TMP40:%.*]] = select i1 [[TMP39]], float [[TMP38]], float [[TMP37]]
; IR-DPP-NEXT: [[TMP41:%.*]] = bitcast float [[TMP40]] to i32
; IR-DPP-NEXT: [[TMP42:%.*]] = call i32 @llvm.amdgcn.update.dpp.i32(i32 -8388608, i32 [[TMP41]], i32 323, i32 12, i32 15, i1 false)
; IR-DPP-NEXT: [[TMP43:%.*]] = bitcast i32 [[TMP42]] to float
; IR-DPP-NEXT: [[TMP44:%.*]] = bitcast i32 [[TMP41]] to float
; IR-DPP-NEXT: [[TMP45:%.*]] = fcmp ugt float [[TMP44]], [[TMP43]]
; IR-DPP-NEXT: [[TMP46:%.*]] = select i1 [[TMP45]], float [[TMP44]], float [[TMP43]]
; IR-DPP-NEXT: [[TMP47:%.*]] = bitcast float [[TMP46]] to i32
arsenmUnsubmitted
Not Done
ReplyInline Actions

Should also test with different scopes, at least default system and agent

arsenm: Should also test with different scopes, at least default system and agent
; IR-DPP-NEXT: [[TMP48:%.*]] = call i32 @llvm.amdgcn.readlane(i32 [[TMP47]], i32 63)
; IR-DPP-NEXT: [[TMP49:%.*]] = bitcast i32 [[TMP48]] to float
; IR-DPP-NEXT: [[TMP50:%.*]] = call float @llvm.amdgcn.strict.wwm.f32(float [[TMP49]])
; IR-DPP-NEXT: [[TMP51:%.*]] = icmp eq i32 [[TMP6]], 0
; IR-DPP-NEXT: br i1 [[TMP51]], label [[TMP52:%.*]], label [[TMP54:%.*]]
; IR-DPP: 52:
; IR-DPP-NEXT: [[TMP53:%.*]] = atomicrmw fmax ptr addrspace(1) [[PTR:%.*]], float [[TMP50]] seq_cst, align 4
; IR-DPP-NEXT: br label [[TMP54]]
; IR-DPP: 54:
; IR-DPP-NEXT: ret void
;
%id.x = call i32 @llvm.amdgcn.workitem.id.x()
%divValue = bitcast i32 %id.x to float
%result = atomicrmw fmax ptr addrspace(1) %ptr, float %divValue seq_cst
ret void
}
attributes #0 = {"target-cpu"="gfx906"} attributes #0 = {"target-cpu"="gfx906"}
arsenmUnsubmitted
Not Done
ReplyInline Actions

Should take this from the command line

arsenm: Should take this from the command line