[NVPTX] Fix ptxas requirements

[boomanaiden154]

#159388 caused some test failures for us internally:

exit status 255
ptxas /tmp/tmpxft_000026cd_00000000-0_stdin, line 37; error : Feature 'used_bytes_mask' requires PTX ISA .version 8.3 or later
ptxas /tmp/tmpxft_000026cd_00000000-0_stdin, line 37; error : Feature 'used_bytes_mask' requires .target sm_50 or higher
ptxas fatal : Ptx assembly aborted due to errors

Add some requirements/versions to get things working.

[llvmbot]

@llvm/pr-subscribers-backend-nvptx

Author: Aiden Grossman (boomanaiden154)

Changes

#159388 caused some test failures for us internally:

exit status 255
ptxas /tmp/tmpxft_000026cd_00000000-0_stdin, line 37; error : Feature 'used_bytes_mask' requires PTX ISA .version 8.3 or later
ptxas /tmp/tmpxft_000026cd_00000000-0_stdin, line 37; error : Feature 'used_bytes_mask' requires .target sm_50 or higher
ptxas fatal : Ptx assembly aborted due to errors

Add some requirements/versions to get things working.

---

Full diff: https://github.com/llvm/llvm-project/pull/171219.diff

2 Files Affected:

• (modified) llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll (+2-2)
• (modified) llvm/test/CodeGen/NVPTX/param-vectorize-device.ll (+1-1)

diff --git a/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll b/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
index a75ddd032d4c0..dd185f5c9f731 100644
--- a/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
+++ b/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
@@ -1,8 +1,8 @@
 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
 ; RUN: llc < %s | FileCheck -check-prefix=ENABLED %s
 ; RUN: llc -disable-nvptx-load-store-vectorizer < %s | FileCheck -check-prefix=DISABLED %s
-; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}
-; RUN: %if ptxas %{ llc -disable-nvptx-load-store-vectorizer < %s | %ptxas-verify %}
+; RUN: %if ptxas-sm_50 && ptxas-isa-8.3 %{ llc -mcpu=sm_50 < %s | %ptxas-verify -arch=sm_50 %}
+; RUN: %if ptxas-sm_50 && ptxas-isa-8.3 %{ llc -mcpu=sm_50 -disable-nvptx-load-store-vectorizer < %s | %ptxas-verify -arch=sm_50 %}
 
 target triple = "nvptx64-nvidia-cuda"
 
diff --git a/llvm/test/CodeGen/NVPTX/param-vectorize-device.ll b/llvm/test/CodeGen/NVPTX/param-vectorize-device.ll
index 643de006f14c4..f0ca80747ccf9 100644
--- a/llvm/test/CodeGen/NVPTX/param-vectorize-device.ll
+++ b/llvm/test/CodeGen/NVPTX/param-vectorize-device.ll
@@ -1,5 +1,5 @@
 ; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-unknown-unknown | %ptxas-verify %}
+; RUN: %if ptxas-sm_50 && ptxas-isa-8.3 %{ llc < %s -mtriple=nvptx64-unknown-unknown -mcpu=sm_50 | %ptxas-verify -arch=sm_50 %}
 ;
 ; Check that parameters of a __device__ function with private or internal
 ; linkage called from a __global__ (kernel) function get increased alignment,

[dakersnar]

Hold on a second actually

[dakersnar]

Sorry for making this mistake again. You beat me to the punch but I put up a fix here too: #171220.

I think you solution might be better (it uses the minimum ptx and arch versions where this feature is allowed, unlike mine which uses a later one) but I think you will see a small change in the check output in LoadStoreVectorizer.ll.

[dakersnar]

This line specifically: https://github.com/llvm/llvm-project/pull/171220/files#diff-5edb4f2e850369cccfb090e314d00d6aa4f9887da10dcac0475247c65eb7318dR55

[boomanaiden154]

This line specifically: https://github.com/llvm/llvm-project/pull/171220/files#diff-5edb4f2e850369cccfb090e314d00d6aa4f9887da10dcac0475247c65eb7318dR55

That shouldn't be necessary here since I didn't change the triple or cpu for the llc invocations that get piped into FileCheck. Should I be updating those too?

[Artem-B]

I've LGTM'ed #171220 as a more comprehensive fix.

[dakersnar]

That shouldn't be necessary here since I didn't change the triple or cpu for the llc invocations that get piped into FileCheck. Should I be updating those too?

Actually, yes, I'm wrong. You're correct, your version should pass. But regardless if we think that the other change is a more comprehensive fix, we should merge that one.

[boomanaiden154]

Sounds good to me to land #171220. It seems like our internal build system isn't actually executing the ptxas-verify lines, so since your fix is validated, let's go with that.

[Artem-B]

Let's continue on #171220. It may need a bit more work.