[NVPTX] Set default version of architecture to SM_30, PTX to 6.0.

[pavelkopyl]

Summary

Support of variadic functions triggers an assertion on several tests from
llvm/test/CodeGen/Generic/ if nvptx64-* is specified as a default triplet:

Support for variadic functions (unsized array parameter) introduced in
PTX ISA version 6.0 and requires target sm_30.

That happens because those tests contain variadic function calls and default
versions of both PTX ISA (3.2) and architecture (sm_20) are below the minimally required.

There were no observable problems with these tests before adding support of variadic
functions, because nvptx backend just didn't handle them properly generating invalid PTX code.

[pavelkopyl]

Example of failing tests: https://lab.llvm.org/buildbot/#/builders/232/builds/2407

[pavelkopyl]

• Update to correct patch version.

[tra]

LGTM with few nits.

llvm/test/CodeGen/NVPTX/sm-version.ll

35–37

Please add the checks for SM21 and SM30, too.

llvm/test/CodeGen/NVPTX/surf-tex.py

2–3

We may as well change it to sm_30, too. sm_20 is gone for all practical purposes. Even sm_30 is on the way out.

[pavelkopyl]

llvm/test/CodeGen/NVPTX/sm-version.ll

35–37

Sure, thank you.

llvm/test/CodeGen/NVPTX/surf-tex.py

2–3

Changing version to sm_30 also changes the generated code. The reason is that sm_20 (Fermi) has no support of image handles, as a result nvptx backed runs special pass NVPTXReplaceImageHandles to workaround this.
Textually such a code diff looks following:

.global .surfref gsurf;

sm_20:

suld.b.1d.b8.trap {%rs1}, [gsurf, {%r1}];

sm_30:

mov.u64 %rd3, gsurf;
suld.b.1d.b8.trap {%rs1}, [%rd3, {%r1}];

We can change this test to support sm_30. But I think it's better to do this in another review. Is that OK?

[pavelkopyl]

• Rebase
• Fixes according to a review notes

[tra]

llvm/test/CodeGen/NVPTX/surf-tex.py

2–3

OK. We can keep sm_20 test around for now.

[asavonic]

llvm/test/CodeGen/NVPTX/surf-tex.py

2–3

The reason is that sm_20 (Fermi) has no support of image handles, as a result nvptx backed runs special pass NVPTXReplaceImageHandles to workaround this.

Can we run NVPTXReplaceImageHandles for other targets as well? These extra mov instructions are unnecessary, and the pass can eliminate them.

[tra]

llvm/test/CodeGen/NVPTX/surf-tex.py

2–3

The moves do not matter all that much. ptxas is an optimizing assembler, which can deal with them. If we really want to get rid of them, a better way would be not to generate them which should be done during lowering, not by running an additional pass.

llvm/lib/Target/NVPTX/NVPTX.td

 
 //===----------------------------------------------------------------------===//
 // NVPTX supported processors.
 //===----------------------------------------------------------------------===//
 
 class Proc<string Name, list<SubtargetFeature> Features>
  : Processor<Name, NoItineraries, Features>;
 
-def : Proc<"sm_20", [SM20]>;
-def : Proc<"sm_21", [SM21]>;
+def : Proc<"sm_20", [SM20, PTX32]>;
+def : Proc<"sm_21", [SM21, PTX32]>;
 def : Proc<"sm_30", [SM30]>;
 def : Proc<"sm_32", [SM32, PTX40]>;
-def : Proc<"sm_35", [SM35]>;
+def : Proc<"sm_35", [SM35, PTX32]>;
 def : Proc<"sm_37", [SM37, PTX41]>;
 def : Proc<"sm_50", [SM50, PTX40]>;
 def : Proc<"sm_52", [SM52, PTX41]>;
 def : Proc<"sm_53", [SM53, PTX42]>;
 def : Proc<"sm_60", [SM60, PTX50]>;
 def : Proc<"sm_61", [SM61, PTX50]>;
 def : Proc<"sm_62", [SM62, PTX50]>;
 def : Proc<"sm_70", [SM70, PTX60]>;

llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp

NoF16Math("nvptx-no-f16-math", cl::Hidden,
               cl::init(false));
 
 // Pin the vtable to this file.
 void NVPTXSubtarget::anchor() {}
 
 NVPTXSubtarget &NVPTXSubtarget::initializeSubtargetDependencies(StringRef CPU,
                                                                 StringRef FS) {
     // Provide the default CPU if we don't have one.
-    TargetName = std::string(CPU.empty() ? "sm_20" : CPU);
+    TargetName = std::string(CPU.empty() ? "sm_30" : CPU);
 
     ParseSubtargetFeatures(TargetName, /*TuneCPU*/ TargetName, FS);
 
-    // Set default to PTX 3.2 (CUDA 5.5)
+    // Set default to PTX 6.0 (CUDA 9.0)
     if (PTXVersion == 0) {
-      PTXVersion = 32;
+      PTXVersion = 60;
   }
 
   return *this;
 }
 
 NVPTXSubtarget::NVPTXSubtarget(const Triple &TT, const std::string &CPU,
                                const std::string &FS,
                                const NVPTXTargetMachine &TM)

