[NVPTX] Add support of ptxas v12.0 to the LIT tests

[pavelkopyl]

Summary

Kepler architecture support was removed from CUDA 12.0.
Skip tests that target to Kepler architecture when using ptxas v12.0.

[tra]

Description could use some details. I.e. what makes sm_35 special? I assume it's due to the fact that sm_35 is no longer supported.

Instead of special-casing ptxas-12.0, perhpas a better way is to just always run ptxas with -arch=sm_60 so it would work with ptxas from all reasonably new CUDA versions.

llvm/test/lit.cfg.py

215–216

There are also 11.7 and 11.8.

[pavelkopyl]

• Rebase

[pavelkopyl]

In D140436#4026879, @tra wrote:

Description could use some details. I.e. what makes sm_35 special? I assume it's due to the fact that sm_35 is no longer supported.

Instead of special-casing ptxas-12.0, perhpas a better way is to just always run ptxas with -arch=sm_60 so it would work with ptxas from all reasonably new CUDA versions.

Thank you. I've added details to the description.

As for choosing to an actual ptxas version, yes we could do this.
But original motivation to use minimal possible version of architecture/PTX was to catch errors related to appearance of new instructions and adding appropriate checks to the codegen. For example, https://reviews.llvm.org/D123040
On the other hand, all such cases seems to have been detected and fixed for old architectures, so we could focus on current ones.

llvm/test/lit.cfg.py

215–216

Done, thank you.

[pavelkopyl]

I've prepared a new review where ptxas is run -march=sm_60: https://reviews.llvm.org/D141736

[pavelkopyl]

Abandoned in favor on https://reviews.llvm.org/D141736

