From 9ab3dd9a8ec1bff75edef6ec7623527c512db397 Mon Sep 17 00:00:00 2001 From: Build Script Date: Fri, 17 Jul 2026 10:13:50 -0700 Subject: [PATCH 1/3] fix: capture #include content in hipRTC cache key via preprocessing MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit computeHiprtcCacheKey() keyed on the raw source, so it never saw the content of headers pulled in by #include — in particular headers resolved from -I filesystem paths (issue #1335). Editing such a header did not invalidate the cache and stale SPIR-V was served silently. When the source may include headers, preprocess it (clang -E -P) and key on the preprocessed translation unit, which embeds all #included content. -P suppresses line markers and -ffile-prefix-map normalizes the temp path so the output is deterministic across runs (verified two cold runs produce identical output). If preprocessing fails for a source that may include headers, caching is disabled for that program (no lookup, no new entry) rather than falling back to a raw-source key: the latter would silently ignore header edits and risk serving stale SPIR-V, i.e. reintroduce the very bug being fixed. To avoid adding a preprocess subprocess to every call, a conservative scan skips it for self-contained sources (no #include/#import): the raw source is then a complete key. The scan de-splices backslash-newline continuations and matches an unanchored '#[ \t]*(include|import)', so it does not miss indented or line-split directives; a match in a comment or string only costs a harmless extra preprocess, never a stale hit. Closes #1335. --- src/spirv_hiprtc.cc | 153 ++++++++++++++++++++++++++++++++++++++------ 1 file changed, 132 insertions(+), 21 deletions(-) diff --git a/src/spirv_hiprtc.cc b/src/spirv_hiprtc.cc index 770fe8f63..27258634e 100644 --- a/src/spirv_hiprtc.cc +++ b/src/spirv_hiprtc.cc @@ -159,7 +159,8 @@ static std::string escapeWithSingleQuotes(const std::string &Str) { static std::string createCompileCommand(const CompileOptions &Options, const fs::path &WorkingDirectory, const fs::path &SourceFile, - const fs::path &OutputFile) { + const fs::path &OutputFile, + bool PreprocessOnly = false) { std::string CompileCommand; @@ -209,7 +210,21 @@ static std::string createCompileCommand(const CompileOptions &Options, if (!Options.HasO) Append("-O2"); - Append("-c"); + if (PreprocessOnly) { + Append("-E"); + // Preprocess-only with options for determinism: + // -P suppresses line markers, which would otherwise embed absolute + // temp paths and make the output non-deterministic. + Append("-P"); + // The source lives in a per-run (random) temp dir. This maps the + // temp dir to "." so those expand deterministically. assert() macro + // expansion to __FILE__ is probably the most likely case where this mapping + // is needed. (Supported by Clang >= 10 / GCC >= 8, within chipStar's + // toolchain baseline.) + Append("-ffile-prefix-map=" + WorkingDirectory.string() + "=."); + } else { + Append("-c"); + } Append(SourceFile.string()); Append("-o"); @@ -248,6 +263,55 @@ static bool executeCommand(const fs::path &WorkingDirectory, return ReturnCode == 0; } +// Runs the compiler in preprocess-only mode over the user program so the cache +// key can reflect the *content* of #included headers in the RTC source. +// Returns the preprocessed translation unit, or nullopt if preprocessing +// did not succeed. +static std::optional +preprocessForCacheKey(const chipstar::Program &Program, + const CompileOptions &Options, + const fs::path &WorkingDirectory) { + auto SourceFile = WorkingDirectory / "pp_input.hip"; + auto OutputFile = WorkingDirectory / "pp_output.i"; + auto LogFile = WorkingDirectory / "pp.log"; + + // Write the in-memory headers and the raw user source. Name-expression + // injection is intentionally omitted: templates are not instantiated by the + // preprocessor, and name expressions are already hashed separately. + if (!createHeaderFiles(Program, WorkingDirectory)) + return std::nullopt; + { + std::ofstream F(SourceFile); + F << Program.getSource() << "\n"; + if (!F.good()) + return std::nullopt; + } + + std::string Cmd = createCompileCommand(Options, WorkingDirectory, SourceFile, + OutputFile, /*PreprocessOnly=*/true); + if (!executeCommand(WorkingDirectory, Cmd, LogFile)) + return std::nullopt; + + return readFromFile(OutputFile); +} + +// Conservative test for whether the source may pull in external header content +// via a preprocessor directive. When it can't, the raw source fully determines +// the compilation output (in-memory headers are hashed separately; toolchain +// headers are covered by the compiler-version component of the key), so we can +// skip the costly preprocess pass and key on the raw source. +// +// This MUST NOT produce false negatives: missing a real directive would let a +// header edit go unnoticed (stale cache). It errs the other way — a match +// inside a comment or string literal merely triggers an unnecessary preprocess, +// which is harmless. Backslash-newline continuations are stripped first so a +// directive split across lines (e.g. "#\\\ninclude") is still detected. +static bool sourceMayIncludeHeaders(const std::string &Source) { + std::string Spliced = std::regex_replace(Source, std::regex(R"(\\\n)"), ""); + static const std::regex Directive(R"(#[ \t]*(include|import))"); + return std::regex_search(Spliced, Directive); +} + static void getLoweredNameExpressions(chipstar::Program &Program, const fs::path &WorkingDirectory, const fs::path &LoweredNamesFile) { @@ -410,11 +474,16 @@ static uint64_t fnv1a64(const std::string &s) { /// source/options but a different set of registered name expressions would /// alias to the same cache entry, leaving some lowered-name lookups unmapped /// on cache hit. -static std::string computeHiprtcCacheKey(const chipstar::Program &Program, - int NumOptions, - const char *const *Options) { +static std::string +computeHiprtcCacheKey(const chipstar::Program &Program, int NumOptions, + const char *const *Options, + const std::optional &PreprocessedSource = + std::nullopt) { std::string combined; - combined += Program.getSource(); + // When available, hash the preprocessed translation unit instead of the raw + // source: it embeds the content of every #included header (including ones + // resolved from -I filesystem paths), so header edits invalidate the cache. + combined += PreprocessedSource ? *PreprocessedSource : Program.getSource(); combined += "\n---headers---\n"; // std::map is sorted by key, so iteration order is deterministic for (auto &[name, content] : Program.getHeaders()) { @@ -632,22 +701,63 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram Prog, int NumOptions, try { auto &Program = *(chipstar::Program *)Prog; - // Check HIPRTC output cache before invoking clang. - auto cacheKey = computeHiprtcCacheKey(Program, NumOptions, Options); - auto t0 = std::chrono::steady_clock::now(); - if (loadHiprtcCache(Program, cacheKey)) { - auto t1 = std::chrono::steady_clock::now(); - double elapsed = std::chrono::duration(t1 - t0).count(); - logInfo("hiprtc: Cache hit — skipped clang compilation ({:.3f}s saved)", elapsed); - return HIPRTC_SUCCESS; + // Only preprocess when the source may pull in header content. For a + // self-contained kernel (no #include/#import) the raw source is a complete + // cache key, so we keep the original zero-I/O fast path: no temp dir, no + // preprocess subprocess. Sources that do include headers pay a preprocess so + // the key reflects header content reached via -I filesystem paths. + std::optional TmpDir; + std::optional Preprocessed; + bool CacheUsable = true; + if (sourceMayIncludeHeaders(Program.getSource())) { + TmpDir = createTemporaryDirectory(); + if (!TmpDir) { + logError( + "hiprtc: Failed to create a temporary directory for compilation."); + return HIPRTC_ERROR_COMPILATION; + } + CompileOptions PPOptions; + processOptions(Program, NumOptions, Options, PPOptions); + Preprocessed = preprocessForCacheKey(Program, PPOptions, *TmpDir); + if (!Preprocessed) { + // We could not build a key that reflects #include content. Keying on + // the raw source instead would ignore header edits and could serve + // stale SPIR-V — the exact bug #1335 is about — so disable the cache + // for this compilation entirely: no lookup, and no new entry written. + logWarn("hiprtc: could not preprocess source for the cache key; " + "compiling without caching for this program."); + CacheUsable = false; + } + } + + // Check the HIPRTC output cache before invoking clang, unless caching was + // disabled above because we have no trustworthy key. + std::string cacheKey; + if (CacheUsable) { + cacheKey = computeHiprtcCacheKey(Program, NumOptions, Options, + Preprocessed); + auto t0 = std::chrono::steady_clock::now(); + if (loadHiprtcCache(Program, cacheKey)) { + auto t1 = std::chrono::steady_clock::now(); + double elapsed = std::chrono::duration(t1 - t0).count(); + logInfo("hiprtc: Cache hit — skipped clang compilation ({:.3f}s saved)", + elapsed); + if (TmpDir && !ChipEnvVars.getSaveTemps()) { + std::error_code IgnoreErrors; + fs::remove_all(*TmpDir, IgnoreErrors); + } + return HIPRTC_SUCCESS; + } } - // Create temporary directory for compilation I/O. - auto TmpDir = createTemporaryDirectory(); + // Miss: ensure a temp dir exists for the real compilation. if (!TmpDir) { - logError( - "hiprtc: Failed to create a temporary directory for compilation."); - return HIPRTC_ERROR_COMPILATION; + TmpDir = createTemporaryDirectory(); + if (!TmpDir) { + logError( + "hiprtc: Failed to create a temporary directory for compilation."); + return HIPRTC_ERROR_COMPILATION; + } } logDebug("hiprtc: Temp directory: '{}'", TmpDir->string()); @@ -682,8 +792,9 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram Prog, int NumOptions, fs::remove_all(*TmpDir, IgnoreErrors); } - // Cache the compiled SPIRV for future runs. - if (Result == HIPRTC_SUCCESS) + // Cache the compiled SPIRV for future runs, unless caching was disabled + // because we could not compute a reliable key for this program. + if (Result == HIPRTC_SUCCESS && CacheUsable) saveHiprtcCache(Program, cacheKey); return Result; From a53cb5d0ae1d488f0877f068d74bd8eb94eba3ed Mon Sep 17 00:00:00 2001 From: Build Script Date: Mon, 20 Jul 2026 07:23:45 -0700 Subject: [PATCH 2/3] tests: -I #include content participates in the hipRTC cache key Adds TestHiprtcCacheIncludes: compiles the same source with the same options four times, changing only the content of a header reached via -I, and reads HIPRTC's own cache decision (hit vs. store) from the compile output: header A -> miss ; A again -> hit ; header B -> miss ; A again -> hit The compile-2 hit proves caching is exercised (the test cannot pass vacuously if caching never happens); the compile-3 miss proves the header content is in the key; the compile-4 hit proves the key tracks content, not timestamps. Compile-only, so it needs no device. Registered through a generic run wrapper (tests/run_test_with_fresh_cache.cmake) that provisions a fresh mktemp'd cache directory per run and sets CHIP_MODULE_CACHE_DIR/CHIP_LOGLEVEL before the process starts (libCHIP reads them at static init). A never-before-used cache name cannot hold a prior run's entries, so the first compile is reliably a cold miss. Verified non-vacuous: with the header omitted from the key, compile 3 hits and the test fails. --- tests/hiprtc/CMakeLists.txt | 17 ++ tests/hiprtc/TestHiprtcCacheIncludes.cc | 198 ++++++++++++++++++++++++ tests/run_test_with_fresh_cache.cmake | 41 +++++ 3 files changed, 256 insertions(+) create mode 100644 tests/hiprtc/TestHiprtcCacheIncludes.cc create mode 100644 tests/run_test_with_fresh_cache.cmake diff --git a/tests/hiprtc/CMakeLists.txt b/tests/hiprtc/CMakeLists.txt index 07a2a5bae..cc099ee93 100644 --- a/tests/hiprtc/CMakeLists.txt +++ b/tests/hiprtc/CMakeLists.txt @@ -36,3 +36,20 @@ add_hip_test(TestHiprtcCacheNameExpressions.cc) # repeat ctest invocations always start from a known-empty cache. set_tests_properties(TestHiprtcCacheNameExpressions PROPERTIES ENVIRONMENT "CHIP_MODULE_CACHE_DIR=${CMAKE_CURRENT_BINARY_DIR}/test_cache_TestHiprtcCacheNameExpressions") + +# TestHiprtcCacheIncludes needs its cache env set before the process starts (see +# tests/run_test_with_fresh_cache.cmake), so instead of add_hip_test's plain +# command we build the executable the same way and register it through the run +# wrapper that provisions a fresh, unique cache directory per run. +add_executable(TestHiprtcCacheIncludes EXCLUDE_FROM_ALL TestHiprtcCacheIncludes.cc) +set_target_properties(TestHiprtcCacheIncludes PROPERTIES CXX_STANDARD_REQUIRED ON) +target_link_libraries(TestHiprtcCacheIncludes CHIP deviceInternal) +target_include_directories(TestHiprtcCacheIncludes + PUBLIC ${CMAKE_SOURCE_DIR}/HIP/include ${CMAKE_SOURCE_DIR}/include) +add_dependencies(build_tests TestHiprtcCacheIncludes) +add_test(NAME TestHiprtcCacheIncludes + COMMAND ${CMAKE_COMMAND} + -DTEST_EXECUTABLE=$ + -P ${CMAKE_SOURCE_DIR}/tests/run_test_with_fresh_cache.cmake) +set_tests_properties(TestHiprtcCacheIncludes PROPERTIES + SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") diff --git a/tests/hiprtc/TestHiprtcCacheIncludes.cc b/tests/hiprtc/TestHiprtcCacheIncludes.cc new file mode 100644 index 000000000..8b1aa629b --- /dev/null +++ b/tests/hiprtc/TestHiprtcCacheIncludes.cc @@ -0,0 +1,198 @@ +/* + * Copyright (c) 2026 chipStar developers + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included + * in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS IN THE SOFTWARE. + */ + +// The content of a header included via an -I search path must participate in the +// HIPRTC compilation cache key: with the source text and compile options held +// constant, a compile is served from cache when (and only when) the -I header +// content matches a previously compiled one. +// +// This is a compile-only test: it never launches a kernel, so it needs no device +// and can run wherever the build runs. It compiles the same source four times +// with the same options, changing only the content of an -I header, and reads +// HIPRTC's own cache decision (hit vs. store) from the compile's stderr output: +// +// compile 1 header = A -> miss (cold: A's content is new) +// compile 2 header = A -> hit (unchanged: A's entry is reused) +// compile 3 header = B -> miss (only the header changed: different key) +// compile 4 header = A -> hit (A's entry is still distinct from B's) +// +// Compile 2's hit proves caching is genuinely exercised (the test cannot pass +// vacuously if caching never happens). Compile 3's miss proves the header +// content is in the key. Compile 4's hit proves the key is a function of the +// header's CONTENT: had the key keyed on, say, the header's timestamp, reverting +// the content would still be a miss (a fresh timestamp), never a hit. +// +// Requirements (provided by the harness, not by this program): +// - CHIP_MODULE_CACHE_DIR points at a cache directory that is EMPTY at start, +// so compile 1 is a cold miss. CMakeLists.txt wipes a dedicated directory via +// a ctest fixture before the test runs. +// - CHIP_LOGLEVEL=info, so HIPRTC's hit/miss reporting is emitted. +// libCHIP reads both at static-init time, before main(); CMakeLists.txt sets them +// for ctest. + +#include "TestCommon.hh" + +#include +#include +#include +#include +#include +#include +#include + +namespace fs = std::filesystem; + +// Fixed across every compile. The only thing that varies between compiles is the +// value defined in the -I header, so the raw source text and options are +// constant while the preprocessed translation unit (and thus the key) differs. +static constexpr auto Source = R"---( +#include "hiprtc_cache_inc.h" +extern "C" __global__ void addConst(int *Out, const int *In) { + *Out = *In + HIPRTC_CACHE_TEST_ADDEND; +} +)---"; + +static const char *IncludeHeaderName = "hiprtc_cache_inc.h"; + +struct CacheDecision { + bool Hit; // HIPRTC served this compile from cache + bool Touched; // HIPRTC reported any cache activity (hit or store) +}; + +// Overwrite the -I header so it defines HIPRTC_CACHE_TEST_ADDEND to Value. The +// numeric value only needs to differ to make the header content differ. +static void writeHeader(const fs::path &HeaderPath, int Value) { + std::ofstream F(HeaderPath, std::ios::binary | std::ios::trunc); + TEST_ASSERT(F.is_open()); + F << "#define HIPRTC_CACHE_TEST_ADDEND " << Value << "\n"; + F.close(); + TEST_ASSERT(F.good()); +} + +// Compile the fixed Source with -I and report HIPRTC's cache decision, +// read from the diagnostics it writes to stderr during the compile. spdlog's +// sink is stderr (fd 2); we redirect it to a temp file for the compile, then +// restore it and echo the captured text so it still appears in the test log. +static CacheDecision compileAndGetCacheDecision(const fs::path &IncDir) { + hiprtcProgram Prog; + HIPRTC_CHECK(hiprtcCreateProgram(&Prog, Source, "cache_inc.hip", 0, nullptr, + nullptr)); + + std::string IncludeOpt = "-I" + IncDir.string(); + std::vector Options = {IncludeOpt.c_str()}; + + std::fflush(stderr); + int SavedStderr = dup(STDERR_FILENO); + TEST_ASSERT(SavedStderr >= 0); + std::FILE *Cap = std::tmpfile(); + TEST_ASSERT(Cap != nullptr); + TEST_ASSERT(dup2(fileno(Cap), STDERR_FILENO) >= 0); + + hiprtcResult R = hiprtcCompileProgram(Prog, Options.size(), Options.data()); + + std::fflush(stderr); + TEST_ASSERT(dup2(SavedStderr, STDERR_FILENO) >= 0); + close(SavedStderr); + + std::string Captured; + std::fseek(Cap, 0, SEEK_END); + long N = std::ftell(Cap); + if (N > 0) { + std::fseek(Cap, 0, SEEK_SET); + Captured.resize(static_cast(N)); + size_t Got = std::fread(&Captured[0], 1, static_cast(N), Cap); + Captured.resize(Got); + } + std::fclose(Cap); + std::cerr << Captured; + + if (R != HIPRTC_SUCCESS) { + size_t LogSize = 0; + HIPRTC_CHECK(hiprtcGetProgramLogSize(Prog, &LogSize)); + if (LogSize) { + std::string Log(LogSize, '\0'); + hiprtcGetProgramLog(Prog, Log.data()); + std::cerr << Log << "\n"; + } + HIPRTC_CHECK(R); + } + HIPRTC_CHECK(hiprtcDestroyProgram(&Prog)); + + bool Hit = Captured.find("Cache hit") != std::string::npos; + bool Stored = Captured.find("Saved SPIRV to cache") != std::string::npos; + bool Loaded = Captured.find("Loaded SPIRV from cache") != std::string::npos; + return {Hit, (Hit || Stored || Loaded)}; +} + +int main() { + // The cache must be enabled for this test to mean anything; the harness points + // CHIP_MODULE_CACHE_DIR at a directory that is empty at start (see the file + // header). We do not manage that directory here. + const char *CacheDir = std::getenv("CHIP_MODULE_CACHE_DIR"); + if (!CacheDir || !*CacheDir) { + std::cerr << "CHIP_MODULE_CACHE_DIR is not set; this test needs an (empty) " + "cache directory. The CMake build provides one for ctest.\n"; + return 1; + } + + // Private -I directory holding the header we mutate between compiles. + std::error_code Ec; + fs::path IncDir = fs::temp_directory_path() / "chipstar_hiprtc_cache_inc_test"; + fs::remove_all(IncDir, Ec); + TEST_ASSERT(fs::create_directories(IncDir, Ec)); + fs::path HeaderPath = IncDir / IncludeHeaderName; + + // Compile 1: header content A. Cold cache -> miss (store). + writeHeader(HeaderPath, 111); + CacheDecision D1 = compileAndGetCacheDecision(IncDir); + std::cerr << "compile #1 (header A): hit=" << D1.Hit << "\n"; + // Cache must actually be exercised, otherwise the remaining checks are vacuous; + // and a fresh cache means this cold compile is a miss (also catches a stale, + // non-empty cache directory). + TEST_ASSERT(D1.Touched); + TEST_ASSERT(!D1.Hit); + + // Compile 2: identical source, options, and header content -> hit. This is what + // proves caching is genuinely happening and the key is stable for equal input. + CacheDecision D2 = compileAndGetCacheDecision(IncDir); + std::cerr << "compile #2 (header A, unchanged): hit=" << D2.Hit << "\n"; + TEST_ASSERT(D2.Hit); + + // Compile 3: only the -I header content changes (B). The key reflects header + // content, so this is a miss. + writeHeader(HeaderPath, 222); + CacheDecision D3 = compileAndGetCacheDecision(IncDir); + std::cerr << "compile #3 (header B): hit=" << D3.Hit << "\n"; + TEST_ASSERT(!D3.Hit); + + // Compile 4: revert to header content A. Its entry is distinct from B's, so + // this is a hit -- proving the key tracks header content, not (e.g.) mtime. + writeHeader(HeaderPath, 111); + CacheDecision D4 = compileAndGetCacheDecision(IncDir); + std::cerr << "compile #4 (header A, reverted): hit=" << D4.Hit << "\n"; + TEST_ASSERT(D4.Hit); + + fs::remove_all(IncDir, Ec); + std::cerr << "Test passed: -I #include content participates in the HIPRTC " + "cache key.\n"; + return 0; +} diff --git a/tests/run_test_with_fresh_cache.cmake b/tests/run_test_with_fresh_cache.cmake new file mode 100644 index 000000000..252f4fe53 --- /dev/null +++ b/tests/run_test_with_fresh_cache.cmake @@ -0,0 +1,41 @@ +# Generic run wrapper: execute a test with a fresh, uniquely named module cache +# directory (CHIP_MODULE_CACHE_DIR) and info-level logging (CHIP_LOGLEVEL=info). +# +# Some tests must observe cache behavior from a known-empty starting state. But +# libCHIP reads CHIP_MODULE_CACHE_DIR and CHIP_LOGLEVEL at static-init time, +# before main(), so they must be set before the test process starts -- which a +# run-time wrapper allows and a fixed configure-time ENVIRONMENT value cannot +# randomize per run. A `mktemp -d` name has never been used before, so it cannot +# contain a prior run's cache entries regardless of how the runtime canonicalizes +# the path (case, symlinks, ...); the first compile is thus reliably a cold miss. +# The directory is removed afterwards. +# +# Invoked as: +# cmake -DTEST_EXECUTABLE= -P run_test_with_fresh_cache.cmake + +if(NOT TEST_EXECUTABLE) + message(FATAL_ERROR "TEST_EXECUTABLE not set") +endif() +get_filename_component(TEST_NAME "${TEST_EXECUTABLE}" NAME) + +execute_process( + COMMAND mktemp -d + OUTPUT_VARIABLE CACHE_DIR + OUTPUT_STRIP_TRAILING_WHITESPACE + RESULT_VARIABLE MKTEMP_RC) +if(NOT MKTEMP_RC EQUAL 0) + message(FATAL_ERROR "mktemp -d failed (${MKTEMP_RC})") +endif() + +execute_process( + COMMAND ${CMAKE_COMMAND} -E env + "CHIP_MODULE_CACHE_DIR=${CACHE_DIR}" + "CHIP_LOGLEVEL=info" + "${TEST_EXECUTABLE}" + RESULT_VARIABLE TEST_RC) + +file(REMOVE_RECURSE "${CACHE_DIR}") + +if(NOT TEST_RC EQUAL 0) + message(FATAL_ERROR "${TEST_NAME} failed (exit ${TEST_RC})") +endif() From 17639c0402935b883cc3cd6fe4f200d0062572ac Mon Sep 17 00:00:00 2001 From: Build Script Date: Wed, 22 Jul 2026 19:07:52 -0700 Subject: [PATCH 3/3] fix: surface preprocessor diagnostics when cache-key preprocessing fails When the preprocess-only pass used to build the hipRTC cache key exits non-zero, preprocessForCacheKey returned nullopt and the caller reported only "could not preprocess source for the cache key" -- the compiler's own error, written to the temp log file, was discarded with the temp dir. Read that log back and emit it via logWarn on failure so the actual reason is visible. This is diagnostic-only: behavior on failure is unchanged (caching stays disabled for that program, never stale). --- src/spirv_hiprtc.cc | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/src/spirv_hiprtc.cc b/src/spirv_hiprtc.cc index 27258634e..552c66e93 100644 --- a/src/spirv_hiprtc.cc +++ b/src/spirv_hiprtc.cc @@ -289,8 +289,17 @@ preprocessForCacheKey(const chipstar::Program &Program, std::string Cmd = createCompileCommand(Options, WorkingDirectory, SourceFile, OutputFile, /*PreprocessOnly=*/true); - if (!executeCommand(WorkingDirectory, Cmd, LogFile)) + if (!executeCommand(WorkingDirectory, Cmd, LogFile)) { + // The caller only reports that caching was disabled; surface the compiler's + // own diagnostics here so the reason (e.g. a toolchain rejecting the + // preprocess-only invocation) is visible instead of silently swallowed. + if (auto Log = readFromFile(LogFile); Log && !Log->empty()) + logWarn("hiprtc: preprocessing for the cache key failed:\n{}", *Log); + else + logWarn("hiprtc: preprocessing for the cache key failed (no compiler " + "output captured)."); return std::nullopt; + } return readFromFile(OutputFile); }