[llvm] [DAGCombiner] Add some very basic folds for ADDRSPACECAST (PR #127733)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 18 17:46:54 PST 2025
https://github.com/AlexMaclean created https://github.com/llvm/llvm-project/pull/127733
Fold an ADDRSPACECAST of another ADDRPSACECAST into a single cast or remove them completely and forward the inner ASC operand.
>From 4d532a56291ba9e97b580b198388ac1c47804036 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Wed, 19 Feb 2025 00:20:32 +0000
Subject: [PATCH 1/2] pre-commit tests
---
.../CodeGen/NVPTX/addrspacecast-folding.ll | 37 +++++++++++++++++++
1 file changed, 37 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
diff --git a/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll b/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
new file mode 100644
index 0000000000000..05eb0385eb571
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
@@ -0,0 +1,37 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mcpu=sm_20 -O0 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 -O0 | %ptxas-verify %}
+
+target triple = "nvptx64-unknown-unknown"
+
+define ptr @test1(ptr %p) {
+; CHECK-LABEL: test1(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test1_param_0];
+; CHECK-NEXT: cvta.to.local.u64 %rd2, %rd1;
+; CHECK-NEXT: cvta.local.u64 %rd3, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %a = addrspacecast ptr %p to ptr addrspace(5)
+ %b = addrspacecast ptr addrspace(5) %a to ptr
+ ret ptr %b
+}
+
+define ptr addrspace(1) @test2(ptr addrspace(5) %p) {
+; CHECK-LABEL: test2(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test2_param_0];
+; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;
+; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %a = addrspacecast ptr addrspace(5) %p to ptr
+ %b = addrspacecast ptr %a to ptr addrspace(1)
+ ret ptr addrspace(1) %b
+}
>From 5a3eaf5daa10536fb8334a659d1c95a4659a8158 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Wed, 19 Feb 2025 00:21:55 +0000
Subject: [PATCH 2/2] [DAGCombiner] Add very basic folds for ADDRSPACECAST
---
llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp | 21 +++++++++++++++++++
.../CodeGen/NVPTX/addrspacecast-folding.ll | 11 +++-------
2 files changed, 24 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
index c6fd72b6b76f4..b0af55722f7ff 100644
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -488,6 +488,7 @@ namespace {
SDValue visitTRUNCATE(SDNode *N);
SDValue visitTRUNCATE_USAT_U(SDNode *N);
SDValue visitBITCAST(SDNode *N);
+ SDValue visitADDRSPACECAST(SDNode *N);
SDValue visitFREEZE(SDNode *N);
SDValue visitBUILD_PAIR(SDNode *N);
SDValue visitFADD(SDNode *N);
@@ -1920,6 +1921,7 @@ SDValue DAGCombiner::visit(SDNode *N) {
case ISD::TRUNCATE: return visitTRUNCATE(N);
case ISD::TRUNCATE_USAT_U: return visitTRUNCATE_USAT_U(N);
case ISD::BITCAST: return visitBITCAST(N);
+ case ISD::ADDRSPACECAST: return visitADDRSPACECAST(N);
case ISD::BUILD_PAIR: return visitBUILD_PAIR(N);
case ISD::FADD: return visitFADD(N);
case ISD::STRICT_FADD: return visitSTRICT_FADD(N);
@@ -16054,6 +16056,25 @@ SDValue DAGCombiner::visitBITCAST(SDNode *N) {
return SDValue();
}
+SDValue DAGCombiner::visitADDRSPACECAST(SDNode *N) {
+ auto *ASCN1 = cast<AddrSpaceCastSDNode>(N);
+
+ if (auto *ASCN2 = dyn_cast<AddrSpaceCastSDNode>(ASCN1->getOperand(0))) {
+ assert(ASCN2->getDestAddressSpace() == ASCN1->getSrcAddressSpace());
+
+ // Fold asc[B -> A](asc[A -> B](x)) -> x
+ if (ASCN1->getDestAddressSpace() == ASCN2->getSrcAddressSpace())
+ return ASCN2->getOperand(0);
+
+ // Fold asc[B -> C](asc[A -> B](x)) -> asc[A -> C](x)
+ return DAG.getAddrSpaceCast(
+ SDLoc(N), N->getValueType(0), ASCN2->getOperand(0),
+ ASCN2->getSrcAddressSpace(), ASCN1->getDestAddressSpace());
+ }
+
+ return SDValue();
+}
+
SDValue DAGCombiner::visitBUILD_PAIR(SDNode *N) {
EVT VT = N->getValueType(0);
return CombineConsecutiveLoads(N, VT);
diff --git a/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll b/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
index 05eb0385eb571..11c2b6782e0d3 100644
--- a/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
+++ b/llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
@@ -7,13 +7,11 @@ target triple = "nvptx64-unknown-unknown"
define ptr @test1(ptr %p) {
; CHECK-LABEL: test1(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [test1_param_0];
-; CHECK-NEXT: cvta.to.local.u64 %rd2, %rd1;
-; CHECK-NEXT: cvta.local.u64 %rd3, %rd2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%a = addrspacecast ptr %p to ptr addrspace(5)
%b = addrspacecast ptr addrspace(5) %a to ptr
@@ -23,13 +21,10 @@ define ptr @test1(ptr %p) {
define ptr addrspace(1) @test2(ptr addrspace(5) %p) {
; CHECK-LABEL: test2(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [test2_param_0];
-; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;
-; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%a = addrspacecast ptr addrspace(5) %p to ptr
%b = addrspacecast ptr %a to ptr addrspace(1)
More information about the llvm-commits
mailing list