[NVPTX] Fix NVPTXReplaceImageHandles for multiple uses of a texref

[slydiman]

Summary

The texsurf_handle is removed by NVPTXReplaceImageHandles.cpp. There are more than one uses of the texsurf_handle, one of them is a regular function call, and one of them is a texture intrinsic. The current hacky logic in NVPTXReplaceImageHandles.cpp for CUDA cannot handle such a mixed use. This patch fixes this issue.

[tra]

LGTM.

llvm/lib/Target/NVPTX/NVPTXReplaceImageHandles.cpp

9

NVIDIA has stopped supporting Fermi(AKA sm_20) few years ago: https://nvidia.custhelp.com/app/answers/detail/a_id/4654/~/support-plan-for-fermi-series-geforce-gpus

We should probably start thinking of removing sm_20 support from LLVM, too.

llvm/lib/Target/NVPTX/NVPTXReplaceImageHandles.cpp

 //===-- NVPTXReplaceImageHandles.cpp - Replace image handles for Fermi ----===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // On Fermi, image handles are not supported. To work around this, we traverse
 // the machine code and replace image handles with concrete symbols. For this
 // to work reliably, inlining of all function call must be performed.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTX.h"
 #include "NVPTXMachineFunctionInfo.h"
 #include "NVPTXSubtarget.h"
bool NVPTXReplaceImageHandles::runOnMachineFunction(MachineFunction &MF) {
   for (MachineBasicBlock &MBB : MF)
     for (MachineInstr &MI : MBB)
       Changed |= processInstr(MI);
 
   // Now clean up any handle-access instructions
   // This is needed in debug mode when code cleanup passes are not executed,
   // but we need the handle access to be eliminated because they are not
   // valid instructions when image handles are disabled.
-  for (MachineInstr *MI : InstrsToRemove)
+  for (MachineInstr *MI : InstrsToRemove) {
+    unsigned DefReg = MI->getOperand(0).getReg();
+    // Only these that are not used can be removed.
+    if (MF.getRegInfo().use_nodbg_empty(DefReg))
       MI->eraseFromParent();
+  }
   return Changed;
 }
 
 static unsigned suldRegisterToIndexOpcode(unsigned RegOC) {
   switch (RegOC) {
   case NVPTX::SULD_1D_I8_CLAMP_R:
     return NVPTX::SULD_1D_I8_CLAMP_I;
   case NVPTX::SULD_1D_I16_CLAMP_R:

llvm/test/CodeGen/NVPTX/tex-read-cuda.ll

; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}]
   %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
   %ret = extractvalue { float, float, float, float } %val, 0
 ; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]]
 ; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]]
   store float %ret, float* %red
   ret void
 }
 
-!nvvm.annotations = !{!1, !2, !3}
+declare float @texfunc(i64)
+
+; SM20-LABEL: .entry baz
+; SM30-LABEL: .entry baz
+define void @baz(float* %red, i32 %idx) {
+; SM30: mov.u64 %rd[[TEXHANDLE:[0-9]+]], tex0
+  %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @tex0)
+; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [tex0, {%r{{[0-9]+}}}]
+; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}]
+  %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
+  %ret = extractvalue { float, float, float, float } %val, 0
+; SM20: call.uni ([[RETVAL:.*]]),
+; SM30: call.uni ([[RETVAL:.*]]),
+; SM20: texfunc,
+; SM30: texfunc,
+  %texcall = tail call float @texfunc(i64 %texHandle)
+; SM20: ld.param.f32 %f[[TEXCALL:[0-9]+]], [[[RETVAL]]+0]
+; SM30: ld.param.f32 %f[[TEXCALL:[0-9]+]], [[[RETVAL]]+0]
+; SM20: add.rn.f32 %f[[RET2:[0-9]+]], %f[[RED]], %f[[TEXCALL]]
+; SM30: add.rn.f32 %f[[RET2:[0-9]+]], %f[[RED]], %f[[TEXCALL]]
+  %ret2 = fadd float %ret, %texcall
+; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RET2]]
+; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RET2]]
+  store float %ret2, float* %red
+  ret void
+}
+
+!nvvm.annotations = !{!1, !2, !3, !4}
 !1 = !{void (i64, float*, i32)* @foo, !"kernel", i32 1}
 !2 = !{void (float*, i32)* @bar, !"kernel", i32 1}
 !3 = !{i64 addrspace(1)* @tex0, !"texture", i32 1}
+!4 = !{void (float*, i32)* @baz, !"kernel", i32 1}