llvm/test/CodeGen/NVPTX/nvvm-reflect-arch.ll

 ; Libdevice in recent CUDA versions relies on __CUDA_ARCH reflecting GPU type.
 ; Verify that __nvvm_reflect() is replaced with an appropriate value.
 ;
-; RUN: opt %s -S -passes='default<O2>' -mtriple=nvptx64 \
+; RUN: opt %s -S -passes='default<O2>' -mtriple=nvptx64 -mcpu=sm_20 \
 ; RUN:   | FileCheck %s --check-prefixes=COMMON,SM20
 ; RUN: opt %s -S -passes='default<O2>' -mtriple=nvptx64 -mcpu=sm_35 \
 ; RUN:   | FileCheck %s --check-prefixes=COMMON,SM35
 
 @"$str" = private addrspace(1) constant [12 x i8] c"__CUDA_ARCH\00"
 
 declare i32 @__nvvm_reflect(ptr)
 

llvm/test/CodeGen/NVPTX/nvvm-reflect-ocl.ll

 ; Verify that __nvvm_reflect_ocl() is replaced with an appropriate value
 ;
-; RUN: opt %s -S -passes='default<O2>' -mtriple=nvptx64 \
+; RUN: opt %s -S -passes='default<O2>' -mtriple=nvptx64 -mcpu=sm_20 \
 ; RUN:   | FileCheck %s --check-prefixes=COMMON,SM20
 ; RUN: opt %s -S -passes='default<O2>' -mtriple=nvptx64 -mcpu=sm_35 \
 ; RUN:   | FileCheck %s --check-prefixes=COMMON,SM35
 
 @"$str" = private addrspace(4) constant [12 x i8] c"__CUDA_ARCH\00"
 
 declare i32 @__nvvm_reflect_ocl(ptr addrspace(4) noundef)
 

llvm/test/CodeGen/NVPTX/sm-version.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 | FileCheck %s --check-prefix=SM60
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_61 | FileCheck %s --check-prefix=SM61
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_62 | FileCheck %s --check-prefix=SM62
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 | FileCheck %s --check-prefix=SM70
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_75 | FileCheck %s --check-prefix=SM75
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 | FileCheck %s --check-prefix=SM80
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_86 | FileCheck %s --check-prefix=SM86
 
-; SM30: .version 3.2
+; SM20: .version 3.2
+; SM21: .version 3.2
+; SM30: .version 6.0
 ; SM32: .version 4.0
 ; SM35: .version 3.2
 ; SM37: .version 4.1
 ; SM50: .version 4.0
 ; SM52: .version 4.1
 ; SM53: .version 4.2
 ; SM60: .version 5.0
 ; SM61: .version 5.0

llvm/test/CodeGen/NVPTX/surf-tex.py

 # RUN: %python %s --target=cuda --tests=suld,sust,tex,tld4 --gen-list=%t.list > %t-cuda.ll
-# RUN: llc %t-cuda.ll -verify-machineinstrs -o - | FileCheck %t-cuda.ll
-# RUN: %if ptxas %{ llc %t-cuda.ll -verify-machineinstrs -o - | %ptxas-verify %}
+# RUN: llc -mcpu=sm_20 %t-cuda.ll -verify-machineinstrs -o - | FileCheck %t-cuda.ll
+# RUN: %if ptxas %{ llc -mcpu=sm_20 %t-cuda.ll -verify-machineinstrs -o - | %ptxas-verify %}
 
 # We only need to run this second time for texture tests, because
 # there is a difference between unified and non-unified intrinsics.
 #
 # RUN: %python %s --target=nvcl --tests=suld,sust,tex,tld4 --gen-list-append --gen-list=%t.list > %t-nvcl.ll
 # RUN: llc %t-nvcl.ll -verify-machineinstrs -o - | FileCheck %t-nvcl.ll
 # RUN: %if ptxas %{ llc %t-nvcl.ll -verify-machineinstrs -o - | %ptxas-verify %}
 

llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll

 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
 ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-nvidia-cuda | %ptxas-verify %}
 
-; CHECK: .target sm_20, debug
+; CHECK: .target sm_30, debug
 
 ; CHECK: .visible .func use_dbg_declare()
 ; CHECK: .local .align 8 .b8 __local_depot0[8];
 ; CHECK: mov.u64 %SPL, __local_depot0;
 ; CHECK: add.u64 %rd1, %SP, 0;
 ; CHECK: .loc 1 5 3                   // t.c:5:3
 ; CHECK: { // callseq 0, 0
 ; CHECK: .reg .b32 temp_param_reg;