[clang] [clang][Tooling] Fix assertion failure when processing CUDA files (PR #173762)

[via cfe-commits]

https://github.com/zeyi2 updated https://github.com/llvm/llvm-project/pull/173762

>From f3b4e516546fc0d43600c8a7409dae9d9857d798 Mon Sep 17 00:00:00 2001
From: mtx <mitchell.xu2 at gmail.com>
Date: Sun, 28 Dec 2025 15:51:55 +0800
Subject: [PATCH 1/3] [clang-tidy] Fix assertion failure when processing CUDA
 files

---
 clang/lib/Tooling/Tooling.cpp | 17 ++++-------------
 1 file changed, 4 insertions(+), 13 deletions(-)

diff --git a/clang/lib/Tooling/Tooling.cpp b/clang/lib/Tooling/Tooling.cpp
index f10aa524674da..f19f2d36bb60e 100644
--- a/clang/lib/Tooling/Tooling.cpp
+++ b/clang/lib/Tooling/Tooling.cpp
@@ -102,19 +102,10 @@ static bool ignoreExtraCC1Commands(const driver::Compilation *Compilation) {
       if (isa<driver::BindArchAction>(A))
         A = *A->input_begin();
       if (isa<driver::OffloadAction>(A)) {
-        // Offload compilation has 2 top-level actions, one (at the front) is
-        // the original host compilation and the other is offload action
-        // composed of at least one device compilation. For such case, general
-        // tooling will consider host-compilation only. For tooling on device
-        // compilation, device compilation only option, such as
-        // `--cuda-device-only`, needs specifying.
-        assert(Actions.size() > 1);
-        assert(
-            isa<driver::CompileJobAction>(Actions.front()) ||
-            // On MacOSX real actions may end up being wrapped in
-            // BindArchAction.
-            (isa<driver::BindArchAction>(Actions.front()) &&
-             isa<driver::CompileJobAction>(*Actions.front()->input_begin())));
+        // For offload compilation, general tooling will consider host
+        // compilation only. For tooling on device compilation, device
+        // compilation only option, such as `--cuda-device-only`, needs
+        // specifying.
         OffloadCompilation = true;
         break;
       }

>From c3b8a75a7fc1d031e974b866a1b12a6dba617557 Mon Sep 17 00:00:00 2001
From: mtx <mitchell.xu2 at gmail.com>
Date: Sun, 28 Dec 2025 23:08:14 +0800
Subject: [PATCH 2/3] Dirty fix

---
 clang/lib/Tooling/Tooling.cpp | 21 +++++++++++++++++----
 1 file changed, 17 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Tooling/Tooling.cpp b/clang/lib/Tooling/Tooling.cpp
index f19f2d36bb60e..4edbfa699151f 100644
--- a/clang/lib/Tooling/Tooling.cpp
+++ b/clang/lib/Tooling/Tooling.cpp
@@ -102,10 +102,23 @@ static bool ignoreExtraCC1Commands(const driver::Compilation *Compilation) {
       if (isa<driver::BindArchAction>(A))
         A = *A->input_begin();
       if (isa<driver::OffloadAction>(A)) {
-        // For offload compilation, general tooling will consider host
-        // compilation only. For tooling on device compilation, device
-        // compilation only option, such as `--cuda-device-only`, needs
-        // specifying.
+        // Offload compilation has 2 top-level actions, one (at the front) is
+        // the original host compilation and the other is offload action
+        // composed of at least one device compilation. For such case, general
+        // tooling will consider host-compilation only. For tooling on device
+        // compilation, device compilation only option, such as
+        // `--cuda-device-only`, needs specifying.
+        if (Actions.size() > 1) {
+          assert(
+              isa<driver::CompileJobAction>(Actions.front()) ||
+              // On MacOSX real actions may end up being wrapped in
+              // BindArchAction.
+              (isa<driver::BindArchAction>(Actions.front()) &&
+               isa<driver::CompileJobAction>(*Actions.front()->input_begin())));
+        }
+        // FIXME: CUDA/HIP can produce a single top-level OffloadAction (e.g.
+        // -fsyntax-only, -E, -M), which contradicts the expectation of at least
+        // two top-level actions.
         OffloadCompilation = true;
         break;
       }

>From 5a35a9d3a15aea0e87e17d0696f1e80cbce62826 Mon Sep 17 00:00:00 2001
From: mtx <mitchell.xu2 at gmail.com>
Date: Mon, 12 Jan 2026 22:54:47 +0800
Subject: [PATCH 3/3] Add testcases and documents

---
 clang/docs/ReleaseNotes.rst             | 5 ++++-
 clang/unittests/Tooling/ToolingTest.cpp | 7 +++++++
 2 files changed, 11 insertions(+), 1 deletion(-)

diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 2319ff13f7864..8b83fa009dad7 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -730,7 +730,10 @@ CUDA/HIP Language Changes
 CUDA Support
 ^^^^^^^^^^^^
 
-Support calling `consteval` function between different target.
+- Fixed an assertion failure when processing CUDA files with `-fsyntax-only` or
+  device architecture flags.
+
+- Support calling `consteval` function between different target.
 
 AIX Support
 ^^^^^^^^^^^
diff --git a/clang/unittests/Tooling/ToolingTest.cpp b/clang/unittests/Tooling/ToolingTest.cpp
index 9a7559405c43c..b25493b29f5df 100644
--- a/clang/unittests/Tooling/ToolingTest.cpp
+++ b/clang/unittests/Tooling/ToolingTest.cpp
@@ -78,6 +78,13 @@ TEST(runToolOnCode, FindsNoTopLevelDeclOnEmptyCode) {
   EXPECT_FALSE(FoundTopLevelDecl);
 }
 
+TEST(runToolOnCode, CudaSyntaxOnly) {
+  EXPECT_TRUE(runToolOnCodeWithArgs(
+      std::make_unique<TestAction>(std::make_unique<clang::ASTConsumer>()),
+      "__global__ void k() {}",
+      {"-fsyntax-only", "-x", "cuda", "--offload-arch=sm_70"}));
+}
+
 namespace {
 class FindClassDeclXConsumer : public clang::ASTConsumer {
  public: