diff --git a/src/spirv_hiprtc.cc b/src/spirv_hiprtc.cc index 770fe8f63..552c66e93 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,64 @@ 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)) { + // 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); +} + +// 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 +483,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 +710,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; + } } - // Create temporary directory for compilation I/O. - auto TmpDir = createTemporaryDirectory(); + // 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; + } + } + + // 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 +801,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; 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()