[NVPTX] Fix failing tests on incompatible PTX version

[jhuber6]

Summary:
It seems some recent changes caused us to emit labels in the debug
sections. This is only allowed in PTX70 which caused the ptxas lines
to fail if present. Simply add +ptx70 on it, maybe this needs to
investigate the change that caused this? But I'm not sure how much we
care about the older targets.

[llvmbot]

@llvm/pr-subscribers-debuginfo

@llvm/pr-subscribers-backend-nvptx

Author: Joseph Huber (jhuber6)

Changes

Summary:
It seems some recent changes caused us to emit labels in the debug
sections. This is only allowed in PTX70 which caused the ptxas lines
to fail if present. Simply add +ptx70 on it, maybe this needs to
investigate the change that caused this? But I'm not sure how much we
care about the older targets.

---

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

3 Files Affected:

• (modified) llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll (+1-1)
• (modified) llvm/test/DebugInfo/NVPTX/debug-info.ll (+1-1)
• (modified) llvm/test/DebugInfo/NVPTX/debug-loc-offset.ll (+1-1)

diff --git a/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll b/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
index d96d5629d03f03..43229975d28799 100644
--- a/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
+++ b/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
@@ -1,5 +1,5 @@
 ; RUN: llc -mtriple=nvptx64-nvidia-cuda < %s | FileCheck %s
-; RUN: %if ptxas %{ llc -mtriple=nvptx64-nvidia-cuda < %s | %ptxas-verify %}
+; RUN: %if ptxas %{ llc -mtriple=nvptx64-nvidia-cuda -mattr=+ptx70 < %s | %ptxas-verify %}
 
 ; Generated with -O1 from:
 ; int f1();
diff --git a/llvm/test/DebugInfo/NVPTX/debug-info.ll b/llvm/test/DebugInfo/NVPTX/debug-info.ll
index 35443738e05379..3fea59923cbe88 100644
--- a/llvm/test/DebugInfo/NVPTX/debug-info.ll
+++ b/llvm/test/DebugInfo/NVPTX/debug-info.ll
@@ -1,5 +1,5 @@
 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-nvidia-cuda | %ptxas-verify %}
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-nvidia-cuda -mattr=+ptx70 | %ptxas-verify %}
 
 ; // Bitcode in this test case is reduced version of compiled code below:
 ;__device__ inline void res(float x, float y, ptr res) { *res = x + y; }
diff --git a/llvm/test/DebugInfo/NVPTX/debug-loc-offset.ll b/llvm/test/DebugInfo/NVPTX/debug-loc-offset.ll
index 45c387fabddacd..8f3c5757adf0e9 100644
--- a/llvm/test/DebugInfo/NVPTX/debug-loc-offset.ll
+++ b/llvm/test/DebugInfo/NVPTX/debug-loc-offset.ll
@@ -1,5 +1,5 @@
 ; RUN: llc -mtriple=nvptx64-nvidia-cuda < %s | FileCheck %s
-; RUN: %if ptxas %{ llc -mtriple=nvptx64-nvidia-cuda < %s | %ptxas-verify %}
+; RUN: %if ptxas %{ llc -mtriple=nvptx64-nvidia-cuda < %s -mattr=+ptx70 | %ptxas-verify %}
 
 ; CHECK: .target sm_{{[0-9]+}}, debug
 

[Artem-B]

But I'm not sure how much we care about the older targets.
ptx7.0 was introduced in cuda-11. Not supporting older versions is OK for the new features. Just document the requirements somewhere.

Also, GPU-side debug has known issues in any case, so anyone relying on it will likely run into issues. E.g. an attempt to use __PRETTY_FUNCTION__ macro on the GPU side will also produce debug info that ptxas will refuse.

[Artem-B]

Do you have a reproducer for the issue? If that would be happening for CUDA those tests would be failing now, yet I don't see any failures.

[jhuber6]

Maybe it was fixed since I made this, but I was just running check-llvm with the PTX stuff set.

[Artem-B]

Ah. It's ptxas that complains. I can reproduce it. This is something that needs fixing. I was planning to look into other debug info issues next week, so I'll take a look at this, too.

Filed #112998