[llvm] 312055d - [NVPTX] Fix `ptxas` failures (NFC) (#125147)

[via llvm-commits]

Author: Justin Fargnoli
Date: 2025-02-01T16:44:37-08:00
New Revision: 312055d1da169f043e65b35fcd62d6d685700114

URL: https://github.com/llvm/llvm-project/commit/312055d1da169f043e65b35fcd62d6d685700114
DIFF: https://github.com/llvm/llvm-project/commit/312055d1da169f043e65b35fcd62d6d685700114.diff

LOG: [NVPTX] Fix `ptxas` failures (NFC) (#125147)

Note:
[lower-args.ll](https://github.com/llvm/llvm-project/compare/main...justinfargnoli:dev/jf/ptxas?expand=1#diff-649d37d1f897d829fb809025437ba5df2e0c8da8395bbac7be713cd8f5bd8237)
and
[kernel-param-align.ll](https://github.com/llvm/llvm-project/compare/main...justinfargnoli:dev/jf/ptxas?expand=1#diff-31f196478b41b95b51298eb8e2efccc8a6f1156f13b648c07db27dd09579f74e)
fail because`ptxas` doesn't support constant pointers in separate
complication mode (`-c`).

Added: 
    

Modified: 
    llvm/test/CodeGen/NVPTX/griddepcontrol.ll
    llvm/test/CodeGen/NVPTX/kernel-param-align.ll
    llvm/test/CodeGen/NVPTX/lower-args.ll

Removed: 
    


################################################################################
diff  --git a/llvm/test/CodeGen/NVPTX/griddepcontrol.ll b/llvm/test/CodeGen/NVPTX/griddepcontrol.ll
index fe15b3fe4afbd9..0bf9196aa2902f 100644
--- a/llvm/test/CodeGen/NVPTX/griddepcontrol.ll
+++ b/llvm/test/CodeGen/NVPTX/griddepcontrol.ll
@@ -1,6 +1,6 @@
 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
 ; RUN: llc < %s -mcpu=sm_90 -march=nvptx64 | FileCheck %s
-; RUN: %if ptxas-11.8 %{ llc < %s -mcpu=sm_90 -march=nvptx64 | %ptxas-verify %}
+; RUN: %if ptxas-11.8 %{ llc < %s -mcpu=sm_90 -march=nvptx64 | %ptxas-verify -arch=sm_90 %}
 
 define void @griddepcontrol() {
 ; CHECK-LABEL: griddepcontrol(

diff  --git a/llvm/test/CodeGen/NVPTX/kernel-param-align.ll b/llvm/test/CodeGen/NVPTX/kernel-param-align.ll
index 2889d2d89a8579..a56b85de801430 100644
--- a/llvm/test/CodeGen/NVPTX/kernel-param-align.ll
+++ b/llvm/test/CodeGen/NVPTX/kernel-param-align.ll
@@ -1,5 +1,5 @@
 ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_60 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 | %ptxas-verify %}
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 | %ptxas -arch=sm_60 - %}
 
 %struct.Large = type { [16 x double] }
 

diff  --git a/llvm/test/CodeGen/NVPTX/lower-args.ll b/llvm/test/CodeGen/NVPTX/lower-args.ll
index 269bba75dc5fb3..81b86c86d40de3 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args.ll
@@ -87,12 +87,12 @@ define ptx_kernel void @ptr_generic(ptr %out, ptr %in) {
 }
 
 ; COMMON-LABEL: ptr_nongeneric
-define ptx_kernel void @ptr_nongeneric(ptr addrspace(1) %out, ptr addrspace(4) %in) {
+define ptx_kernel void @ptr_nongeneric(ptr addrspace(1) %out, ptr addrspace(3) %in) {
 ; IR-NOT: addrspacecast
 ; PTX-NOT: cvta.to.global
-; PTX:  ld.const.u32
+; PTX:  ld.shared.u32
 ; PTX   st.global.u32
-  %v = load i32, ptr addrspace(4) %in, align 4
+  %v = load i32, ptr addrspace(3) %in, align 4
   store i32 %v, ptr addrspace(1) %out, align 4
   ret void
 }