address comments

Author: AlexMaclean

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

@@ -363,15 +363,8 @@ static std::optional<unsigned> convertAS(unsigned AS) {
 }
 
 static unsigned int getCodeAddrSpace(const MemSDNode *N) {
-  if (const Value *Src = N->getMemOperand()->getValue())
-    if (auto *PT = dyn_cast<PointerType>(Src->getType()))
-      if (auto AS = convertAS(PT->getAddressSpace()))
-        return AS.value();
-
-  if (auto AS = convertAS(N->getMemOperand()->getAddrSpace()))
-    return AS.value();
-
-  return NVPTX::AddressSpace::Generic;
+  return convertAS(N->getMemOperand()->getAddrSpace())
+      .value_or(NVPTX::AddressSpace::Generic);
 }
 
 namespace {

llvm/test/CodeGen/NVPTX/variadics-backend.ll

@@ -214,32 +214,34 @@ define dso_local i32 @bar() {
 ; CHECK-PTX-NEXT:    .reg .b64 %SPL;
 ; CHECK-PTX-NEXT:    .reg .b16 %rs<8>;
 ; CHECK-PTX-NEXT:    .reg .b32 %r<4>;
-; CHECK-PTX-NEXT:    .reg .b64 %rd<4>;
+; CHECK-PTX-NEXT:    .reg .b64 %rd<6>;
 ; CHECK-PTX-EMPTY:
 ; CHECK-PTX-NEXT:  // %bb.0: // %entry
 ; CHECK-PTX-NEXT:    mov.u64 %SPL, __local_depot3;
 ; CHECK-PTX-NEXT:    cvta.local.u64 %SP, %SPL;
-; CHECK-PTX-NEXT:    mov.u64 %rd1, __const_$_bar_$_s1;
-; CHECK-PTX-NEXT:    ld.global.nc.u8 %rs1, [%rd1+7];
+; CHECK-PTX-NEXT:    add.u64 %rd2, %SPL, 0;
+; CHECK-PTX-NEXT:    mov.u64 %rd3, __const_$_bar_$_s1;
+; CHECK-PTX-NEXT:    ld.global.nc.u8 %rs1, [%rd3+7];
 ; CHECK-PTX-NEXT:    cvt.u16.u8 %rs2, %rs1;
-; CHECK-PTX-NEXT:    st.u8 [%SP+2], %rs2;
-; CHECK-PTX-NEXT:    ld.global.nc.u8 %rs3, [%rd1+5];
+; CHECK-PTX-NEXT:    st.local.u8 [%rd2+2], %rs2;
+; CHECK-PTX-NEXT:    ld.global.nc.u8 %rs3, [%rd3+6];
 ; CHECK-PTX-NEXT:    cvt.u16.u8 %rs4, %rs3;
-; CHECK-PTX-NEXT:    ld.global.nc.u8 %rs5, [%rd1+6];
+; CHECK-PTX-NEXT:    st.local.u8 [%rd2+1], %rs4;
+; CHECK-PTX-NEXT:    ld.global.nc.u8 %rs5, [%rd3+5];
 ; CHECK-PTX-NEXT:    cvt.u16.u8 %rs6, %rs5;
 ; CHECK-PTX-NEXT:    st.local.u8 [%rd2], %rs6;
 ; CHECK-PTX-NEXT:    mov.b32 %r1, 1;
 ; CHECK-PTX-NEXT:    st.u32 [%SP+8], %r1;
-; CHECK-PTX-NEXT:    mov.b16 %rs9, 1;
-; CHECK-PTX-NEXT:    st.u8 [%SP+12], %rs9;
-; CHECK-PTX-NEXT:    mov.b64 %rd2, 1;
-; CHECK-PTX-NEXT:    st.u64 [%SP+16], %rd2;
-; CHECK-PTX-NEXT:    add.u64 %rd3, %SP, 8;
+; CHECK-PTX-NEXT:    mov.b16 %rs7, 1;
+; CHECK-PTX-NEXT:    st.u8 [%SP+12], %rs7;
+; CHECK-PTX-NEXT:    mov.b64 %rd4, 1;
+; CHECK-PTX-NEXT:    st.u64 [%SP+16], %rd4;
+; CHECK-PTX-NEXT:    add.u64 %rd5, %SP, 8;
 ; CHECK-PTX-NEXT:    { // callseq 1, 0
 ; CHECK-PTX-NEXT:    .param .b32 param0;
 ; CHECK-PTX-NEXT:    st.param.b32 [param0], 1;
 ; CHECK-PTX-NEXT:    .param .b64 param1;
-; CHECK-PTX-NEXT:    st.param.b64 [param1], %rd3;
+; CHECK-PTX-NEXT:    st.param.b64 [param1], %rd5;
 ; CHECK-PTX-NEXT:    .param .b32 retval0;
 ; CHECK-PTX-NEXT:    call.uni (retval0),
 ; CHECK-PTX-NEXT:    variadics2,
@@ -378,25 +380,28 @@ define dso_local void @qux() {
 ; CHECK-PTX-NEXT:    .reg .b64 %SP;
 ; CHECK-PTX-NEXT:    .reg .b64 %SPL;
 ; CHECK-PTX-NEXT:    .reg .b32 %r<3>;
-; CHECK-PTX-NEXT:    .reg .b64 %rd<6>;
+; CHECK-PTX-NEXT:    .reg .b64 %rd<10>;
 ; CHECK-PTX-EMPTY:
 ; CHECK-PTX-NEXT:  // %bb.0: // %entry
 ; CHECK-PTX-NEXT:    mov.u64 %SPL, __local_depot7;
 ; CHECK-PTX-NEXT:    cvta.local.u64 %SP, %SPL;
-; CHECK-PTX-NEXT:    ld.global.nc.u64 %rd1, [__const_$_qux_$_s];
-; CHECK-PTX-NEXT:    st.u64 [%SP], %rd1;
-; CHECK-PTX-NEXT:    mov.u64 %rd2, __const_$_qux_$_s;
-; CHECK-PTX-NEXT:    ld.global.nc.u64 %rd3, [%rd2+8];
-; CHECK-PTX-NEXT:    st.u64 [%SP+8], %rd3;
-; CHECK-PTX-NEXT:    mov.b64 %rd4, 1;
-; CHECK-PTX-NEXT:    st.u64 [%SP+16], %rd4;
-; CHECK-PTX-NEXT:    add.u64 %rd5, %SP, 16;
+; CHECK-PTX-NEXT:    add.u64 %rd2, %SPL, 0;
+; CHECK-PTX-NEXT:    mov.u64 %rd3, __const_$_qux_$_s;
+; CHECK-PTX-NEXT:    ld.global.nc.u64 %rd4, [%rd3+8];
+; CHECK-PTX-NEXT:    st.local.u64 [%rd2+8], %rd4;
+; CHECK-PTX-NEXT:    ld.global.nc.u64 %rd5, [__const_$_qux_$_s];
+; CHECK-PTX-NEXT:    st.local.u64 [%rd2], %rd5;
+; CHECK-PTX-NEXT:    mov.b64 %rd6, 1;
+; CHECK-PTX-NEXT:    st.u64 [%SP+16], %rd6;
+; CHECK-PTX-NEXT:    ld.local.u64 %rd7, [%rd2];
+; CHECK-PTX-NEXT:    ld.local.u64 %rd8, [%rd2+8];
+; CHECK-PTX-NEXT:    add.u64 %rd9, %SP, 16;
 ; CHECK-PTX-NEXT:    { // callseq 3, 0
 ; CHECK-PTX-NEXT:    .param .align 8 .b8 param0[16];
-; CHECK-PTX-NEXT:    st.param.b64 [param0], %rd1;
-; CHECK-PTX-NEXT:    st.param.b64 [param0+8], %rd3;
+; CHECK-PTX-NEXT:    st.param.b64 [param0], %rd7;
+; CHECK-PTX-NEXT:    st.param.b64 [param0+8], %rd8;
 ; CHECK-PTX-NEXT:    .param .b64 param1;
-; CHECK-PTX-NEXT:    st.param.b64 [param1], %rd5;
+; CHECK-PTX-NEXT:    st.param.b64 [param1], %rd9;
 ; CHECK-PTX-NEXT:    .param .b32 retval0;
 ; CHECK-PTX-NEXT:    call.uni (retval0),
 ; CHECK-PTX-NEXT:    variadics4,