[CUDA] device side asserts crash clang when using -g

[OgnianM]

$ clang++ -g -x cuda -c test.cu
ptxas /tmp/test-sm_52-55e123.s, line 841; fatal   : Parsing error near '.': syntax error
ptxas fatal   : Ptx assembly aborted due to errors
clang++: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 20.0.0git (https://github.com/llvm/llvm-project.git 5bd38a98d7585841c1688f6b9eee8ce150dc429c)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/local/bin
clang++: note: diagnostic msg: 
********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang++: note: diagnostic msg: /tmp/test-731cd8.cu
clang++: note: diagnostic msg: /tmp/test-sm_52-df7b05.cu
clang++: note: diagnostic msg: /tmp/test-731cd8.sh
clang++: note: diagnostic msg: 

********************

test.cu:

#include <cassert>

__global__ void test(int x) { assert(x > 10); }


int main() {
	test<<<1,1>>>(5);
}

[jhuber6]

PTX doesn't accept . as a separator in names unlike the rest of the LLVM. We normally replace all of these . characters with $ when it's emitted. However, in this case the transformation is not applied to the debug symbol, which both causes the error and does not correctly reference the actual assert string.

.global .align 1 .b8 __PRETTY_FUNCTION___$__Z4testi[15] = {118, 111, 105, 100, 32, 116, 101, 115, 116, 40, 105, 110, 116, 41};
...
.b64 __PRETTY_FUNCTION__._Z4testi

See llvm/lib/Target/NVPTX/NVPTXAssignValidGlobalNames.cpp and try to find where we also emit the debug symbols.

[llvmbot]

Hi!

This issue may be a good introductory issue for people new to working on LLVM. If you would like to work on this issue, your first steps are:

1. Check that no other contributor has already been assigned to this issue. If you believe that no one is actually working on it despite an assignment, ping the person. After one week without a response, the assignee may be changed.
2. In the comments of this issue, request for it to be assigned to you, or just create a pull request after following the steps below. Mention this issue in the description of the pull request.
3. Fix the issue locally.
4. Run the test suite locally. Remember that the subdirectories under test/ create fine-grained testing targets, so you can e.g. use make check-clang-ast to only run Clang's AST tests.
5. Create a Git commit.
6. Run git clang-format HEAD~1 to format your changes.
7. Open a pull request to the upstream repository on GitHub. Detailed instructions can be found in GitHub's documentation. Mention this issue in the description of the pull request.

If you have any further questions about this issue, don't hesitate to ask via a comment in the thread below.

[llvmbot]

@llvm/issue-subscribers-good-first-issue

Author: Ognyan Mirev (OgnianM)

``` $ clang++ -g -x cuda -c test.cu ptxas /tmp/test-sm_52-55e123.s, line 841; fatal : Parsing error near '.': syntax error ptxas fatal : Ptx assembly aborted due to errors clang++: error: ptxas command failed with exit code 255 (use -v to see invocation) clang version 20.0.0git (https://github.com/llvm/llvm-project.git 5bd38a9) Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /usr/local/bin clang++: note: diagnostic msg: ********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang++: note: diagnostic msg: /tmp/test-731cd8.cu
clang++: note: diagnostic msg: /tmp/test-sm_52-df7b05.cu
clang++: note: diagnostic msg: /tmp/test-731cd8.sh
clang++: note: diagnostic msg:

---

test.cu:

#include <cassert>

global void test(int x) { assert(x > 10); }

int main() {
test<<<1,1>>>(5);
}

</details>

[vyom1611]

I can take a look at this, please assign it to me

[Artem-B]

This is the same issue as #77009 and #58491