llvm/test/CodeGen/NVPTX/aggregate-return.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 declare <2 x float> @barv(<2 x float> %input)
 declare <3 x float> @barv3(<3 x float> %input)
 declare [2 x float] @bara([2 x float] %input)
 declare {float, float} @bars({float, float} %input)
 
 define void @test_v2f32(<2 x float> %input, ptr %output) {
 ; CHECK-LABEL: @test_v2f32

llvm/test/CodeGen/NVPTX/branch-fold.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -disable-cgp -verify-machineinstrs | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 -disable-cgp -verify-machineinstrs | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 -disable-cgp -verify-machineinstrs | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; Disable CGP which also folds branches, so that only BranchFolding is under
 ; the spotlight.
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
 define void @foo(i32 %x, ptr %output) {

llvm/test/CodeGen/NVPTX/bug21465.ll

 ; RUN: opt < %s -nvptx-lower-args -S | FileCheck %s
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix PTX
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 %struct.S = type { i32, i32 }
 
 ; Function Attrs: nounwind
 define void @_Z11TakesStruct1SPi(ptr byval(%struct.S) nocapture readonly %input, ptr nocapture %output) #0 {

llvm/test/CodeGen/NVPTX/bug26185-2.ll

 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; Verify that we correctly emit code for extending ldg/ldu. We do not expose
 ; extending variants in the backend, but the ldg/ldu selection code may pick
 ; extending loads as candidates. We do want to support this, so make sure we
 ; emit the necessary cvt.* instructions to implement the extension and let ptxas
 ; emit the real extending loads.
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"

llvm/test/CodeGen/NVPTX/bug26185.ll

 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit
 ; registers in the backend, so these loads need special handling.
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 ; CHECK-LABEL: ex_zext

llvm/test/CodeGen/NVPTX/bypass-div.ll

 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; 64-bit divides and rems should be split into a fast and slow path where
 ; the fast path uses a 32-bit operation.
 
 define void @sdiv64(i64 %a, i64 %b, ptr %retptr) {
 ; CHECK-LABEL: sdiv64(
 ; CHECK:        div.s64
 ; CHECK:        div.u32

llvm/test/CodeGen/NVPTX/divrem-combine.ll

 ; RUN: llc -O2 < %s -march=nvptx -mcpu=sm_35 | FileCheck %s --check-prefix=O2 --check-prefix=CHECK
 ; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_35 | FileCheck %s --check-prefix=O0 --check-prefix=CHECK
-; RUN: %if ptxas %{ llc -O2 < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
-; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc -O2 < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
+; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; The following IR
 ;
 ;   quot = n / d
 ;   rem  = n % d
 ;
 ; should be transformed into
 ;

llvm/test/CodeGen/NVPTX/extloadv.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 define void @foo(ptr nocapture readonly %x_value, ptr nocapture %output) #0 {
   %1 = load <4 x float>, ptr %x_value, align 16
   %2 = fpext <4 x float> %1 to <4 x double>
 ; CHECK-NOT: ld.v2.f32 {%fd{{[0-9]+}}, %fd{{[0-9]+}}}, [%rd{{[0-9]+}}];
 ; CHECK:  cvt.f64.f32
 ; CHECK:  cvt.f64.f32
 ; CHECK:  cvt.f64.f32
 ; CHECK:  cvt.f64.f32
   store <4 x double> %2, ptr %output
   ret void
 }

llvm/test/CodeGen/NVPTX/ldg-invariant.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; Check that invariant loads from the global addrspace are lowered to
 ; ld.global.nc.
 
 ; CHECK-LABEL: @ld_global
 define i32 @ld_global(ptr addrspace(1) %ptr) {
 ; CHECK: ld.global.nc.{{[a-z]}}32
   %a = load i32, ptr addrspace(1) %ptr, !invariant.load !0

llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck -check-prefix=SM20 %s
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck -check-prefix=SM35 %s
 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 ; SM20-LABEL: .visible .entry foo1(
 ; SM20: ld.global.f32
 ; SM35-LABEL: .visible .entry foo1(
 ; SM35: ld.global.nc.f32

llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -O0 | FileCheck %s --check-prefix PTX
 ; RUN: opt < %s -S -nvptx-lower-aggr-copies | FileCheck %s --check-prefix IR
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 -O0 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 -O0 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; Verify that the NVPTXLowerAggrCopies pass works as expected - calls to
 ; llvm.mem* intrinsics get lowered to loops.
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "nvptx64-unknown-unknown"
 
 declare void @llvm.memcpy.p0.p0.i64(ptr nocapture, ptr nocapture readonly, i64, i1) #1

llvm/test/CodeGen/NVPTX/lower-alloca.ll

 ; RUN: opt < %s -S -nvptx-lower-alloca -infer-address-spaces | FileCheck %s
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix PTX
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 define void @kernel() {
 ; LABEL: @lower_alloca
 ; PTX-LABEL: .visible .entry kernel(
   %A = alloca i32

llvm/test/CodeGen/NVPTX/param-load-store.ll

 ; Verifies correctness of load/store of parameters and return values.
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -O0 -verify-machineinstrs | FileCheck -allow-deprecated-dag-overlap %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 -O0 -verify-machineinstrs | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 -O0 -verify-machineinstrs | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 %s_i1 = type { i1 }
 %s_i8 = type { i8 }
 %s_i16 = type { i16 }
 %s_f16 = type { half }
 %s_i32 = type { i32 }
 %s_f32 = type { float }
 %s_i64 = type { i64 }

llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll

 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 ; Check load from constant global variables.  These loads should be
 ; ld.global.nc (aka ldg).
 
 @gv_float = external constant float
 @gv_float2 = external constant <2 x float>
 @gv_float4 = external constant <4 x float>
 

llvm/test/CodeGen/NVPTX/reg-copy.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
 
 define void @PR24303(ptr %f) {
 ; CHECK-LABEL: .visible .entry PR24303(
 ; Do not use mov.f or mov.u to convert between float and int.
 ; CHECK-NOT: mov.{{f|u}}{{32|64}} %f{{[0-9]+}}, %r{{[0-9]+}}

llvm/test/CodeGen/NVPTX/rotate.ll

 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck --check-prefix=SM20 %s
 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck --check-prefix=SM35 %s
 ; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify -arch=sm_35 %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %if !ptxas-12.0 %{-arch=sm_35%} %}
 
 
 declare i32 @llvm.nvvm.rotate.b32(i32, i32)
 declare i64 @llvm.nvvm.rotate.b64(i64, i32)
 declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)
 
 ; SM20: rotate32
 ; SM35: rotate32

llvm/test/lit.cfg.py

def enable_ptxas(ptxas_executable):
     version = ptxas_version(ptxas_executable)
     if version:
         # ptxas is supposed to be backward compatible with previous
         # versions, so add a feature for every known version prior to
         # the current one.
         ptxas_known_versions = [
             (9, 0), (9, 1), (9, 2),
             (10, 0), (10, 1), (10, 2),
             (11, 0), (11, 1), (11, 2), (11, 3), (11, 4), (11, 5), (11, 6),
+            (11, 7), (11, 8), (12, 0),
         ]
 
         def version_int(ver):
             return ver[0] * 100 + ver[1]
 
         # ignore ptxas if its version is below the minimum supported
         # version
         min_version = ptxas_known_versions[0]