[llvm] cb24422 - [NVPTX] remove store.params of undef (#96940)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 28 18:00:55 PDT 2024
Author: Alex MacLean
Date: 2024-06-28T18:00:51-07:00
New Revision: cb24422b6f27bbdb9e55d5b3d0c4e8eb7e57d1d1
URL: https://github.com/llvm/llvm-project/commit/cb24422b6f27bbdb9e55d5b3d0c4e8eb7e57d1d1
DIFF: https://github.com/llvm/llvm-project/commit/cb24422b6f27bbdb9e55d5b3d0c4e8eb7e57d1d1.diff
LOG: [NVPTX] remove store.params of undef (#96940)
Extend the logic for `NVPTXISD::StoreRetval` to `NVPTXISD::StoreParam`
to remove stores of undef values.
Added:
llvm/test/CodeGen/NVPTX/store-undef.ll
Modified:
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e3201516e4e7f..476a532db0a37 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5345,15 +5345,26 @@ PerformFADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
return SDValue();
}
+static SDValue PerformStoreCombineHelper(SDNode *N, std::size_t Front,
+ std::size_t Back) {
+ if (all_of(N->ops().drop_front(Front).drop_back(Back),
+ [](const SDUse &U) { return U.get()->isUndef(); }))
+ // Operand 0 is the previous value in the chain. Cannot return EntryToken
+ // as the previous value will become unused and eliminated later.
+ return N->getOperand(0);
+
+ return SDValue();
+}
+
+static SDValue PerformStoreParamCombine(SDNode *N) {
+ // Operands from the 3rd to the 2nd last one are the values to be stored.
+ // {Chain, ArgID, Offset, Val, Glue}
+ return PerformStoreCombineHelper(N, 3, 1);
+}
+
static SDValue PerformStoreRetvalCombine(SDNode *N) {
// Operands from the 2nd to the last one are the values to be stored
- for (std::size_t I = 2, OpsCount = N->ops().size(); I != OpsCount; ++I)
- if (!N->getOperand(I).isUndef())
- return SDValue();
-
- // Operand 0 is the previous value in the chain. Cannot return EntryToken
- // as the previous value will become unused and eliminated later.
- return N->getOperand(0);
+ return PerformStoreCombineHelper(N, 2, 0);
}
/// PerformADDCombine - Target-specific dag combine xforms for ISD::ADD.
@@ -5822,6 +5833,10 @@ static SDValue PerformEXTRACTCombine(SDNode *N,
VectorVT == MVT::v4i8 || VectorVT == MVT::v8i8)
return SDValue();
+ // Don't mess with undef values as sra may be simplified to 0, not undef.
+ if (Vector->isUndef() || ISD::allOperandsUndef(Vector.getNode()))
+ return SDValue();
+
uint64_t VectorBits = VectorVT.getSizeInBits();
// We only handle the types we can extract in-register.
if (!(VectorBits == 16 || VectorBits == 32 || VectorBits == 64))
@@ -5950,6 +5965,10 @@ SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N,
case NVPTXISD::StoreRetvalV2:
case NVPTXISD::StoreRetvalV4:
return PerformStoreRetvalCombine(N);
+ case NVPTXISD::StoreParam:
+ case NVPTXISD::StoreParamV2:
+ case NVPTXISD::StoreParamV4:
+ return PerformStoreParamCombine(N);
case ISD::EXTRACT_VECTOR_ELT:
return PerformEXTRACTCombine(N, DCI);
case ISD::VSELECT:
diff --git a/llvm/test/CodeGen/NVPTX/store-undef.ll b/llvm/test/CodeGen/NVPTX/store-undef.ll
new file mode 100644
index 0000000000000..4941760a78c79
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/store-undef.ll
@@ -0,0 +1,92 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+%struct.T = type { i64, <2 x i32>, <4 x i32> }
+
+declare void @test_call(%struct.T)
+
+define void @test_store_param_undef() {
+; CHECK-LABEL: test_store_param_undef(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: { // callseq 0, 0
+; CHECK-NEXT: .param .align 16 .b8 param0[32];
+; CHECK-NEXT: call.uni
+; CHECK-NEXT: test_call,
+; CHECK-NEXT: (
+; CHECK-NEXT: param0
+; CHECK-NEXT: );
+; CHECK-NEXT: } // callseq 0
+; CHECK-NEXT: ret;
+ call void @test_call(%struct.T undef)
+ ret void
+}
+
+define void @test_store_param_def(i64 %param0, i32 %param1) {
+; CHECK-LABEL: test_store_param_def(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<6>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_store_param_def_param_0];
+; CHECK-NEXT: ld.param.u32 %r1, [test_store_param_def_param_1];
+; CHECK-NEXT: { // callseq 1, 0
+; CHECK-NEXT: .param .align 16 .b8 param0[32];
+; CHECK-NEXT: st.param.b64 [param0+0], %rd1;
+; CHECK-NEXT: st.param.v2.b32 [param0+8], {%r2, %r1};
+; CHECK-NEXT: st.param.v4.b32 [param0+16], {%r3, %r1, %r4, %r5};
+; CHECK-NEXT: call.uni
+; CHECK-NEXT: test_call,
+; CHECK-NEXT: (
+; CHECK-NEXT: param0
+; CHECK-NEXT: );
+; CHECK-NEXT: } // callseq 1
+; CHECK-NEXT: ret;
+ %V2 = insertelement <2 x i32> undef, i32 %param1, i32 1
+ %V4 = insertelement <4 x i32> undef, i32 %param1, i32 1
+ %S0 = insertvalue %struct.T undef, i64 %param0, 0
+ %S1 = insertvalue %struct.T %S0, <2 x i32> %V2, 1
+ %S2 = insertvalue %struct.T %S1, <4 x i32> %V4, 2
+ call void @test_call(%struct.T %S2)
+ ret void
+}
+
+define void @test_store_undef(ptr %out) {
+; CHECK-LABEL: test_store_undef(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ret;
+ store %struct.T undef, ptr %out
+ ret void
+}
+
+define void @test_store_def(i64 %param0, i32 %param1, ptr %out) {
+; CHECK-LABEL: test_store_def(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<6>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_store_def_param_0];
+; CHECK-NEXT: ld.param.u32 %r1, [test_store_def_param_1];
+; CHECK-NEXT: ld.param.u64 %rd2, [test_store_def_param_2];
+; CHECK-NEXT: st.v4.u32 [%rd2+16], {%r2, %r1, %r3, %r4};
+; CHECK-NEXT: st.v2.u32 [%rd2+8], {%r5, %r1};
+; CHECK-NEXT: st.u64 [%rd2], %rd1;
+; CHECK-NEXT: ret;
+ %V2 = insertelement <2 x i32> undef, i32 %param1, i32 1
+ %V4 = insertelement <4 x i32> undef, i32 %param1, i32 1
+ %S0 = insertvalue %struct.T undef, i64 %param0, 0
+ %S1 = insertvalue %struct.T %S0, <2 x i32> %V2, 1
+ %S2 = insertvalue %struct.T %S1, <4 x i32> %V4, 2
+ store %struct.T %S2, ptr %out
+ ret void
+}
More information about the llvm-commits
mailing list