[llvm] r323516 - [AMDGPU] fix LDS f32 intrinsics
Daniil Fukalov via llvm-commits
llvm-commits at lists.llvm.org
Fri Jan 26 03:09:38 PST 2018
Author: dfukalov
Date: Fri Jan 26 03:09:38 2018
New Revision: 323516
URL: http://llvm.org/viewvc/llvm-project?rev=323516&view=rev
Log:
[AMDGPU] fix LDS f32 intrinsics
- using qualified pointer addrspace in intrinsics class to avoid .f32 mangling
- changed too common atomic mangling to ds
- added missing intrinsics to AMDGPUTTIImpl::getTgtMemIntrinsic
Reviewed by: b-sumner
Differential Revision: https://reviews.llvm.org/D42383
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/trunk/lib/Transforms/Scalar/InferAddressSpaces.cpp
llvm/trunk/test/CodeGen/AMDGPU/lds_atomic_f32.ll
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=323516&r1=323515&r2=323516&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Fri Jan 26 03:09:38 2018
@@ -295,10 +295,10 @@ class AMDGPUAtomicIncIntrin : Intrinsic<
def int_amdgcn_atomic_inc : AMDGPUAtomicIncIntrin;
def int_amdgcn_atomic_dec : AMDGPUAtomicIncIntrin;
-class AMDGPUAtomicF32Intrin<string clang_builtin> :
+class AMDGPULDSF32Intrin<string clang_builtin> :
GCCBuiltin<clang_builtin>,
Intrinsic<[llvm_float_ty],
- [LLVMAnyPointerType<llvm_float_ty>,
+ [LLVMQualPointerType<llvm_float_ty, 3>,
llvm_float_ty,
llvm_i32_ty, // ordering
llvm_i32_ty, // scope
@@ -306,9 +306,9 @@ class AMDGPUAtomicF32Intrin<string clang
[IntrArgMemOnly, NoCapture<0>]
>;
-def int_amdgcn_atomic_fadd : AMDGPUAtomicF32Intrin<"__builtin_amdgcn_ds_fadd">;
-def int_amdgcn_atomic_fmin : AMDGPUAtomicF32Intrin<"__builtin_amdgcn_ds_fmin">;
-def int_amdgcn_atomic_fmax : AMDGPUAtomicF32Intrin<"__builtin_amdgcn_ds_fmax">;
+def int_amdgcn_ds_fadd : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fadd">;
+def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmin">;
+def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmax">;
class AMDGPUImageLoad<bit NoMem = 0> : Intrinsic <
[llvm_anyfloat_ty], // vdata(VGPR)
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp?rev=323516&r1=323515&r2=323516&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp Fri Jan 26 03:09:38 2018
@@ -292,7 +292,10 @@ bool AMDGPUTTIImpl::getTgtMemIntrinsic(I
MemIntrinsicInfo &Info) const {
switch (Inst->getIntrinsicID()) {
case Intrinsic::amdgcn_atomic_inc:
- case Intrinsic::amdgcn_atomic_dec: {
+ case Intrinsic::amdgcn_atomic_dec:
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax: {
auto *Ordering = dyn_cast<ConstantInt>(Inst->getArgOperand(2));
auto *Volatile = dyn_cast<ConstantInt>(Inst->getArgOperand(4));
if (!Ordering || !Volatile)
@@ -475,9 +478,9 @@ static bool isIntrinsicSourceOfDivergenc
case Intrinsic::r600_read_tidig_z:
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
- case Intrinsic::amdgcn_atomic_fadd:
- case Intrinsic::amdgcn_atomic_fmin:
- case Intrinsic::amdgcn_atomic_fmax:
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax:
case Intrinsic::amdgcn_image_atomic_swap:
case Intrinsic::amdgcn_image_atomic_add:
case Intrinsic::amdgcn_image_atomic_sub:
Modified: llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp?rev=323516&r1=323515&r2=323516&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp Fri Jan 26 03:09:38 2018
@@ -566,9 +566,9 @@ bool SITargetLowering::getTgtMemIntrinsi
switch (IntrID) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
- case Intrinsic::amdgcn_atomic_fadd:
- case Intrinsic::amdgcn_atomic_fmin:
- case Intrinsic::amdgcn_atomic_fmax: {
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::getVT(CI.getType());
Info.ptrVal = CI.getOperand(0);
@@ -807,9 +807,9 @@ bool SITargetLowering::getAddrModeArgume
switch (II->getIntrinsicID()) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
- case Intrinsic::amdgcn_atomic_fadd:
- case Intrinsic::amdgcn_atomic_fmin:
- case Intrinsic::amdgcn_atomic_fmax: {
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax: {
Value *Ptr = II->getArgOperand(0);
AccessTy = II->getType();
Ops.push_back(Ptr);
@@ -4827,9 +4827,9 @@ SDValue SITargetLowering::LowerINTRINSIC
switch (IntrID) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
- case Intrinsic::amdgcn_atomic_fadd:
- case Intrinsic::amdgcn_atomic_fmin:
- case Intrinsic::amdgcn_atomic_fmax: {
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax: {
MemSDNode *M = cast<MemSDNode>(Op);
unsigned Opc;
switch (IntrID) {
@@ -4839,13 +4839,13 @@ SDValue SITargetLowering::LowerINTRINSIC
case Intrinsic::amdgcn_atomic_dec:
Opc = AMDGPUISD::ATOMIC_DEC;
break;
- case Intrinsic::amdgcn_atomic_fadd:
+ case Intrinsic::amdgcn_ds_fadd:
Opc = AMDGPUISD::ATOMIC_LOAD_FADD;
break;
- case Intrinsic::amdgcn_atomic_fmin:
+ case Intrinsic::amdgcn_ds_fmin:
Opc = AMDGPUISD::ATOMIC_LOAD_FMIN;
break;
- case Intrinsic::amdgcn_atomic_fmax:
+ case Intrinsic::amdgcn_ds_fmax:
Opc = AMDGPUISD::ATOMIC_LOAD_FMAX;
break;
default:
Modified: llvm/trunk/lib/Transforms/Scalar/InferAddressSpaces.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Transforms/Scalar/InferAddressSpaces.cpp?rev=323516&r1=323515&r2=323516&view=diff
==============================================================================
--- llvm/trunk/lib/Transforms/Scalar/InferAddressSpaces.cpp (original)
+++ llvm/trunk/lib/Transforms/Scalar/InferAddressSpaces.cpp Fri Jan 26 03:09:38 2018
@@ -261,9 +261,9 @@ bool InferAddressSpaces::rewriteIntrinsi
switch (II->getIntrinsicID()) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
- case Intrinsic::amdgcn_atomic_fadd:
- case Intrinsic::amdgcn_atomic_fmin:
- case Intrinsic::amdgcn_atomic_fmax: {
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax: {
const ConstantInt *IsVolatile = dyn_cast<ConstantInt>(II->getArgOperand(4));
if (!IsVolatile || !IsVolatile->isZero())
return false;
@@ -292,9 +292,9 @@ void InferAddressSpaces::collectRewritab
case Intrinsic::objectsize:
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
- case Intrinsic::amdgcn_atomic_fadd:
- case Intrinsic::amdgcn_atomic_fmin:
- case Intrinsic::amdgcn_atomic_fmax:
+ case Intrinsic::amdgcn_ds_fadd:
+ case Intrinsic::amdgcn_ds_fmin:
+ case Intrinsic::amdgcn_ds_fmax:
appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(0),
PostorderStack, Visited);
break;
Modified: llvm/trunk/test/CodeGen/AMDGPU/lds_atomic_f32.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/lds_atomic_f32.ll?rev=323516&r1=323515&r2=323516&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/lds_atomic_f32.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/lds_atomic_f32.ll Fri Jan 26 03:09:38 2018
@@ -1,11 +1,11 @@
; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,VI %s
; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX9 %s
-declare float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* nocapture, float, i32, i32, i1)
-declare float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* nocapture, float, i32, i32, i1)
-declare float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* nocapture, float, i32, i32, i1)
+declare float @llvm.amdgcn.ds.fadd(float addrspace(3)* nocapture, float, i32, i32, i1)
+declare float @llvm.amdgcn.ds.fmin(float addrspace(3)* nocapture, float, i32, i32, i1)
+declare float @llvm.amdgcn.ds.fmax(float addrspace(3)* nocapture, float, i32, i32, i1)
-; GCN-LABEL: {{^}}lds_atomic_fadd_f32:
+; GCN-LABEL: {{^}}lds_ds_fadd:
; VI-DAG: s_mov_b32 m0
; GFX9-NOT: m0
; GCN-DAG: v_mov_b32_e32 [[V0:v[0-9]+]], 0x42280000
@@ -13,20 +13,20 @@ declare float @llvm.amdgcn.atomic.fmax.f
; GCN: ds_add_f32 [[V3:v[0-9]+]], [[V0]] offset:64
; GCN: s_waitcnt lgkmcnt(1)
; GCN: ds_add_rtn_f32 {{v[0-9]+}}, {{v[0-9]+}}, [[V2]]
-define amdgpu_kernel void @lds_atomic_fadd_f32(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
+define amdgpu_kernel void @lds_ds_fadd(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
%idx.add = add nuw i32 %idx, 4
%shl0 = shl i32 %idx.add, 3
%shl1 = shl i32 %idx.add, 4
%ptr0 = inttoptr i32 %shl0 to float addrspace(3)*
%ptr1 = inttoptr i32 %shl1 to float addrspace(3)*
- %a1 = call float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
- %a2 = call float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
- %a3 = call float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
+ %a1 = call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
+ %a2 = call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
+ %a3 = call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
store float %a3, float addrspace(1)* %out
ret void
}
-; GCN-LABEL: {{^}}lds_atomic_fmin_f32:
+; GCN-LABEL: {{^}}lds_ds_fmin:
; VI-DAG: s_mov_b32 m0
; GFX9-NOT: m0
; GCN-DAG: v_mov_b32_e32 [[V0:v[0-9]+]], 0x42280000
@@ -34,20 +34,20 @@ define amdgpu_kernel void @lds_atomic_fa
; GCN: ds_min_f32 [[V3:v[0-9]+]], [[V0]] offset:64
; GCN: s_waitcnt lgkmcnt(1)
; GCN: ds_min_rtn_f32 {{v[0-9]+}}, {{v[0-9]+}}, [[V2]]
-define amdgpu_kernel void @lds_atomic_fmin_f32(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
+define amdgpu_kernel void @lds_ds_fmin(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
%idx.add = add nuw i32 %idx, 4
%shl0 = shl i32 %idx.add, 3
%shl1 = shl i32 %idx.add, 4
%ptr0 = inttoptr i32 %shl0 to float addrspace(3)*
%ptr1 = inttoptr i32 %shl1 to float addrspace(3)*
- %a1 = call float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
- %a2 = call float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
- %a3 = call float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
+ %a1 = call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
+ %a2 = call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
+ %a3 = call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
store float %a3, float addrspace(1)* %out
ret void
}
-; GCN-LABEL: {{^}}lds_atomic_fmax_f32:
+; GCN-LABEL: {{^}}lds_ds_fmax:
; VI-DAG: s_mov_b32 m0
; GFX9-NOT: m0
; GCN-DAG: v_mov_b32_e32 [[V0:v[0-9]+]], 0x42280000
@@ -55,15 +55,15 @@ define amdgpu_kernel void @lds_atomic_fm
; GCN: ds_max_f32 [[V3:v[0-9]+]], [[V0]] offset:64
; GCN: s_waitcnt lgkmcnt(1)
; GCN: ds_max_rtn_f32 {{v[0-9]+}}, {{v[0-9]+}}, [[V2]]
-define amdgpu_kernel void @lds_atomic_fmax_f32(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
+define amdgpu_kernel void @lds_ds_fmax(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
%idx.add = add nuw i32 %idx, 4
%shl0 = shl i32 %idx.add, 3
%shl1 = shl i32 %idx.add, 4
%ptr0 = inttoptr i32 %shl0 to float addrspace(3)*
%ptr1 = inttoptr i32 %shl1 to float addrspace(3)*
- %a1 = call float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
- %a2 = call float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
- %a3 = call float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
+ %a1 = call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
+ %a2 = call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
+ %a3 = call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
store float %a3, float addrspace(1)* %out
ret void
}
More information about the llvm-commits
mailing list