From f9fe90c460c50644de6e6234ea8aabed3b519cf8 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Mon, 23 Feb 2026 14:48:59 +0530 Subject: [PATCH 01/10] [HIPIFY][feature] Local header hipification with preceding include injection via ArgumentsAdjusters --- src/LocalHeader.cpp | 89 +++++++++++++++++-- src/LocalHeader.h | 25 +++++- src/main.cpp | 24 ++++- .../local_headers/injection_has_cmath.h | 16 ++++ .../headers/local_headers/injection_helper.h | 22 +++++ .../headers/local_headers/injection_inner.h | 14 +++ .../headers/local_headers/injection_multi.cu | 22 +++++ .../local_headers/injection_no_duplicate.cu | 14 +++ .../headers/local_headers/injection_outer.h | 16 ++++ .../local_headers/injection_pragma_header.h | 11 +++ .../local_headers/injection_pragma_once.cu | 17 ++++ .../local_headers/injection_recursive.cu | 17 ++++ .../headers/local_headers/injection_test.cu | 17 ++++ .../local_headers/injection_uses_cmath.h | 18 ++++ .../headers/local_headers/vector_math.h | 12 +++ 15 files changed, 320 insertions(+), 14 deletions(-) create mode 100644 tests/unit_tests/headers/local_headers/injection_has_cmath.h create mode 100644 tests/unit_tests/headers/local_headers/injection_helper.h create mode 100644 tests/unit_tests/headers/local_headers/injection_inner.h create mode 100644 tests/unit_tests/headers/local_headers/injection_multi.cu create mode 100644 tests/unit_tests/headers/local_headers/injection_no_duplicate.cu create mode 100644 tests/unit_tests/headers/local_headers/injection_outer.h create mode 100644 tests/unit_tests/headers/local_headers/injection_pragma_header.h create mode 100644 tests/unit_tests/headers/local_headers/injection_pragma_once.cu create mode 100644 tests/unit_tests/headers/local_headers/injection_recursive.cu create mode 100644 tests/unit_tests/headers/local_headers/injection_test.cu create mode 100644 tests/unit_tests/headers/local_headers/injection_uses_cmath.h create mode 100644 tests/unit_tests/headers/local_headers/vector_math.h diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 2d5e97545..056f89024 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -1,4 +1,5 @@ #include "LocalHeader.h" +#include "LLVMCompat.h" #include #include @@ -11,8 +12,6 @@ #include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/raw_ostream.h" -#include "LLVMCompat.h" - using namespace clang; using namespace clang::tooling; using namespace llvm; @@ -46,6 +45,9 @@ namespace { static const std::regex LocalIncludeRe( R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?$)", std::regex::ECMAScript); + static const std::regex SystemIncludeRe( + R"(^\s*#\s*include\s*<([^>\n]+)>)", std::regex::ECMAScript); + bool readFile(const std::string &path, std::string &out) { auto MBOrErr = llvm::MemoryBuffer::getFile(path); if (!MBOrErr) return false; @@ -67,7 +69,41 @@ namespace { } return false; } -} + + bool collectPrecedingIncludes(const std::string &mainSourceAbspath, + const std::string &targetHeaderAbspath, + std::vector &outIncludes) { + std::string mainSourceContent; + if (!readFile(mainSourceAbspath, mainSourceContent)) { + errs() << sHipify << sError << "Cannot read source files: " + << mainSourceAbspath << "\n"; + return false; + } + + std::string targetFileName = std::string(sys::path::filename(targetHeaderAbspath)); + std::istringstream iss(mainSourceContent); + std::string line; + std::smatch m; + + while (std::getline(iss, line)) { + if (std::regex_match(line, m, LocalIncludeRe)) { + std::string quotedName = m[1].str(); + std::string quotedFileName = std::string(sys::path::filename(quotedName)); + if (quotedFileName == targetFileName) + break; + std::string absPath; + if (resolveLocalIncludeInternal(mainSourceAbspath, quotedName, absPath)) + outIncludes.push_back(absPath); + } + + if (std::regex_search(line, m, SystemIncludeRe)) { + outIncludes.push_back(m[1].str()); + } + } + + return true; + } +} bool resolveLocalInclude(const std::string &mainSourceAbsPath, const std::string &includeToken, @@ -116,52 +152,87 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, } if (initial.empty()) { - outs() << sHipify << "No local headers detected in " << mainSourceAbsPath << "\n"; + outs() << "\n" << sHipify << "No local headers detected in " + << sys::path::filename(mainSourceAbsPath) << "\n"; return true; } + outs() << "\n" << sHipify << "Local headers found: " << initial.size() + << " in " << sys::path::filename(mainSourceAbsPath) << "\n"; + for (size_t i = 0; i < initial.size(); ++i) { + outs() << (i + 1) << "/" << initial.size() + << ": " << sys::path::filename(initial[i]) << "\n"; + } + std::vector work(initial.begin(), initial.end()); std::set processed; + size_t total = initial.size(); + size_t current = 0; while (!work.empty()) { std::string hdr = work.back(); work.pop_back(); if (processed.count(hdr)) { - errs() << sHipify << sWarning << "Duplicate local header reference ignored: " << hdr << "\n"; + outs() << sHipify << sWarning + << "Duplicate local header reference ignored: " + << sys::path::filename(hdr) << "\n"; continue; } processed.insert(hdr); + ++current; std::string original; if (!readFile(hdr, original)) { - errs() << sHipify << sError << "Cannot read header: " << hdr << "\n"; + errs() << "\n" << sHipify << sError + << "Cannot read header: " << sys::path::filename(hdr) << "\n"; continue; } std::string hipOut = hdr + ".hip"; + std::vector precedingIncludes; + collectPrecedingIncludes(mainSourceAbsPath, hdr, precedingIncludes); + + outs() << "\n" << sHipify << "Hipifying local header [" << current + << "/" << total << "]: " << sys::path::filename(hdr) << "\n"; + bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, - hipify_exe, mainSourceAbsPath, false); + hipify_exe, mainSourceAbsPath, false, + precedingIncludes); if (!ok) { - errs() << sHipify << sError << "Hipify failed for header: " << hdr << "\n"; + errs() << "\n" << sHipify << sError + << "Hipify failed for header [" << current << "/" << total + << "]: " << sys::path::filename(hdr) << "\n"; return false; } + outs() << sHipify << "Successfully hipified header file" << "\n"; if (recursive) { std::smatch m; std::istringstream iss(original); std::string line; + std::vector newHeaders; while (std::getline(iss, line)) { if (std::regex_match(line, m, LocalIncludeRe)) { std::string rel = m[1].str(); std::string abs; if (resolveLocalIncludeInternal(hdr, rel, abs) && - !processed.count(abs)) + !processed.count(abs)) { work.push_back(abs); + newHeaders.push_back(abs); + } } } + if (!newHeaders.empty()) { + total += newHeaders.size(); + outs() << sHipify << " Recursive: found " << newHeaders.size() + << " additional local header(s) in " + << sys::path::filename(hdr) << "\n"; + } } } + outs() << "\n" << sHipify << "Local header hipification complete: " + << processed.size() << " header(s) processed.\n"; return true; } diff --git a/src/LocalHeader.h b/src/LocalHeader.h index 31763c697..097c1df61 100644 --- a/src/LocalHeader.h +++ b/src/LocalHeader.h @@ -1,3 +1,25 @@ +/* +Copyright (c) 2026 - present Advanced Micro Devices, Inc. All rights reserved. + +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. +*/ + #pragma once #include @@ -12,7 +34,8 @@ extern bool hipifySingleSource(const std::string &srcPath, ct::CommonOptionsParser *OptionsParserPtr, const char *hipify_exe_path, const std::string &mainContextPath, - bool preserveTemp); + bool preserveTemp, + const std::vector &additionalIncludes = {}); bool hipifyLocalHeaders(const std::string &srcPath, const ct::CompilationDatabase *compDB, diff --git a/src/main.cpp b/src/main.cpp index 0675bdfce..8c23a1085 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -21,6 +21,7 @@ THE SOFTWARE. */ #include +#include #include "CUDA2HIP.h" #include "CUDA2HIP_Scripting.h" #include "LLVMCompat.h" @@ -241,7 +242,8 @@ bool hipifySingleSource(const std::string &srcPath, ct::CommonOptionsParser *OptionsParserPtr, const char *hipify_exe_path, const std::string &mainContextPath, - bool preserveTemp) { + bool preserveTemp, + const std::vector &additionalIncludes) { std::error_code EC; SmallString<128> tmpFile; StringRef srcFileName = sys::path::filename(srcPath); @@ -276,6 +278,15 @@ bool hipifySingleSource(const std::string &srcPath, return false; } + for (auto it = additionalIncludes.rbegin(); it != additionalIncludes.rend(); ++it) { + Tool.appendArgumentsAdjuster( + ct::getInsertArgumentAdjuster(it->c_str(), + ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster( + ct::getInsertArgumentAdjuster("-include", + ct::ArgumentInsertPosition::BEGIN)); + } + // Hipify _all_ the things! if (Tool.runAndSave(&actionFactory)) { llvm::errs() << "\n" << sHipify << sError @@ -516,15 +527,20 @@ int main(int argc, const char **argv) { } // Initialise the statistics counters for this file. Statistics::setActive(src); - // Checks the local headers if --local-headers/--local-header-recursive specified. if (OptLocalHeaders || OptLocalHeadersRecursive) { + llvm::outs() << "\n" << sHipify + << "Local header hipification enabled (" + << (OptLocalHeadersRecursive ? "recursive" : "non-recursive") + << ") for: " << sys::path::filename(sSourceAbsPath) << "\n"; if (!hipifyLocalHeaders(sSourceAbsPath, compilationDatabase.get(), &OptionsParser, argv[0], OptLocalHeadersRecursive)) { + llvm::errs() << "\n" << sHipify << sError + << "Local header hipification failed for: " + << sys::path::filename(sSourceAbsPath) << "\n"; Statistics::current().hasErrors = true; - LLVM_DEBUG(llvm::dbgs() << "Local header hipification failed for: " << sSourceAbsPath << "\n"); Result = 1; } } @@ -535,7 +551,7 @@ int main(int argc, const char **argv) { &OptionsParser, argv[0], sSourceAbsPath, - false)) { + false, {})) { Statistics::current().hasErrors = true; Result = 1; LLVM_DEBUG(llvm::dbgs() << "Hipification failed for: " << src << "\n"); diff --git a/tests/unit_tests/headers/local_headers/injection_has_cmath.h b/tests/unit_tests/headers/local_headers/injection_has_cmath.h new file mode 100644 index 000000000..5fddca53f --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_has_cmath.h @@ -0,0 +1,16 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +#ifndef INJECTION_HAS_CMATH_H +#define INJECTION_HAS_CMATH_H + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +#include +#include + +inline float compute_value(float x) { + return sqrtf(x) + 1.0f; +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/injection_helper.h b/tests/unit_tests/headers/local_headers/injection_helper.h new file mode 100644 index 000000000..5185c6e78 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_helper.h @@ -0,0 +1,22 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +#ifndef INJECTION_HELPER_H +#define INJECTION_HELPER_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline __device__ float3 add_vectors(float3 a, float3 b) { + return make_float3(0.0f, 0.0f, 0.0f); +} + +inline __device__ void accumulate(float3* sum, float3 val) { + return; +} + +inline __device__ float3 scale_and_diff(float3 a, float3 b, float s) { + return make_float3(0.0f, 0.0f, 0.0f); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/injection_inner.h b/tests/unit_tests/headers/local_headers/injection_inner.h new file mode 100644 index 000000000..d0ebb3fcf --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_inner.h @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args + +#ifndef INJECTION_INNER_H +#define INJECTION_INNER_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline __device__ void inner_add(float3* data, int idx) { + return; +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/injection_multi.cu b/tests/unit_tests/headers/local_headers/injection_multi.cu new file mode 100644 index 000000000..20341f306 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_multi.cu @@ -0,0 +1,22 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +// CHECK: #include "vector_math.h" +// CHECK: #include "injection_helper.h" +// CHECK: #include "injection_uses_cmath.h" +#include +#include +#include "vector_math.h" +#include "injection_helper.h" +#include "injection_uses_cmath.h" + +__global__ void multiKernel(float3* data, float* vals) { + int idx = threadIdx.x; + dummy_vector_op(); + add_vectors(data[idx], data[idx + 1]); + vals[idx] = compute_sqrt(vals[idx]); +} + +int main() { return 0; } diff --git a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu new file mode 100644 index 000000000..195ecddf6 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +// CHECK: #include "injection_has_cmath.h" +#include +#include +#include "injection_has_cmath.h" + +int main() { + float x = compute_value(4.0f); + return (int)x; +} diff --git a/tests/unit_tests/headers/local_headers/injection_outer.h b/tests/unit_tests/headers/local_headers/injection_outer.h new file mode 100644 index 000000000..5ae240f14 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_outer.h @@ -0,0 +1,16 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args + +#ifndef INJECTION_OUTER_H +#define INJECTION_OUTER_H + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "injection_inner.h" +#include +#include "injection_inner.h" + +inline __device__ void outer_process(float3* data, int idx) { + inner_add(data, idx); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/injection_pragma_header.h b/tests/unit_tests/headers/local_headers/injection_pragma_header.h new file mode 100644 index 000000000..09792ce53 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_pragma_header.h @@ -0,0 +1,11 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +#pragma once + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline __device__ void pragma_add(float3* data, int idx) { + return; +} diff --git a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu new file mode 100644 index 000000000..f9643f1e4 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "vector_math.h" +// CHECK: #include "injection_pragma_header.h" +#include +#include "vector_math.h" +#include "injection_pragma_header.h" + +__global__ void pragmaKernel(float3* data) { + int idx = threadIdx.x; + dummy_vector_op(); + pragma_add(data, idx); +} + +int main() { return 0; } diff --git a/tests/unit_tests/headers/local_headers/injection_recursive.cu b/tests/unit_tests/headers/local_headers/injection_recursive.cu new file mode 100644 index 000000000..dfa6ef52d --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_recursive.cu @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "vector_math.h" +// CHECK: #include "injection_outer.h" +#include +#include "vector_math.h" +#include "injection_outer.h" + +__global__ void recursiveKernel(float3* data) { + int idx = threadIdx.x; + dummy_vector_op(); + outer_process(data, idx); +} + +int main() { return 0; } diff --git a/tests/unit_tests/headers/local_headers/injection_test.cu b/tests/unit_tests/headers/local_headers/injection_test.cu new file mode 100644 index 000000000..ced4b5066 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_test.cu @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "vector_math.h" +// CHECK: #include "injection_helper.h" +#include +#include "vector_math.h" +#include "injection_helper.h" + +__global__ void testKernel(float3* data) { + int idx = threadIdx.x; + dummy_vector_op(); + add_vectors(data[idx], data[idx + 1]); +} + +int main() { return 0; } diff --git a/tests/unit_tests/headers/local_headers/injection_uses_cmath.h b/tests/unit_tests/headers/local_headers/injection_uses_cmath.h new file mode 100644 index 000000000..d7b5b6a54 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_uses_cmath.h @@ -0,0 +1,18 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +#ifndef INJECTION_USES_CMATH_H +#define INJECTION_USES_CMATH_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline __device__ float compute_sqrt(float x) { + return sqrtf(x); +} + +inline __device__ float compute_magnitude(float x, float y) { + return sqrtf(x * x + y * y); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/vector_math.h b/tests/unit_tests/headers/local_headers/vector_math.h new file mode 100644 index 000000000..12ed82cbe --- /dev/null +++ b/tests/unit_tests/headers/local_headers/vector_math.h @@ -0,0 +1,12 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args + +#ifndef VECTOR_MATH_H +#define VECTOR_MATH_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline __host__ __device__ void dummy_vector_op() { return; } + +#endif From 1cca1b9fc894605045b5d5ad061a474ebe586371 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Wed, 25 Feb 2026 11:42:54 +0530 Subject: [PATCH 02/10] [HIPIFY][doc][NFC] Updated the licence --- src/LocalHeader.cpp | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 056f89024..7c76b7929 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -1,3 +1,25 @@ +/* +Copyright (c) 2026 - present Advanced Micro Devices, Inc. All rights reserved. + +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. +*/ + #include "LocalHeader.h" #include "LLVMCompat.h" From 89e87c1fc994169572a5a25702f3f4b71dce8520 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Wed, 4 Mar 2026 11:50:37 +0530 Subject: [PATCH 03/10] [HIPIFY][feature] Implicit hipification of local header files --- src/ArgParse.cpp | 12 +++--------- src/ArgParse.h | 3 +-- src/LocalHeader.cpp | 13 ++++++++----- src/main.cpp | 8 ++------ .../headers/local_headers/angle_bracket.cu | 2 +- tests/unit_tests/headers/local_headers/common.h | 2 +- .../headers/local_headers/cyclic_header.cu | 2 +- tests/unit_tests/headers/local_headers/duplicate.h | 2 +- .../headers/local_headers/duplicate_header.cu | 2 +- .../headers/local_headers/injection_has_cmath.h | 2 +- .../headers/local_headers/injection_helper.h | 2 +- .../headers/local_headers/injection_inner.h | 2 +- .../headers/local_headers/injection_multi.cu | 2 +- .../headers/local_headers/injection_no_duplicate.cu | 2 +- .../headers/local_headers/injection_outer.h | 2 +- .../headers/local_headers/injection_pragma_header.h | 2 +- .../headers/local_headers/injection_pragma_once.cu | 2 +- .../headers/local_headers/injection_recursive.cu | 2 +- .../headers/local_headers/injection_test.cu | 2 +- .../headers/local_headers/injection_uses_cmath.h | 2 +- .../unit_tests/headers/local_headers/local_angle.h | 2 +- tests/unit_tests/headers/local_headers/main.cu | 2 +- tests/unit_tests/headers/local_headers/mixed.h | 2 +- .../headers/local_headers/no_local_headers.cu | 2 +- tests/unit_tests/headers/local_headers/non_cuda.h | 2 +- .../headers/local_headers/non_cuda_header.cu | 2 +- tests/unit_tests/headers/local_headers/rec_1.h | 2 +- .../headers/local_headers/recursive-header.cu | 2 +- .../headers/local_headers/relative/common1.h | 2 +- .../headers/local_headers/relative/sub/common2.h | 2 +- .../headers/local_headers/relative_path.cu | 2 +- tests/unit_tests/headers/local_headers/shared.h | 2 +- .../headers/local_headers/shared_header_1.cu | 2 +- .../headers/local_headers/shared_header_2.cu | 2 +- .../headers/local_headers/single_header.cu | 2 +- .../headers/local_headers/single_header.h | 2 +- .../unit_tests/headers/local_headers/stress_test.cu | 2 +- .../unit_tests/headers/local_headers/vector_math.h | 2 +- 38 files changed, 48 insertions(+), 56 deletions(-) diff --git a/src/ArgParse.cpp b/src/ArgParse.cpp index de10b7f82..f70bf08da 100644 --- a/src/ArgParse.cpp +++ b/src/ArgParse.cpp @@ -226,13 +226,8 @@ cl::opt HipDnnSupport("hipdnn", cl::init(false), cl::cat(ToolTemplateCategory)); -cl::opt OptLocalHeaders("local-headers", - cl::desc("Enable hipification of quoted local headers (non-recursive)"), - cl::init(false), - cl::cat(ToolTemplateCategory)); - -cl::opt OptLocalHeadersRecursive("local-headers-recursive", - cl::desc("Enable hipification of quoted local headers recursively"), +cl::opt SkipLocalHeaders("skip-local-headers", + cl::desc("Skip implicit hipification of local (quoted) headers included in the main source file"), cl::init(false), cl::cat(ToolTemplateCategory)); @@ -264,8 +259,7 @@ const std::vector hipifyOptions { std::string(NoWarningsUndocumented.ArgStr), std::string(HipifyAMAP.ArgStr), std::string(HipDnnSupport.ArgStr), - std::string(OptLocalHeaders.ArgStr), - std::string(OptLocalHeadersRecursive.ArgStr), + std::string(SkipLocalHeaders.ArgStr), }; const std::vector hipifyOptionsWithTwoArgs { diff --git a/src/ArgParse.h b/src/ArgParse.h index 5a32c386f..96405f4b9 100644 --- a/src/ArgParse.h +++ b/src/ArgParse.h @@ -70,5 +70,4 @@ extern cl::opt NoUndocumented; extern cl::opt NoWarningsUndocumented; extern cl::opt HipifyAMAP; extern cl::opt HipDnnSupport; -extern cl::opt OptLocalHeaders; -extern cl::opt OptLocalHeadersRecursive; +extern cl::opt SkipLocalHeaders; diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 7c76b7929..d5d43edcc 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -186,13 +186,16 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, << ": " << sys::path::filename(initial[i]) << "\n"; } - std::vector work(initial.begin(), initial.end()); + std::vector> work; + for (const auto &h : initial) { + work.push_back({h, mainSourceAbsPath}); + } std::set processed; size_t total = initial.size(); size_t current = 0; while (!work.empty()) { - std::string hdr = work.back(); + auto [hdr, parentPath] = work.back(); work.pop_back(); if (processed.count(hdr)) { outs() << sHipify << sWarning @@ -212,13 +215,13 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, std::string hipOut = hdr + ".hip"; std::vector precedingIncludes; - collectPrecedingIncludes(mainSourceAbsPath, hdr, precedingIncludes); + collectPrecedingIncludes(parentPath, hdr, precedingIncludes); outs() << "\n" << sHipify << "Hipifying local header [" << current << "/" << total << "]: " << sys::path::filename(hdr) << "\n"; bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, - hipify_exe, mainSourceAbsPath, false, + hipify_exe, parentPath, false, precedingIncludes); if (!ok) { @@ -240,7 +243,7 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, std::string abs; if (resolveLocalIncludeInternal(hdr, rel, abs) && !processed.count(abs)) { - work.push_back(abs); + work.push_back({abs, hdr}); newHeaders.push_back(abs); } } diff --git a/src/main.cpp b/src/main.cpp index 8c23a1085..6b8a3c219 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -527,16 +527,12 @@ int main(int argc, const char **argv) { } // Initialise the statistics counters for this file. Statistics::setActive(src); - if (OptLocalHeaders || OptLocalHeadersRecursive) { - llvm::outs() << "\n" << sHipify - << "Local header hipification enabled (" - << (OptLocalHeadersRecursive ? "recursive" : "non-recursive") - << ") for: " << sys::path::filename(sSourceAbsPath) << "\n"; + if (!SkipLocalHeaders) { if (!hipifyLocalHeaders(sSourceAbsPath, compilationDatabase.get(), &OptionsParser, argv[0], - OptLocalHeadersRecursive)) { + true)) { llvm::errs() << "\n" << sHipify << sError << "Local header hipification failed for: " << sys::path::filename(sSourceAbsPath) << "\n"; diff --git a/tests/unit_tests/headers/local_headers/angle_bracket.cu b/tests/unit_tests/headers/local_headers/angle_bracket.cu index e9628f061..ef8313513 100644 --- a/tests/unit_tests/headers/local_headers/angle_bracket.cu +++ b/tests/unit_tests/headers/local_headers/angle_bracket.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/common.h b/tests/unit_tests/headers/local_headers/common.h index cb5844475..363b0dd54 100644 --- a/tests/unit_tests/headers/local_headers/common.h +++ b/tests/unit_tests/headers/local_headers/common.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef COMMON_H #define COMMON_H diff --git a/tests/unit_tests/headers/local_headers/cyclic_header.cu b/tests/unit_tests/headers/local_headers/cyclic_header.cu index 3d3e6c692..9ab52d89d 100644 --- a/tests/unit_tests/headers/local_headers/cyclic_header.cu +++ b/tests/unit_tests/headers/local_headers/cyclic_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/duplicate.h b/tests/unit_tests/headers/local_headers/duplicate.h index 030c7dbba..8ef0d1ea2 100644 --- a/tests/unit_tests/headers/local_headers/duplicate.h +++ b/tests/unit_tests/headers/local_headers/duplicate.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef DUP_H #define DUP_H diff --git a/tests/unit_tests/headers/local_headers/duplicate_header.cu b/tests/unit_tests/headers/local_headers/duplicate_header.cu index d07f7f74d..d64a9c084 100644 --- a/tests/unit_tests/headers/local_headers/duplicate_header.cu +++ b/tests/unit_tests/headers/local_headers/duplicate_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK: #include "duplicate.h" diff --git a/tests/unit_tests/headers/local_headers/injection_has_cmath.h b/tests/unit_tests/headers/local_headers/injection_has_cmath.h index 5fddca53f..fcc8ae96e 100644 --- a/tests/unit_tests/headers/local_headers/injection_has_cmath.h +++ b/tests/unit_tests/headers/local_headers/injection_has_cmath.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef INJECTION_HAS_CMATH_H #define INJECTION_HAS_CMATH_H diff --git a/tests/unit_tests/headers/local_headers/injection_helper.h b/tests/unit_tests/headers/local_headers/injection_helper.h index 5185c6e78..150920889 100644 --- a/tests/unit_tests/headers/local_headers/injection_helper.h +++ b/tests/unit_tests/headers/local_headers/injection_helper.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef INJECTION_HELPER_H #define INJECTION_HELPER_H diff --git a/tests/unit_tests/headers/local_headers/injection_inner.h b/tests/unit_tests/headers/local_headers/injection_inner.h index d0ebb3fcf..0c213c7fe 100644 --- a/tests/unit_tests/headers/local_headers/injection_inner.h +++ b/tests/unit_tests/headers/local_headers/injection_inner.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef INJECTION_INNER_H #define INJECTION_INNER_H diff --git a/tests/unit_tests/headers/local_headers/injection_multi.cu b/tests/unit_tests/headers/local_headers/injection_multi.cu index 20341f306..f4dd3c7aa 100644 --- a/tests/unit_tests/headers/local_headers/injection_multi.cu +++ b/tests/unit_tests/headers/local_headers/injection_multi.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu index 195ecddf6..cebb191c1 100644 --- a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu +++ b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_outer.h b/tests/unit_tests/headers/local_headers/injection_outer.h index 5ae240f14..e742b9699 100644 --- a/tests/unit_tests/headers/local_headers/injection_outer.h +++ b/tests/unit_tests/headers/local_headers/injection_outer.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef INJECTION_OUTER_H #define INJECTION_OUTER_H diff --git a/tests/unit_tests/headers/local_headers/injection_pragma_header.h b/tests/unit_tests/headers/local_headers/injection_pragma_header.h index 09792ce53..86f68bc53 100644 --- a/tests/unit_tests/headers/local_headers/injection_pragma_header.h +++ b/tests/unit_tests/headers/local_headers/injection_pragma_header.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #pragma once diff --git a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu index f9643f1e4..67e5e370b 100644 --- a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu +++ b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_recursive.cu b/tests/unit_tests/headers/local_headers/injection_recursive.cu index dfa6ef52d..0d705509c 100644 --- a/tests/unit_tests/headers/local_headers/injection_recursive.cu +++ b/tests/unit_tests/headers/local_headers/injection_recursive.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_test.cu b/tests/unit_tests/headers/local_headers/injection_test.cu index ced4b5066..23a500ff1 100644 --- a/tests/unit_tests/headers/local_headers/injection_test.cu +++ b/tests/unit_tests/headers/local_headers/injection_test.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_uses_cmath.h b/tests/unit_tests/headers/local_headers/injection_uses_cmath.h index d7b5b6a54..8ea1b31f2 100644 --- a/tests/unit_tests/headers/local_headers/injection_uses_cmath.h +++ b/tests/unit_tests/headers/local_headers/injection_uses_cmath.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef INJECTION_USES_CMATH_H #define INJECTION_USES_CMATH_H diff --git a/tests/unit_tests/headers/local_headers/local_angle.h b/tests/unit_tests/headers/local_headers/local_angle.h index 37450404b..7a8cbba1f 100644 --- a/tests/unit_tests/headers/local_headers/local_angle.h +++ b/tests/unit_tests/headers/local_headers/local_angle.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef LOCAL_ANGLE_H #define LOCAL_ANGLE_H // CHECK: #include diff --git a/tests/unit_tests/headers/local_headers/main.cu b/tests/unit_tests/headers/local_headers/main.cu index 17850a78e..2dc0e8273 100644 --- a/tests/unit_tests/headers/local_headers/main.cu +++ b/tests/unit_tests/headers/local_headers/main.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/mixed.h b/tests/unit_tests/headers/local_headers/mixed.h index 285a79150..b493046ad 100644 --- a/tests/unit_tests/headers/local_headers/mixed.h +++ b/tests/unit_tests/headers/local_headers/mixed.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef MIXED_H #define MIXED_H // CHECK: #include diff --git a/tests/unit_tests/headers/local_headers/no_local_headers.cu b/tests/unit_tests/headers/local_headers/no_local_headers.cu index e6de738eb..0701fcb53 100644 --- a/tests/unit_tests/headers/local_headers/no_local_headers.cu +++ b/tests/unit_tests/headers/local_headers/no_local_headers.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/non_cuda.h b/tests/unit_tests/headers/local_headers/non_cuda.h index 6680ef38f..78943ad00 100644 --- a/tests/unit_tests/headers/local_headers/non_cuda.h +++ b/tests/unit_tests/headers/local_headers/non_cuda.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef NON_CUDA_H #define NON_CUDA_H diff --git a/tests/unit_tests/headers/local_headers/non_cuda_header.cu b/tests/unit_tests/headers/local_headers/non_cuda_header.cu index 5ae03f26b..04943f554 100644 --- a/tests/unit_tests/headers/local_headers/non_cuda_header.cu +++ b/tests/unit_tests/headers/local_headers/non_cuda_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK: #include "non_cuda.h" diff --git a/tests/unit_tests/headers/local_headers/rec_1.h b/tests/unit_tests/headers/local_headers/rec_1.h index ddfc81b6f..8690e0ee4 100644 --- a/tests/unit_tests/headers/local_headers/rec_1.h +++ b/tests/unit_tests/headers/local_headers/rec_1.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef REC_H #define REC_H diff --git a/tests/unit_tests/headers/local_headers/recursive-header.cu b/tests/unit_tests/headers/local_headers/recursive-header.cu index 542d4d42f..89b25c650 100644 --- a/tests/unit_tests/headers/local_headers/recursive-header.cu +++ b/tests/unit_tests/headers/local_headers/recursive-header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/relative/common1.h b/tests/unit_tests/headers/local_headers/relative/common1.h index d8ade3b73..70ac3ef44 100644 --- a/tests/unit_tests/headers/local_headers/relative/common1.h +++ b/tests/unit_tests/headers/local_headers/relative/common1.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef COMMON1_H #define COMMON1_H diff --git a/tests/unit_tests/headers/local_headers/relative/sub/common2.h b/tests/unit_tests/headers/local_headers/relative/sub/common2.h index 6cc5db120..2e208c192 100644 --- a/tests/unit_tests/headers/local_headers/relative/sub/common2.h +++ b/tests/unit_tests/headers/local_headers/relative/sub/common2.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef COMMON2_H #define COMMON2_H diff --git a/tests/unit_tests/headers/local_headers/relative_path.cu b/tests/unit_tests/headers/local_headers/relative_path.cu index b5a75791c..13cde5bad 100644 --- a/tests/unit_tests/headers/local_headers/relative_path.cu +++ b/tests/unit_tests/headers/local_headers/relative_path.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/shared.h b/tests/unit_tests/headers/local_headers/shared.h index 45693dfb5..088605f34 100644 --- a/tests/unit_tests/headers/local_headers/shared.h +++ b/tests/unit_tests/headers/local_headers/shared.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef SHARED_H #define SHARED_H // CHECK: #include diff --git a/tests/unit_tests/headers/local_headers/shared_header_1.cu b/tests/unit_tests/headers/local_headers/shared_header_1.cu index 669bedbad..8f6ab12b1 100644 --- a/tests/unit_tests/headers/local_headers/shared_header_1.cu +++ b/tests/unit_tests/headers/local_headers/shared_header_1.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/shared_header_2.cu b/tests/unit_tests/headers/local_headers/shared_header_2.cu index 669bedbad..8f6ab12b1 100644 --- a/tests/unit_tests/headers/local_headers/shared_header_2.cu +++ b/tests/unit_tests/headers/local_headers/shared_header_2.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/single_header.cu b/tests/unit_tests/headers/local_headers/single_header.cu index d8efd6154..399742509 100644 --- a/tests/unit_tests/headers/local_headers/single_header.cu +++ b/tests/unit_tests/headers/local_headers/single_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/single_header.h b/tests/unit_tests/headers/local_headers/single_header.h index 97570f3ea..7291c7315 100644 --- a/tests/unit_tests/headers/local_headers/single_header.h +++ b/tests/unit_tests/headers/local_headers/single_header.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef SINGLE_HEADER_H #define SINGLE_HEADER_H // CHECK: #include diff --git a/tests/unit_tests/headers/local_headers/stress_test.cu b/tests/unit_tests/headers/local_headers/stress_test.cu index 95022efea..b65bc3e95 100644 --- a/tests/unit_tests/headers/local_headers/stress_test.cu +++ b/tests/unit_tests/headers/local_headers/stress_test.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/vector_math.h b/tests/unit_tests/headers/local_headers/vector_math.h index 12ed82cbe..ae4eb4a25 100644 --- a/tests/unit_tests/headers/local_headers/vector_math.h +++ b/tests/unit_tests/headers/local_headers/vector_math.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #ifndef VECTOR_MATH_H #define VECTOR_MATH_H From 7d267344499565abc8251be0e26de282e53e08e7 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Mon, 6 Apr 2026 22:39:15 +0530 Subject: [PATCH 04/10] [HIPIFY][fix] Fix regex, premature loop break, and duplicate warnings --- src/LocalHeader.cpp | 18 ++++++++++++------ .../local_headers/block_comment_include.h | 15 +++++++++++++++ tests/unit_tests/headers/local_headers/main.cu | 15 ++++++++++++++- .../headers/local_headers/parent_a.h | 17 +++++++++++++++++ .../headers/local_headers/parent_b.h | 17 +++++++++++++++++ .../headers/local_headers/shared_dep.h | 15 +++++++++++++++ .../headers/local_headers/subdir_a/dup_name.h | 15 +++++++++++++++ .../headers/local_headers/subdir_b/dup_name.h | 15 +++++++++++++++ 8 files changed, 120 insertions(+), 7 deletions(-) create mode 100644 tests/unit_tests/headers/local_headers/block_comment_include.h create mode 100644 tests/unit_tests/headers/local_headers/parent_a.h create mode 100644 tests/unit_tests/headers/local_headers/parent_b.h create mode 100644 tests/unit_tests/headers/local_headers/shared_dep.h create mode 100644 tests/unit_tests/headers/local_headers/subdir_a/dup_name.h create mode 100644 tests/unit_tests/headers/local_headers/subdir_b/dup_name.h diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index d5d43edcc..a45600c1e 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -65,7 +65,7 @@ static bool pathExists(const std::string &p) { namespace { static const std::regex LocalIncludeRe( - R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?$)", std::regex::ECMAScript); + R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?)", std::regex::ECMAScript); static const std::regex SystemIncludeRe( R"(^\s*#\s*include\s*<([^>\n]+)>)", std::regex::ECMAScript); @@ -110,12 +110,12 @@ namespace { while (std::getline(iss, line)) { if (std::regex_match(line, m, LocalIncludeRe)) { std::string quotedName = m[1].str(); - std::string quotedFileName = std::string(sys::path::filename(quotedName)); - if (quotedFileName == targetFileName) - break; std::string absPath; - if (resolveLocalIncludeInternal(mainSourceAbspath, quotedName, absPath)) + if (resolveLocalIncludeInternal(mainSourceAbspath, quotedName, absPath)) { + if (absPath == targetHeaderAbspath) + break; outIncludes.push_back(absPath); + } } if (std::regex_search(line, m, SystemIncludeRe)) { @@ -191,9 +191,14 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, work.push_back({h, mainSourceAbsPath}); } std::set processed; + std::set queued; size_t total = initial.size(); size_t current = 0; + for (const auto &h: initial) { + queued.insert(h); + } + while (!work.empty()) { auto [hdr, parentPath] = work.back(); work.pop_back(); @@ -242,7 +247,8 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, std::string rel = m[1].str(); std::string abs; if (resolveLocalIncludeInternal(hdr, rel, abs) && - !processed.count(abs)) { + !processed.count(abs) && !queued.count(abs)) { + queued.insert(abs); work.push_back({abs, hdr}); newHeaders.push_back(abs); } diff --git a/tests/unit_tests/headers/local_headers/block_comment_include.h b/tests/unit_tests/headers/local_headers/block_comment_include.h new file mode 100644 index 000000000..ba14e41bd --- /dev/null +++ b/tests/unit_tests/headers/local_headers/block_comment_include.h @@ -0,0 +1,15 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef BLOCK_COMMENT_INCLUDE_H +#define BLOCK_COMMENT_INCLUDE_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline void block_comment_sync() { + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/main.cu b/tests/unit_tests/headers/local_headers/main.cu index 2dc0e8273..84ce17992 100644 --- a/tests/unit_tests/headers/local_headers/main.cu +++ b/tests/unit_tests/headers/local_headers/main.cu @@ -5,8 +5,21 @@ // CHECK: #include "common.h" // CHECK: #include "single_header.h" // CHECK: #include "shared.h" +// CHECK: #include "block_comment_include.h" +// CHECK: #include "parent_a.h" +// CHECK: #include "parent_b.h" +// CHECK: #include "subdir_a/dup_name.h" +// CHECK: #include "subdir_b/dup_name.h" #include #include "common.h" #include "single_header.h" -#include "shared.h" \ No newline at end of file +#include "shared.h" + +#include "block_comment_include.h" /* block comment after include */ + +#include "parent_a.h" +#include "parent_b.h" + +#include "subdir_a/dup_name.h" +#include "subdir_b/dup_name.h" diff --git a/tests/unit_tests/headers/local_headers/parent_a.h b/tests/unit_tests/headers/local_headers/parent_a.h new file mode 100644 index 000000000..8c67626d7 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/parent_a.h @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef PARENT_A_H +#define PARENT_A_H + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "shared_dep.h" +#include +#include "shared_dep.h" + +inline void parent_a_malloc(void **p) { + // CHECK: hipMalloc(p, 32); + cudaMalloc(p, 32); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/parent_b.h b/tests/unit_tests/headers/local_headers/parent_b.h new file mode 100644 index 000000000..e5fa7b2a5 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/parent_b.h @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef PARENT_B_H +#define PARENT_B_H + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "shared_dep.h" +#include +#include "shared_dep.h" + +inline void parent_b_free(void *p) { + // CHECK: hipFree(p); + cudaFree(p); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/shared_dep.h b/tests/unit_tests/headers/local_headers/shared_dep.h new file mode 100644 index 000000000..bf7dd270b --- /dev/null +++ b/tests/unit_tests/headers/local_headers/shared_dep.h @@ -0,0 +1,15 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef SHARED_DEP_H +#define SHARED_DEP_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline void shared_dep_sync() { + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/subdir_a/dup_name.h b/tests/unit_tests/headers/local_headers/subdir_a/dup_name.h new file mode 100644 index 000000000..f4c641b2f --- /dev/null +++ b/tests/unit_tests/headers/local_headers/subdir_a/dup_name.h @@ -0,0 +1,15 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef SUBDIR_A_DUP_NAME_H +#define SUBDIR_A_DUP_NAME_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline void subdir_a_sync() { + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/subdir_b/dup_name.h b/tests/unit_tests/headers/local_headers/subdir_b/dup_name.h new file mode 100644 index 000000000..9220166d5 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/subdir_b/dup_name.h @@ -0,0 +1,15 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef SUBDIR_B_DUP_NAME_H +#define SUBDIR_B_DUP_NAME_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +inline void subdir_b_sync() { + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); +} + +#endif From fdc1148ddef8160a128f975dca05a20382fe1782 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Tue, 7 Apr 2026 16:04:16 +0530 Subject: [PATCH 05/10] [HIPIFY][fix] Fix compile_commands json issue --- src/LocalHeader.cpp | 21 +++++++++++++++++---- 1 file changed, 17 insertions(+), 4 deletions(-) diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index a45600c1e..62d8e3935 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -102,7 +102,6 @@ namespace { return false; } - std::string targetFileName = std::string(sys::path::filename(targetHeaderAbspath)); std::istringstream iss(mainSourceContent); std::string line; std::smatch m; @@ -127,6 +126,19 @@ namespace { } } +static std::string +resolveCompileContext(const std::string &parentPath, + const std::string &mainSourceAbsPath, + const clang::tooling::CompilationDatabase *compDB) { + if (!compDB) + return mainSourceAbsPath; + + if (!compDB->getCompileCommands(parentPath).empty()) + return parentPath; + + return mainSourceAbsPath; +} + bool resolveLocalInclude(const std::string &mainSourceAbsPath, const std::string &includeToken, std::string &outAbsPath) { @@ -225,9 +237,10 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, outs() << "\n" << sHipify << "Hipifying local header [" << current << "/" << total << "]: " << sys::path::filename(hdr) << "\n"; - bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, - hipify_exe, parentPath, false, - precedingIncludes); + bool ok = hipifySingleSource( + hdr, hipOut, compDB, OptionsParserPtr, hipify_exe, + resolveCompileContext(parentPath, mainSourceAbsPath, compDB), false, + precedingIncludes); if (!ok) { errs() << "\n" << sHipify << sError From df4257c164893bd0c6de69d33a02001d059b0a31 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Mon, 13 Apr 2026 11:24:09 +0530 Subject: [PATCH 06/10] [HIPIFY][fix] Fix regex and structured bindings --- src/LocalHeader.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 62d8e3935..6a0103926 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -64,8 +64,8 @@ static bool pathExists(const std::string &p) { } namespace { - static const std::regex LocalIncludeRe( - R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?)", std::regex::ECMAScript); + static const std::regex LocalIncludeRe + (R"re(^\s*#\s*include\s*"([^"\n]+)")re", std::regex::ECMAScript); static const std::regex SystemIncludeRe( R"(^\s*#\s*include\s*<([^>\n]+)>)", std::regex::ECMAScript); @@ -107,7 +107,7 @@ namespace { std::smatch m; while (std::getline(iss, line)) { - if (std::regex_match(line, m, LocalIncludeRe)) { + if (std::regex_search(line, m, LocalIncludeRe)) { std::string quotedName = m[1].str(); std::string absPath; if (resolveLocalIncludeInternal(mainSourceAbspath, quotedName, absPath)) { @@ -158,7 +158,7 @@ bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, std::istringstream iss(content); std::string line; while (std::getline(iss, line)) { - if (std::regex_match(line, m, LocalIncludeRe)) { + if (std::regex_search(line, m, LocalIncludeRe)) { std::string rel = m[1].str(); std::string abs; if (resolveLocalIncludeInternal(mainSourceAbsPath, rel, abs)){ @@ -212,7 +212,8 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, } while (!work.empty()) { - auto [hdr, parentPath] = work.back(); + std::string hdr = work.back().first; + std::string parentPath = work.back().second; work.pop_back(); if (processed.count(hdr)) { outs() << sHipify << sWarning @@ -256,7 +257,7 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, std::string line; std::vector newHeaders; while (std::getline(iss, line)) { - if (std::regex_match(line, m, LocalIncludeRe)) { + if (std::regex_search(line, m, LocalIncludeRe)) { std::string rel = m[1].str(); std::string abs; if (resolveLocalIncludeInternal(hdr, rel, abs) && From 1562bb2b646598bba962bbb955f7626e6c438cb1 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Mon, 20 Apr 2026 15:42:00 +0530 Subject: [PATCH 07/10] [HIPIFY][fix] Refactor include scanning and add depth-2 transitive context test --- src/LocalHeader.cpp | 74 +++++++++++++------ tests/lit.cfg | 2 + .../headers/local_headers/transitive_child.h | 9 +++ .../headers/local_headers/transitive_parent.h | 15 ++++ .../transitive_system_include.cu | 10 +++ 5 files changed, 87 insertions(+), 23 deletions(-) create mode 100644 tests/unit_tests/headers/local_headers/transitive_child.h create mode 100644 tests/unit_tests/headers/local_headers/transitive_parent.h create mode 100644 tests/unit_tests/headers/local_headers/transitive_system_include.cu diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 6a0103926..19c98790a 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -91,39 +91,63 @@ namespace { } return false; } - - bool collectPrecedingIncludes(const std::string &mainSourceAbspath, - const std::string &targetHeaderAbspath, - std::vector &outIncludes) { - std::string mainSourceContent; - if (!readFile(mainSourceAbspath, mainSourceContent)) { - errs() << sHipify << sError << "Cannot read source files: " - << mainSourceAbspath << "\n"; + + static bool collectIncludesBefore(const std::string &filePath, + const std::string &stopAtAbsPath, + bool collectLocal, + std::set &seen, + std::vector &outIncludes) { + std::string content; + if (!readFile(filePath, content)) return false; - } - - std::istringstream iss(mainSourceContent); + + std::istringstream iss(content); std::string line; std::smatch m; while (std::getline(iss, line)) { if (std::regex_search(line, m, LocalIncludeRe)) { - std::string quotedName = m[1].str(); - std::string absPath; - if (resolveLocalIncludeInternal(mainSourceAbspath, quotedName, absPath)) { - if (absPath == targetHeaderAbspath) + std::string abspath; + if (resolveLocalIncludeInternal(filePath, m[1].str(), abspath)) { + if (abspath == stopAtAbsPath) break; - outIncludes.push_back(absPath); + if (collectLocal && seen.insert(abspath).second) + outIncludes.push_back(abspath); } } - if (std::regex_search(line, m, SystemIncludeRe)) { - outIncludes.push_back(m[1].str()); + std::string sysInclude = m[1].str(); + if (seen.insert(sysInclude).second) + outIncludes.push_back(sysInclude); } } + return true; + } + bool collectPrecedingIncludes(const std::string &mainSourceAbspath, + const std::string &targetHeaderAbspath, + std::vector &outIncludes) { + + std::set seen; + + if (!collectIncludesBefore(mainSourceAbspath, targetHeaderAbspath, true, seen, outIncludes)) { + errs() << sHipify << sError << "Cannot read source files: " + << mainSourceAbspath << "\n"; + return false; + } return true; } + + void collectAncestorSystemIncludes( + const std::vector &ancestorChain, + std::vector &outIncludes) { + + std::set seen(outIncludes.begin(), outIncludes.end()); + + for (size_t i = 1; i < ancestorChain.size(); ++i) { + collectIncludesBefore(ancestorChain[i], ancestorChain[i - 1], false, seen, outIncludes); + } + } } static std::string @@ -198,9 +222,9 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, << ": " << sys::path::filename(initial[i]) << "\n"; } - std::vector> work; + std::vector>> work; for (const auto &h : initial) { - work.push_back({h, mainSourceAbsPath}); + work.push_back({h, std::vector{mainSourceAbsPath}}); } std::set processed; std::set queued; @@ -213,8 +237,9 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, while (!work.empty()) { std::string hdr = work.back().first; - std::string parentPath = work.back().second; + std::vector ancestorChain = work.back().second; work.pop_back(); + std::string parentPath = ancestorChain[0]; if (processed.count(hdr)) { outs() << sHipify << sWarning << "Duplicate local header reference ignored: " @@ -234,7 +259,7 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, std::string hipOut = hdr + ".hip"; std::vector precedingIncludes; collectPrecedingIncludes(parentPath, hdr, precedingIncludes); - + collectAncestorSystemIncludes(ancestorChain, precedingIncludes); outs() << "\n" << sHipify << "Hipifying local header [" << current << "/" << total << "]: " << sys::path::filename(hdr) << "\n"; @@ -263,7 +288,10 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, if (resolveLocalIncludeInternal(hdr, rel, abs) && !processed.count(abs) && !queued.count(abs)) { queued.insert(abs); - work.push_back({abs, hdr}); + std::vector childChain; + childChain.push_back(hdr); + childChain.insert(childChain.end(), ancestorChain.begin(), ancestorChain.end()); + work.push_back({abs, childChain}); newHeaders.push_back(abs); } } diff --git a/tests/lit.cfg b/tests/lit.cfg index df8b465b5..199c5047c 100644 --- a/tests/lit.cfg +++ b/tests/lit.cfg @@ -19,6 +19,8 @@ config.excludes.append('inc.h') config.excludes.append('inet.h') config.excludes.append('types.h') config.excludes.append('socket.h') +config.excludes.append('transitive_parent.h') +config.excludes.append('transitive_child.h') delimiter = "==============================================================="; print(delimiter) diff --git a/tests/unit_tests/headers/local_headers/transitive_child.h b/tests/unit_tests/headers/local_headers/transitive_child.h new file mode 100644 index 000000000..50549fa6f --- /dev/null +++ b/tests/unit_tests/headers/local_headers/transitive_child.h @@ -0,0 +1,9 @@ +#ifndef TRANSITIVE_CHILD_H +#define TRANSITIVE_CHILD_H + +inline void child_sort_alloc(int *data, int n, void **p) { + std::sort(data, data + n); + cudaMalloc(p, n * sizeof(int)); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/transitive_parent.h b/tests/unit_tests/headers/local_headers/transitive_parent.h new file mode 100644 index 000000000..34b473e62 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/transitive_parent.h @@ -0,0 +1,15 @@ +#ifndef TRANSITIVE_PARENT_H +#define TRANSITIVE_PARENT_H + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "transitive_child.h" +#include +#include "transitive_child.h" + +inline void parent_sync() { + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/transitive_system_include.cu b/tests/unit_tests/headers/local_headers/transitive_system_include.cu new file mode 100644 index 000000000..6ff75418d --- /dev/null +++ b/tests/unit_tests/headers/local_headers/transitive_system_include.cu @@ -0,0 +1,10 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "transitive_parent.h" +#include +#include +#include "transitive_parent.h" + +int main() { return 0; } From ffa0b40a1aada67915a6691a197748cccf6a5c0d Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Tue, 12 May 2026 12:49:05 +0530 Subject: [PATCH 08/10] [HIPIFY] Replace regex-based include scanning with Clang PPCallbacks and restore opt-in flags --- src/ArgParse.cpp | 68 ++-- src/ArgParse.h | 3 +- src/LocalHeader.cpp | 374 +++++++++--------- src/LocalHeader.h | 27 +- src/main.cpp | 12 +- .../headers/local_headers/angle_bracket.cu | 2 +- .../unit_tests/headers/local_headers/common.h | 2 +- .../headers/local_headers/cyclic_header.cu | 3 +- .../headers/local_headers/duplicate_header.cu | 2 +- .../headers/local_headers/injection_multi.cu | 2 +- .../local_headers/injection_no_duplicate.cu | 2 +- .../local_headers/injection_pragma_once.cu | 2 +- .../local_headers/injection_recursive.cu | 3 +- .../headers/local_headers/injection_test.cu | 2 +- .../unit_tests/headers/local_headers/main.cu | 2 +- .../unit_tests/headers/local_headers/mixed.h | 2 +- .../headers/local_headers/no_local_headers.cu | 2 +- .../headers/local_headers/non_cuda_header.cu | 2 +- .../headers/local_headers/recursive-header.cu | 3 +- .../headers/local_headers/relative_path.cu | 2 +- .../headers/local_headers/shared_header_1.cu | 2 +- .../headers/local_headers/shared_header_2.cu | 2 +- .../headers/local_headers/single_header.cu | 2 +- .../headers/local_headers/stress_test.cu | 2 +- .../transitive_system_include.cu | 2 +- 25 files changed, 267 insertions(+), 260 deletions(-) diff --git a/src/ArgParse.cpp b/src/ArgParse.cpp index f70bf08da..6c19a2c50 100644 --- a/src/ArgParse.cpp +++ b/src/ArgParse.cpp @@ -226,40 +226,46 @@ cl::opt HipDnnSupport("hipdnn", cl::init(false), cl::cat(ToolTemplateCategory)); -cl::opt SkipLocalHeaders("skip-local-headers", - cl::desc("Skip implicit hipification of local (quoted) headers included in the main source file"), - cl::init(false), - cl::cat(ToolTemplateCategory)); +cl::opt OptLocalHeaders( + "local-headers", + cl::desc("Enable hipification of quoted local headers (non-recursive)"), + cl::init(false), cl::cat(ToolTemplateCategory)); + +cl::opt OptLocalHeadersRecursive( + "local-headers-recursive", + cl::desc("Enable hipification of quoted local headers recursively"), + cl::init(false), cl::cat(ToolTemplateCategory)); cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage); -const std::vector hipifyOptions { - std::string(PrintStatsCSV.ArgStr), - std::string(PrintStats.ArgStr), - std::string(SkipExcludedPPConditionalBlocks.ArgStr), - std::string(DefaultPreprocessor.ArgStr), - std::string(HipKernelExecutionSyntax.ArgStr), - std::string(CudaKernelExecutionSyntax.ArgStr), - std::string(GeneratePerl.ArgStr), - std::string(GeneratePython.ArgStr), - std::string(TranslateToRoc.ArgStr), - std::string(TranslateToMIOpen.ArgStr), - std::string(GenerateMarkdown.ArgStr), - std::string(GenerateCSV.ArgStr), - std::string(NoBackup.ArgStr), - std::string(NoOutput.ArgStr), - std::string(Inplace.ArgStr), - std::string(Examine.ArgStr), - std::string(SaveTemps.ArgStr), - std::string(DocFormat.ArgStr), - std::string(DocRoc.ArgStr), - std::string(Experimental.ArgStr), - std::string(Versions.ArgStr), - std::string(NoUndocumented.ArgStr), - std::string(NoWarningsUndocumented.ArgStr), - std::string(HipifyAMAP.ArgStr), - std::string(HipDnnSupport.ArgStr), - std::string(SkipLocalHeaders.ArgStr), +const std::vector hipifyOptions{ + std::string(PrintStatsCSV.ArgStr), + std::string(PrintStats.ArgStr), + std::string(SkipExcludedPPConditionalBlocks.ArgStr), + std::string(DefaultPreprocessor.ArgStr), + std::string(HipKernelExecutionSyntax.ArgStr), + std::string(CudaKernelExecutionSyntax.ArgStr), + std::string(GeneratePerl.ArgStr), + std::string(GeneratePython.ArgStr), + std::string(TranslateToRoc.ArgStr), + std::string(TranslateToMIOpen.ArgStr), + std::string(GenerateMarkdown.ArgStr), + std::string(GenerateCSV.ArgStr), + std::string(NoBackup.ArgStr), + std::string(NoOutput.ArgStr), + std::string(Inplace.ArgStr), + std::string(Examine.ArgStr), + std::string(SaveTemps.ArgStr), + std::string(DocFormat.ArgStr), + std::string(DocRoc.ArgStr), + std::string(Experimental.ArgStr), + std::string(Versions.ArgStr), + std::string(NoUndocumented.ArgStr), + std::string(NoWarningsUndocumented.ArgStr), + std::string(HipifyAMAP.ArgStr), + std::string(HipDnnSupport.ArgStr), + std::string(OptLocalHeaders.ArgStr), + std::string(OptLocalHeadersRecursive.ArgStr), }; const std::vector hipifyOptionsWithTwoArgs { diff --git a/src/ArgParse.h b/src/ArgParse.h index 96405f4b9..5a32c386f 100644 --- a/src/ArgParse.h +++ b/src/ArgParse.h @@ -70,4 +70,5 @@ extern cl::opt NoUndocumented; extern cl::opt NoWarningsUndocumented; extern cl::opt HipifyAMAP; extern cl::opt HipDnnSupport; -extern cl::opt SkipLocalHeaders; +extern cl::opt OptLocalHeaders; +extern cl::opt OptLocalHeadersRecursive; diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 19c98790a..734ea53b0 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -23,15 +23,16 @@ THE SOFTWARE. #include "LocalHeader.h" #include "LLVMCompat.h" -#include -#include #include #include -#include -#include "llvm/Support/FileSystem.h" +#include "clang/Frontend/CompilerInstance.h" +#include "clang/Frontend/FrontendActions.h" +#include "clang/Lex/PPCallbacks.h" +#include "clang/Lex/Preprocessor.h" +#include "clang/Tooling/CompilationDatabase.h" +#include "clang/Tooling/Tooling.h" #include "llvm/Support/Path.h" -#include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/raw_ostream.h" using namespace clang; @@ -39,176 +40,187 @@ using namespace clang::tooling; using namespace llvm; using namespace std; -static std::string normalizeSmallStringPath(SmallString<256> &p) { - llvm::sys::path::remove_dots(p, true); - - SmallString<256> realBuf; - std::error_code ec = llvm::sys::fs::real_path(p, realBuf); - if (!ec) { - return std::string(realBuf.str()); - } - - return std::string(p.str()); -} - -static bool pathExists(const std::string &p) { - SmallString<256> in(p.begin(), p.end()); - - SmallString<256> realBuf; - std::error_code ec = llvm::sys::fs::real_path(in, realBuf); - if (!ec) return true; - - SmallString<256> norm = in; - llvm::sys::path::remove_dots(norm, true); - return llvm::sys::fs::exists(norm); -} - namespace { - static const std::regex LocalIncludeRe - (R"re(^\s*#\s*include\s*"([^"\n]+)")re", std::regex::ECMAScript); - - static const std::regex SystemIncludeRe( - R"(^\s*#\s*include\s*<([^>\n]+)>)", std::regex::ECMAScript); - - bool readFile(const std::string &path, std::string &out) { - auto MBOrErr = llvm::MemoryBuffer::getFile(path); - if (!MBOrErr) return false; - out = MBOrErr->get()->getBuffer().str(); - return true; - } - bool resolveLocalIncludeInternal(const std::string &mainSourceAbsPath, - const std::string &includeTok, - std::string &outAbs) { - SmallString<256> base(mainSourceAbsPath); - sys::path::remove_filename(base); - SmallString<256> candidate(base); - sys::path::append(candidate, includeTok); - sys::path::remove_dots(candidate, true); - if (pathExists(std::string(candidate.str()))) { - outAbs = normalizeSmallStringPath(candidate); - return true; +class IncludeCollectorCallbacks : public clang::PPCallbacks { + const clang::SourceManager &SM; + std::vector &Entries; + +public: + IncludeCollectorCallbacks(const clang::SourceManager &SM, + std::vector &entries) + : SM(SM), Entries(entries) {} + + void InclusionDirective(clang::SourceLocation hash_loc, + const clang::Token &include_token, + StringRef file_name, bool is_angled, + clang::CharSourceRange filename_range, +#if LLVM_VERSION_MAJOR < 15 + const clang::FileEntry *file, +#elif LLVM_VERSION_MAJOR == 15 + Optional file, +#else + clang::OptionalFileEntryRef file, +#endif + StringRef search_path, StringRef relative_path, +#if LLVM_VERSION_MAJOR < 19 + const clang::Module *SuggestedModule +#else + const clang::Module *SuggestedModule, + bool ModuleImported +#endif +#if LLVM_VERSION_MAJOR > 6 + , + clang::SrcMgr::CharacteristicKind FileType +#endif + ) override { + if (!SM.isWrittenInMainFile(hash_loc)) + return; + + IncludeEntry entry; + entry.fileName = file_name.str(); + entry.isAngled = is_angled; + + if (file) { +#if LLVM_VERSION_MAJOR < 15 + entry.resolvedPath = file->tryGetRealPathName().str(); + if (entry.resolvedPath.empty()) + entry.resolvedPath = file->getName().str(); +#else + entry.resolvedPath = file->getFileEntry().tryGetRealPathName().str(); + if (entry.resolvedPath.empty()) + entry.resolvedPath = file->getName().str(); +#endif } - return false; + + Entries.push_back(std::move(entry)); } - - static bool collectIncludesBefore(const std::string &filePath, - const std::string &stopAtAbsPath, - bool collectLocal, - std::set &seen, - std::vector &outIncludes) { - std::string content; - if (!readFile(filePath, content)) - return false; - - std::istringstream iss(content); - std::string line; - std::smatch m; - - while (std::getline(iss, line)) { - if (std::regex_search(line, m, LocalIncludeRe)) { - std::string abspath; - if (resolveLocalIncludeInternal(filePath, m[1].str(), abspath)) { - if (abspath == stopAtAbsPath) - break; - if (collectLocal && seen.insert(abspath).second) - outIncludes.push_back(abspath); - } - } - if (std::regex_search(line, m, SystemIncludeRe)) { - std::string sysInclude = m[1].str(); - if (seen.insert(sysInclude).second) - outIncludes.push_back(sysInclude); - } - } - return true; +}; + +class IncludeCollectorAction : public clang::PreprocessorFrontendAction { + std::vector &Entries; + +public: + explicit IncludeCollectorAction(std::vector &entries) + : Entries(entries) {} + + void ExecuteAction() override { + clang::CompilerInstance &CI = getCompilerInstance(); + clang::Preprocessor &PP = CI.getPreprocessor(); + PP.addPPCallbacks(std::make_unique( + CI.getSourceManager(), Entries)); + + PP.EnterMainSourceFile(); + clang::Token Tok; + do { + PP.Lex(Tok); + } while (Tok.isNot(clang::tok::eof)); } +}; - bool collectPrecedingIncludes(const std::string &mainSourceAbspath, - const std::string &targetHeaderAbspath, - std::vector &outIncludes) { +class IncludeCollectorActionFactory : public FrontendActionFactory { + std::vector &Entries; - std::set seen; +public: + explicit IncludeCollectorActionFactory(std::vector &entries) + : Entries(entries) {} - if (!collectIncludesBefore(mainSourceAbspath, targetHeaderAbspath, true, seen, outIncludes)) { - errs() << sHipify << sError << "Cannot read source files: " - << mainSourceAbspath << "\n"; - return false; - } - return true; +#if LLVM_VERSION_MAJOR >= 10 + std::unique_ptr create() override { + return std::make_unique(Entries); } - - void collectAncestorSystemIncludes( - const std::vector &ancestorChain, - std::vector &outIncludes) { - - std::set seen(outIncludes.begin(), outIncludes.end()); - - for (size_t i = 1; i < ancestorChain.size(); ++i) { - collectIncludesBefore(ancestorChain[i], ancestorChain[i - 1], false, seen, outIncludes); - } +#else + clang::FrontendAction *create() override { + return new IncludeCollectorAction(Entries); + } +#endif +}; + +} // namespace + +bool collectIncludeTree(const std::string &srcPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe, + const std::string &mainContextPath, + std::vector &outEntries) { + outEntries.clear(); + + const ct::CompilationDatabase &baseDB = + compDB ? *compDB : OptionsParserPtr->getCompilations(); + + // If srcPath has no entry in the compilation database, fall back to a + // FixedCompilationDatabase rooted at the mainContextPath's directory so that + // the tool doesn't skip the file. + std::vector cmds = baseDB.getCompileCommands(srcPath); + std::unique_ptr fallbackDB; + if (cmds.empty()) { + std::string dir = sys::path::parent_path(mainContextPath).str(); + fallbackDB = std::make_unique( + dir, std::vector()); } -} -static std::string -resolveCompileContext(const std::string &parentPath, - const std::string &mainSourceAbsPath, - const clang::tooling::CompilationDatabase *compDB) { - if (!compDB) - return mainSourceAbsPath; + ct::RefactoringTool Tool(fallbackDB ? *fallbackDB : baseDB, {srcPath}); - if (!compDB->getCompileCommands(parentPath).empty()) - return parentPath; + if (!appendArgumentsAdjusters(Tool, mainContextPath, hipify_exe)) { + return false; + } - return mainSourceAbsPath; -} + // Strip the implicit CUDA header that appendArgumentsAdjusters adds — + // not needed for include scanning and would pollute the results. + Tool.appendArgumentsAdjuster( + [](const ct::CommandLineArguments &Args, StringRef) { + ct::CommandLineArguments filtered; + for (size_t i = 0; i < Args.size(); ++i) { + if (Args[i] == "-include" && i + 1 < Args.size() && + Args[i + 1] == "cuda_runtime.h") { + ++i; + continue; + } + filtered.push_back(Args[i]); + } + return filtered; + }); -bool resolveLocalInclude(const std::string &mainSourceAbsPath, - const std::string &includeToken, - std::string &outAbsPath) { - return resolveLocalIncludeInternal(mainSourceAbsPath, includeToken, outAbsPath); + IncludeCollectorActionFactory factory(outEntries); + Tool.run(&factory); + return true; } bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe, std::vector &outHeaders) { - std::string content; - if (!readFile(mainSourceAbsPath, content)) { - errs() << "\n" << sHipify << sError << "Cannot read source file: " << mainSourceAbsPath << "\n"; + std::vector entries; + if (!collectIncludeTree(mainSourceAbsPath, compDB, OptionsParserPtr, + hipify_exe, mainSourceAbsPath, entries)) { + errs() << "\n" + << sHipify << sError + << "Failed to collect includes from: " << mainSourceAbsPath << "\n"; return false; } std::set uniq; - std::smatch m; - std::istringstream iss(content); - std::string line; - while (std::getline(iss, line)) { - if (std::regex_search(line, m, LocalIncludeRe)) { - std::string rel = m[1].str(); - std::string abs; - if (resolveLocalIncludeInternal(mainSourceAbsPath, rel, abs)){ - uniq.insert(abs); - } else { - errs() << sHipify << sWarning - << "Missing local header referenced: \"" << rel - << "\" in " << mainSourceAbsPath << "\n"; - } - } + for (const auto &e : entries) { + if (!e.isAngled && !e.resolvedPath.empty()) + uniq.insert(e.resolvedPath); } outHeaders.assign(uniq.begin(), uniq.end()); return true; } bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, - const ct::CompilationDatabase *compDB, - ct::CommonOptionsParser *OptionsParserPtr, - const char *hipify_exe, - bool recursive) { + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe, bool recursive) { std::vector initial; - if (!collectLocalQuotedIncludes(mainSourceAbsPath, initial)) { + if (!collectLocalQuotedIncludes(mainSourceAbsPath, compDB, OptionsParserPtr, + hipify_exe, initial)) { return false; } - + if (initial.empty()) { outs() << "\n" << sHipify << "No local headers detected in " << sys::path::filename(mainSourceAbsPath) << "\n"; @@ -222,51 +234,24 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, << ": " << sys::path::filename(initial[i]) << "\n"; } - std::vector>> work; - for (const auto &h : initial) { - work.push_back({h, std::vector{mainSourceAbsPath}}); - } + std::vector work(initial.begin(), initial.end()); std::set processed; - std::set queued; + std::set queued(initial.begin(), initial.end()); size_t total = initial.size(); size_t current = 0; - for (const auto &h: initial) { - queued.insert(h); - } - while (!work.empty()) { - std::string hdr = work.back().first; - std::vector ancestorChain = work.back().second; + std::string hdr = work.back(); work.pop_back(); - std::string parentPath = ancestorChain[0]; if (processed.count(hdr)) { - outs() << sHipify << sWarning - << "Duplicate local header reference ignored: " - << sys::path::filename(hdr) << "\n"; continue; } processed.insert(hdr); ++current; - std::string original; - if (!readFile(hdr, original)) { - errs() << "\n" << sHipify << sError - << "Cannot read header: " << sys::path::filename(hdr) << "\n"; - continue; - } - std::string hipOut = hdr + ".hip"; - std::vector precedingIncludes; - collectPrecedingIncludes(parentPath, hdr, precedingIncludes); - collectAncestorSystemIncludes(ancestorChain, precedingIncludes); - outs() << "\n" << sHipify << "Hipifying local header [" << current - << "/" << total << "]: " << sys::path::filename(hdr) << "\n"; - - bool ok = hipifySingleSource( - hdr, hipOut, compDB, OptionsParserPtr, hipify_exe, - resolveCompileContext(parentPath, mainSourceAbsPath, compDB), false, - precedingIncludes); + bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, + hipify_exe, mainSourceAbsPath, false); if (!ok) { errs() << "\n" << sHipify << sError @@ -277,30 +262,25 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, outs() << sHipify << "Successfully hipified header file" << "\n"; if (recursive) { - std::smatch m; - std::istringstream iss(original); - std::string line; - std::vector newHeaders; - while (std::getline(iss, line)) { - if (std::regex_search(line, m, LocalIncludeRe)) { - std::string rel = m[1].str(); - std::string abs; - if (resolveLocalIncludeInternal(hdr, rel, abs) && - !processed.count(abs) && !queued.count(abs)) { - queued.insert(abs); - std::vector childChain; - childChain.push_back(hdr); - childChain.insert(childChain.end(), ancestorChain.begin(), ancestorChain.end()); - work.push_back({abs, childChain}); - newHeaders.push_back(abs); + std::vector childEntries; + if (collectIncludeTree(hdr, compDB, OptionsParserPtr, hipify_exe, + mainSourceAbsPath, childEntries)) { + std::vector newHeaders; + for (const auto &e : childEntries) { + if (!e.isAngled && !e.resolvedPath.empty() && + !processed.count(e.resolvedPath) && + !queued.count(e.resolvedPath)) { + newHeaders.push_back(e.resolvedPath); + work.push_back(e.resolvedPath); + queued.insert(e.resolvedPath); } } - } - if (!newHeaders.empty()) { - total += newHeaders.size(); - outs() << sHipify << " Recursive: found " << newHeaders.size() - << " additional local header(s) in " - << sys::path::filename(hdr) << "\n"; + if (!newHeaders.empty()) { + total += newHeaders.size(); + outs() << sHipify << " Recursive: found " << newHeaders.size() + << " additional local header(s) in " + << sys::path::filename(hdr) << "\n"; + } } } } diff --git a/src/LocalHeader.h b/src/LocalHeader.h index 097c1df61..b3668bf6d 100644 --- a/src/LocalHeader.h +++ b/src/LocalHeader.h @@ -22,12 +22,25 @@ THE SOFTWARE. #pragma once +#include "clang/Frontend/FrontendActions.h" +#include "clang/Lex/PPCallbacks.h" +#include "clang/Tooling/CommonOptionsParser.h" +#include "clang/Tooling/Refactoring.h" #include #include -#include "clang/Tooling/CommonOptionsParser.h" namespace ct = clang::tooling; +extern bool appendArgumentsAdjusters(ct::RefactoringTool &Tool, + const std::string &sSourceAbsPath, + const char *hipify_exe); + +struct IncludeEntry { + std::string fileName; + std::string resolvedPath; + bool isAngled; +}; + extern bool hipifySingleSource(const std::string &srcPath, const std::string &dstPath, const ct::CompilationDatabase *compDB, @@ -43,9 +56,15 @@ bool hipifyLocalHeaders(const std::string &srcPath, const char *hipify_exe, bool recursive = false); -bool resolveLocalInclude(const std::string &mainSourceAbsPath, - const std::string &includeToken, - std::string &outAbsPath); +bool collectIncludeTree(const std::string &srcPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe, + const std::string &mainContextPath, + std::vector &outEntries); bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe, std::vector &outHeaders); diff --git a/src/main.cpp b/src/main.cpp index 6b8a3c219..572872284 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -527,12 +527,10 @@ int main(int argc, const char **argv) { } // Initialise the statistics counters for this file. Statistics::setActive(src); - if (!SkipLocalHeaders) { - if (!hipifyLocalHeaders(sSourceAbsPath, - compilationDatabase.get(), - &OptionsParser, - argv[0], - true)) { + if (OptLocalHeaders || OptLocalHeadersRecursive) { + if (!hipifyLocalHeaders(sSourceAbsPath, compilationDatabase.get(), + &OptionsParser, argv[0], + OptLocalHeadersRecursive)) { llvm::errs() << "\n" << sHipify << sError << "Local header hipification failed for: " << sys::path::filename(sSourceAbsPath) << "\n"; @@ -540,7 +538,7 @@ int main(int argc, const char **argv) { Result = 1; } } - + std::string outputPath = NoOutput ? "" : dst; if (!hipifySingleSource(src, outputPath, compilationDatabase.get(), diff --git a/tests/unit_tests/headers/local_headers/angle_bracket.cu b/tests/unit_tests/headers/local_headers/angle_bracket.cu index ef8313513..e9628f061 100644 --- a/tests/unit_tests/headers/local_headers/angle_bracket.cu +++ b/tests/unit_tests/headers/local_headers/angle_bracket.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/common.h b/tests/unit_tests/headers/local_headers/common.h index 363b0dd54..cb5844475 100644 --- a/tests/unit_tests/headers/local_headers/common.h +++ b/tests/unit_tests/headers/local_headers/common.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args #ifndef COMMON_H #define COMMON_H diff --git a/tests/unit_tests/headers/local_headers/cyclic_header.cu b/tests/unit_tests/headers/local_headers/cyclic_header.cu index 9ab52d89d..9a0c3afff 100644 --- a/tests/unit_tests/headers/local_headers/cyclic_header.cu +++ b/tests/unit_tests/headers/local_headers/cyclic_header.cu @@ -1,4 +1,5 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive +// %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/duplicate_header.cu b/tests/unit_tests/headers/local_headers/duplicate_header.cu index d64a9c084..d07f7f74d 100644 --- a/tests/unit_tests/headers/local_headers/duplicate_header.cu +++ b/tests/unit_tests/headers/local_headers/duplicate_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK: #include "duplicate.h" diff --git a/tests/unit_tests/headers/local_headers/injection_multi.cu b/tests/unit_tests/headers/local_headers/injection_multi.cu index f4dd3c7aa..20341f306 100644 --- a/tests/unit_tests/headers/local_headers/injection_multi.cu +++ b/tests/unit_tests/headers/local_headers/injection_multi.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu index cebb191c1..195ecddf6 100644 --- a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu +++ b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu index 67e5e370b..f9643f1e4 100644 --- a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu +++ b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_recursive.cu b/tests/unit_tests/headers/local_headers/injection_recursive.cu index 0d705509c..e9d0c9a32 100644 --- a/tests/unit_tests/headers/local_headers/injection_recursive.cu +++ b/tests/unit_tests/headers/local_headers/injection_recursive.cu @@ -1,4 +1,5 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive +// %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_test.cu b/tests/unit_tests/headers/local_headers/injection_test.cu index 23a500ff1..ced4b5066 100644 --- a/tests/unit_tests/headers/local_headers/injection_test.cu +++ b/tests/unit_tests/headers/local_headers/injection_test.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/main.cu b/tests/unit_tests/headers/local_headers/main.cu index 84ce17992..8108d2020 100644 --- a/tests/unit_tests/headers/local_headers/main.cu +++ b/tests/unit_tests/headers/local_headers/main.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/mixed.h b/tests/unit_tests/headers/local_headers/mixed.h index b493046ad..285a79150 100644 --- a/tests/unit_tests/headers/local_headers/mixed.h +++ b/tests/unit_tests/headers/local_headers/mixed.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args #ifndef MIXED_H #define MIXED_H // CHECK: #include diff --git a/tests/unit_tests/headers/local_headers/no_local_headers.cu b/tests/unit_tests/headers/local_headers/no_local_headers.cu index 0701fcb53..e6de738eb 100644 --- a/tests/unit_tests/headers/local_headers/no_local_headers.cu +++ b/tests/unit_tests/headers/local_headers/no_local_headers.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/non_cuda_header.cu b/tests/unit_tests/headers/local_headers/non_cuda_header.cu index 04943f554..5ae03f26b 100644 --- a/tests/unit_tests/headers/local_headers/non_cuda_header.cu +++ b/tests/unit_tests/headers/local_headers/non_cuda_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK: #include "non_cuda.h" diff --git a/tests/unit_tests/headers/local_headers/recursive-header.cu b/tests/unit_tests/headers/local_headers/recursive-header.cu index 89b25c650..f712105ad 100644 --- a/tests/unit_tests/headers/local_headers/recursive-header.cu +++ b/tests/unit_tests/headers/local_headers/recursive-header.cu @@ -1,4 +1,5 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive +// %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/relative_path.cu b/tests/unit_tests/headers/local_headers/relative_path.cu index 13cde5bad..b5a75791c 100644 --- a/tests/unit_tests/headers/local_headers/relative_path.cu +++ b/tests/unit_tests/headers/local_headers/relative_path.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/shared_header_1.cu b/tests/unit_tests/headers/local_headers/shared_header_1.cu index 8f6ab12b1..669bedbad 100644 --- a/tests/unit_tests/headers/local_headers/shared_header_1.cu +++ b/tests/unit_tests/headers/local_headers/shared_header_1.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/shared_header_2.cu b/tests/unit_tests/headers/local_headers/shared_header_2.cu index 8f6ab12b1..669bedbad 100644 --- a/tests/unit_tests/headers/local_headers/shared_header_2.cu +++ b/tests/unit_tests/headers/local_headers/shared_header_2.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/single_header.cu b/tests/unit_tests/headers/local_headers/single_header.cu index 399742509..d8efd6154 100644 --- a/tests/unit_tests/headers/local_headers/single_header.cu +++ b/tests/unit_tests/headers/local_headers/single_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/stress_test.cu b/tests/unit_tests/headers/local_headers/stress_test.cu index b65bc3e95..95022efea 100644 --- a/tests/unit_tests/headers/local_headers/stress_test.cu +++ b/tests/unit_tests/headers/local_headers/stress_test.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/transitive_system_include.cu b/tests/unit_tests/headers/local_headers/transitive_system_include.cu index 6ff75418d..3967917b8 100644 --- a/tests/unit_tests/headers/local_headers/transitive_system_include.cu +++ b/tests/unit_tests/headers/local_headers/transitive_system_include.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include From 8892c93641f60fb4ccc6494acf1266e4d8c7aca2 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Tue, 26 May 2026 10:47:19 +0530 Subject: [PATCH 09/10] [HIPIFY][fix] Apply review fixes: forward decl, , lit RUN args, depth-2 test --- src/ArgParse.cpp | 72 +++++++++---------- src/LocalHeader.cpp | 21 ++---- src/LocalHeader.h | 4 -- tests/lit.cfg | 3 + .../headers/local_headers/angle_bracket.cu | 2 +- .../unit_tests/headers/local_headers/common.h | 2 +- .../headers/local_headers/cyclic_header.cu | 3 +- .../headers/local_headers/diamond_dep.cu | 13 ++++ .../headers/local_headers/diamond_left.h | 11 +++ .../headers/local_headers/diamond_right.h | 11 +++ .../headers/local_headers/diamond_shared.h | 9 +++ .../headers/local_headers/duplicate_header.cu | 2 +- .../headers/local_headers/injection_multi.cu | 2 +- .../local_headers/injection_no_duplicate.cu | 2 +- .../local_headers/injection_pragma_once.cu | 2 +- .../local_headers/injection_recursive.cu | 3 +- .../headers/local_headers/injection_test.cu | 2 +- .../unit_tests/headers/local_headers/main.cu | 2 +- .../unit_tests/headers/local_headers/mixed.h | 2 +- .../headers/local_headers/no_local_headers.cu | 2 +- .../headers/local_headers/non_cuda_header.cu | 2 +- .../headers/local_headers/recursive-header.cu | 3 +- .../headers/local_headers/relative_path.cu | 2 +- .../headers/local_headers/shared_header_1.cu | 2 +- .../headers/local_headers/shared_header_2.cu | 2 +- .../headers/local_headers/single_header.cu | 2 +- .../headers/local_headers/stress_test.cu | 2 +- .../headers/local_headers/transitive_child.h | 3 + .../transitive_system_include.cu | 2 +- 29 files changed, 111 insertions(+), 79 deletions(-) create mode 100644 tests/unit_tests/headers/local_headers/diamond_dep.cu create mode 100644 tests/unit_tests/headers/local_headers/diamond_left.h create mode 100644 tests/unit_tests/headers/local_headers/diamond_right.h create mode 100644 tests/unit_tests/headers/local_headers/diamond_shared.h diff --git a/src/ArgParse.cpp b/src/ArgParse.cpp index 6c19a2c50..de10b7f82 100644 --- a/src/ArgParse.cpp +++ b/src/ArgParse.cpp @@ -226,46 +226,46 @@ cl::opt HipDnnSupport("hipdnn", cl::init(false), cl::cat(ToolTemplateCategory)); -cl::opt OptLocalHeaders( - "local-headers", - cl::desc("Enable hipification of quoted local headers (non-recursive)"), - cl::init(false), cl::cat(ToolTemplateCategory)); +cl::opt OptLocalHeaders("local-headers", + cl::desc("Enable hipification of quoted local headers (non-recursive)"), + cl::init(false), + cl::cat(ToolTemplateCategory)); -cl::opt OptLocalHeadersRecursive( - "local-headers-recursive", - cl::desc("Enable hipification of quoted local headers recursively"), - cl::init(false), cl::cat(ToolTemplateCategory)); +cl::opt OptLocalHeadersRecursive("local-headers-recursive", + cl::desc("Enable hipification of quoted local headers recursively"), + cl::init(false), + cl::cat(ToolTemplateCategory)); cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage); -const std::vector hipifyOptions{ - std::string(PrintStatsCSV.ArgStr), - std::string(PrintStats.ArgStr), - std::string(SkipExcludedPPConditionalBlocks.ArgStr), - std::string(DefaultPreprocessor.ArgStr), - std::string(HipKernelExecutionSyntax.ArgStr), - std::string(CudaKernelExecutionSyntax.ArgStr), - std::string(GeneratePerl.ArgStr), - std::string(GeneratePython.ArgStr), - std::string(TranslateToRoc.ArgStr), - std::string(TranslateToMIOpen.ArgStr), - std::string(GenerateMarkdown.ArgStr), - std::string(GenerateCSV.ArgStr), - std::string(NoBackup.ArgStr), - std::string(NoOutput.ArgStr), - std::string(Inplace.ArgStr), - std::string(Examine.ArgStr), - std::string(SaveTemps.ArgStr), - std::string(DocFormat.ArgStr), - std::string(DocRoc.ArgStr), - std::string(Experimental.ArgStr), - std::string(Versions.ArgStr), - std::string(NoUndocumented.ArgStr), - std::string(NoWarningsUndocumented.ArgStr), - std::string(HipifyAMAP.ArgStr), - std::string(HipDnnSupport.ArgStr), - std::string(OptLocalHeaders.ArgStr), - std::string(OptLocalHeadersRecursive.ArgStr), +const std::vector hipifyOptions { + std::string(PrintStatsCSV.ArgStr), + std::string(PrintStats.ArgStr), + std::string(SkipExcludedPPConditionalBlocks.ArgStr), + std::string(DefaultPreprocessor.ArgStr), + std::string(HipKernelExecutionSyntax.ArgStr), + std::string(CudaKernelExecutionSyntax.ArgStr), + std::string(GeneratePerl.ArgStr), + std::string(GeneratePython.ArgStr), + std::string(TranslateToRoc.ArgStr), + std::string(TranslateToMIOpen.ArgStr), + std::string(GenerateMarkdown.ArgStr), + std::string(GenerateCSV.ArgStr), + std::string(NoBackup.ArgStr), + std::string(NoOutput.ArgStr), + std::string(Inplace.ArgStr), + std::string(Examine.ArgStr), + std::string(SaveTemps.ArgStr), + std::string(DocFormat.ArgStr), + std::string(DocRoc.ArgStr), + std::string(Experimental.ArgStr), + std::string(Versions.ArgStr), + std::string(NoUndocumented.ArgStr), + std::string(NoWarningsUndocumented.ArgStr), + std::string(HipifyAMAP.ArgStr), + std::string(HipDnnSupport.ArgStr), + std::string(OptLocalHeaders.ArgStr), + std::string(OptLocalHeadersRecursive.ArgStr), }; const std::vector hipifyOptionsWithTwoArgs { diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 734ea53b0..6a36705e4 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -23,6 +23,7 @@ THE SOFTWARE. #include "LocalHeader.h" #include "LLVMCompat.h" +#include #include #include @@ -138,6 +139,10 @@ class IncludeCollectorActionFactory : public FrontendActionFactory { } // namespace +bool appendArgumentsAdjusters(ct::RefactoringTool &Tool, + const std::string &sSourceAbsPath, + const char *hipify_exe); + bool collectIncludeTree(const std::string &srcPath, const ct::CompilationDatabase *compDB, ct::CommonOptionsParser *OptionsParserPtr, @@ -166,22 +171,6 @@ bool collectIncludeTree(const std::string &srcPath, return false; } - // Strip the implicit CUDA header that appendArgumentsAdjusters adds — - // not needed for include scanning and would pollute the results. - Tool.appendArgumentsAdjuster( - [](const ct::CommandLineArguments &Args, StringRef) { - ct::CommandLineArguments filtered; - for (size_t i = 0; i < Args.size(); ++i) { - if (Args[i] == "-include" && i + 1 < Args.size() && - Args[i + 1] == "cuda_runtime.h") { - ++i; - continue; - } - filtered.push_back(Args[i]); - } - return filtered; - }); - IncludeCollectorActionFactory factory(outEntries); Tool.run(&factory); return true; diff --git a/src/LocalHeader.h b/src/LocalHeader.h index b3668bf6d..caaba0970 100644 --- a/src/LocalHeader.h +++ b/src/LocalHeader.h @@ -31,10 +31,6 @@ THE SOFTWARE. namespace ct = clang::tooling; -extern bool appendArgumentsAdjusters(ct::RefactoringTool &Tool, - const std::string &sSourceAbsPath, - const char *hipify_exe); - struct IncludeEntry { std::string fileName; std::string resolvedPath; diff --git a/tests/lit.cfg b/tests/lit.cfg index 199c5047c..332a64637 100644 --- a/tests/lit.cfg +++ b/tests/lit.cfg @@ -21,6 +21,9 @@ config.excludes.append('types.h') config.excludes.append('socket.h') config.excludes.append('transitive_parent.h') config.excludes.append('transitive_child.h') +config.excludes.append('diamond_left.h') +config.excludes.append('diamond_right.h') +config.excludes.append('diamond_shared.h') delimiter = "==============================================================="; print(delimiter) diff --git a/tests/unit_tests/headers/local_headers/angle_bracket.cu b/tests/unit_tests/headers/local_headers/angle_bracket.cu index e9628f061..040d9367d 100644 --- a/tests/unit_tests/headers/local_headers/angle_bracket.cu +++ b/tests/unit_tests/headers/local_headers/angle_bracket.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/common.h b/tests/unit_tests/headers/local_headers/common.h index cb5844475..3903ec5c7 100644 --- a/tests/unit_tests/headers/local_headers/common.h +++ b/tests/unit_tests/headers/local_headers/common.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args #ifndef COMMON_H #define COMMON_H diff --git a/tests/unit_tests/headers/local_headers/cyclic_header.cu b/tests/unit_tests/headers/local_headers/cyclic_header.cu index 9a0c3afff..059d1651a 100644 --- a/tests/unit_tests/headers/local_headers/cyclic_header.cu +++ b/tests/unit_tests/headers/local_headers/cyclic_header.cu @@ -1,5 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive -// %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers-recursive %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/diamond_dep.cu b/tests/unit_tests/headers/local_headers/diamond_dep.cu new file mode 100644 index 000000000..83e7ad02e --- /dev/null +++ b/tests/unit_tests/headers/local_headers/diamond_dep.cu @@ -0,0 +1,13 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers-recursive %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "diamond_left.h" +// CHECK: #include "diamond_right.h" +#include +#include "diamond_left.h" +#include "diamond_right.h" + +int main() { + return 0; +} diff --git a/tests/unit_tests/headers/local_headers/diamond_left.h b/tests/unit_tests/headers/local_headers/diamond_left.h new file mode 100644 index 000000000..919e6c8ba --- /dev/null +++ b/tests/unit_tests/headers/local_headers/diamond_left.h @@ -0,0 +1,11 @@ +#ifndef DIAMOND_LEFT_H +#define DIAMOND_LEFT_H + +#include "diamond_shared.h" + +inline void left_sync() { + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/diamond_right.h b/tests/unit_tests/headers/local_headers/diamond_right.h new file mode 100644 index 000000000..0762747d5 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/diamond_right.h @@ -0,0 +1,11 @@ +#ifndef DIAMOND_RIGHT_H +#define DIAMOND_RIGHT_H + +#include "diamond_shared.h" + +inline void right_free(void *p) { + // CHECK: hipFree(p); + cudaFree(p); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/diamond_shared.h b/tests/unit_tests/headers/local_headers/diamond_shared.h new file mode 100644 index 000000000..ffb326f9e --- /dev/null +++ b/tests/unit_tests/headers/local_headers/diamond_shared.h @@ -0,0 +1,9 @@ +#ifndef DIAMOND_SHARED_H +#define DIAMOND_SHARED_H + +inline void shared_alloc(void **p, size_t size) { + // CHECK: hipMalloc(p, size); + cudaMalloc(p, size); +} + +#endif diff --git a/tests/unit_tests/headers/local_headers/duplicate_header.cu b/tests/unit_tests/headers/local_headers/duplicate_header.cu index d07f7f74d..d87270043 100644 --- a/tests/unit_tests/headers/local_headers/duplicate_header.cu +++ b/tests/unit_tests/headers/local_headers/duplicate_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK: #include "duplicate.h" diff --git a/tests/unit_tests/headers/local_headers/injection_multi.cu b/tests/unit_tests/headers/local_headers/injection_multi.cu index 20341f306..0283d1047 100644 --- a/tests/unit_tests/headers/local_headers/injection_multi.cu +++ b/tests/unit_tests/headers/local_headers/injection_multi.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu index 195ecddf6..05747193d 100644 --- a/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu +++ b/tests/unit_tests/headers/local_headers/injection_no_duplicate.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu index f9643f1e4..4887bd592 100644 --- a/tests/unit_tests/headers/local_headers/injection_pragma_once.cu +++ b/tests/unit_tests/headers/local_headers/injection_pragma_once.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_recursive.cu b/tests/unit_tests/headers/local_headers/injection_recursive.cu index e9d0c9a32..e943933b5 100644 --- a/tests/unit_tests/headers/local_headers/injection_recursive.cu +++ b/tests/unit_tests/headers/local_headers/injection_recursive.cu @@ -1,5 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive -// %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers-recursive %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/injection_test.cu b/tests/unit_tests/headers/local_headers/injection_test.cu index ced4b5066..8bcc93801 100644 --- a/tests/unit_tests/headers/local_headers/injection_test.cu +++ b/tests/unit_tests/headers/local_headers/injection_test.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/main.cu b/tests/unit_tests/headers/local_headers/main.cu index 8108d2020..79b1bb617 100644 --- a/tests/unit_tests/headers/local_headers/main.cu +++ b/tests/unit_tests/headers/local_headers/main.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/mixed.h b/tests/unit_tests/headers/local_headers/mixed.h index 285a79150..7e6413cff 100644 --- a/tests/unit_tests/headers/local_headers/mixed.h +++ b/tests/unit_tests/headers/local_headers/mixed.h @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args #ifndef MIXED_H #define MIXED_H // CHECK: #include diff --git a/tests/unit_tests/headers/local_headers/no_local_headers.cu b/tests/unit_tests/headers/local_headers/no_local_headers.cu index e6de738eb..24b2d11cc 100644 --- a/tests/unit_tests/headers/local_headers/no_local_headers.cu +++ b/tests/unit_tests/headers/local_headers/no_local_headers.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/non_cuda_header.cu b/tests/unit_tests/headers/local_headers/non_cuda_header.cu index 5ae03f26b..9d1dbce17 100644 --- a/tests/unit_tests/headers/local_headers/non_cuda_header.cu +++ b/tests/unit_tests/headers/local_headers/non_cuda_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK: #include "non_cuda.h" diff --git a/tests/unit_tests/headers/local_headers/recursive-header.cu b/tests/unit_tests/headers/local_headers/recursive-header.cu index f712105ad..9f2c48fde 100644 --- a/tests/unit_tests/headers/local_headers/recursive-header.cu +++ b/tests/unit_tests/headers/local_headers/recursive-header.cu @@ -1,5 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers-recursive -// %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers-recursive %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/relative_path.cu b/tests/unit_tests/headers/local_headers/relative_path.cu index b5a75791c..800e1dd01 100644 --- a/tests/unit_tests/headers/local_headers/relative_path.cu +++ b/tests/unit_tests/headers/local_headers/relative_path.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/shared_header_1.cu b/tests/unit_tests/headers/local_headers/shared_header_1.cu index 669bedbad..f3815e6d2 100644 --- a/tests/unit_tests/headers/local_headers/shared_header_1.cu +++ b/tests/unit_tests/headers/local_headers/shared_header_1.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/shared_header_2.cu b/tests/unit_tests/headers/local_headers/shared_header_2.cu index 669bedbad..f3815e6d2 100644 --- a/tests/unit_tests/headers/local_headers/shared_header_2.cu +++ b/tests/unit_tests/headers/local_headers/shared_header_2.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/single_header.cu b/tests/unit_tests/headers/local_headers/single_header.cu index d8efd6154..d1f415042 100644 --- a/tests/unit_tests/headers/local_headers/single_header.cu +++ b/tests/unit_tests/headers/local_headers/single_header.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/stress_test.cu b/tests/unit_tests/headers/local_headers/stress_test.cu index 95022efea..ffbf59347 100644 --- a/tests/unit_tests/headers/local_headers/stress_test.cu +++ b/tests/unit_tests/headers/local_headers/stress_test.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/unit_tests/headers/local_headers/transitive_child.h b/tests/unit_tests/headers/local_headers/transitive_child.h index 50549fa6f..ff86eaeb8 100644 --- a/tests/unit_tests/headers/local_headers/transitive_child.h +++ b/tests/unit_tests/headers/local_headers/transitive_child.h @@ -1,8 +1,11 @@ #ifndef TRANSITIVE_CHILD_H #define TRANSITIVE_CHILD_H +#include + inline void child_sort_alloc(int *data, int n, void **p) { std::sort(data, data + n); + // CHECK: hipMalloc(p, n * sizeof(int)); cudaMalloc(p, n * sizeof(int)); } diff --git a/tests/unit_tests/headers/local_headers/transitive_system_include.cu b/tests/unit_tests/headers/local_headers/transitive_system_include.cu index 3967917b8..235e9e8a5 100644 --- a/tests/unit_tests/headers/local_headers/transitive_system_include.cu +++ b/tests/unit_tests/headers/local_headers/transitive_system_include.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %hipify_args --local-headers %clang_args +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers-recursive %clang_args // CHECK: #include // CHECK-NOT: #include From 7bc6edfe99775f1cc83da2d8081844f0ceba8758 Mon Sep 17 00:00:00 2001 From: ranapratap55 Date: Mon, 3 Aug 2026 13:38:52 +0530 Subject: [PATCH 10/10] [HIPIFY][fix] Single-pass include collection and local header context injection --- src/LocalHeader.cpp | 271 +++++++++--------- src/LocalHeader.h | 16 +- src/main.cpp | 26 +- tests/lit.cfg | 2 + .../local_headers/non_self_contained.cu | 23 ++ .../headers/local_headers/nsc_context.h | 15 + .../headers/local_headers/nsc_dependent.h | 12 + 7 files changed, 214 insertions(+), 151 deletions(-) create mode 100644 tests/unit_tests/headers/local_headers/non_self_contained.cu create mode 100644 tests/unit_tests/headers/local_headers/nsc_context.h create mode 100644 tests/unit_tests/headers/local_headers/nsc_dependent.h diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 6a36705e4..e44172865 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -24,15 +24,17 @@ THE SOFTWARE. #include "LLVMCompat.h" #include -#include +#include #include +#include "clang/Basic/SourceManager.h" #include "clang/Frontend/CompilerInstance.h" #include "clang/Frontend/FrontendActions.h" #include "clang/Lex/PPCallbacks.h" #include "clang/Lex/Preprocessor.h" #include "clang/Tooling/CompilationDatabase.h" #include "clang/Tooling/Tooling.h" +#include "llvm/ADT/StringSet.h" #include "llvm/Support/Path.h" #include "llvm/Support/raw_ostream.h" @@ -43,79 +45,97 @@ using namespace std; namespace { +std::string getFilePathForID(const clang::SourceManager &SM, + clang::FileID FID) { + const clang::FileEntry *FE = SM.getFileEntryForID(FID); + if (!FE) + return std::string(); + StringRef RealPath = FE->tryGetRealPathName(); + if (!RealPath.empty()) + return RealPath.str(); + return SM.getFilename(SM.getLocForStartOfFile(FID)).str(); +} + +// Records every `#include` of a translation unit, except those issued from a +// system header or from the command line. class IncludeCollectorCallbacks : public clang::PPCallbacks { const clang::SourceManager &SM; std::vector &Entries; public: IncludeCollectorCallbacks(const clang::SourceManager &SM, - std::vector &entries) - : SM(SM), Entries(entries) {} + std::vector &Entries) + : SM(SM), Entries(Entries) {} - void InclusionDirective(clang::SourceLocation hash_loc, - const clang::Token &include_token, - StringRef file_name, bool is_angled, - clang::CharSourceRange filename_range, + void InclusionDirective(clang::SourceLocation HashLoc, const clang::Token &, + StringRef FileName, bool IsAngled, + clang::CharSourceRange, #if LLVM_VERSION_MAJOR < 15 - const clang::FileEntry *file, + const clang::FileEntry *File, #elif LLVM_VERSION_MAJOR == 15 - Optional file, + Optional File, #else - clang::OptionalFileEntryRef file, + clang::OptionalFileEntryRef File, #endif - StringRef search_path, StringRef relative_path, + StringRef, StringRef, #if LLVM_VERSION_MAJOR < 19 - const clang::Module *SuggestedModule + const clang::Module * #else - const clang::Module *SuggestedModule, - bool ModuleImported + const clang::Module *, bool #endif #if LLVM_VERSION_MAJOR > 6 , clang::SrcMgr::CharacteristicKind FileType #endif ) override { - if (!SM.isWrittenInMainFile(hash_loc)) + const clang::FileID IncluderID = SM.getFileID(HashLoc); + std::string IncluderPath = getFilePathForID(SM, IncluderID); + if (IncluderPath.empty() || SM.isInSystemHeader(HashLoc)) return; - IncludeEntry entry; - entry.fileName = file_name.str(); - entry.isAngled = is_angled; + IncludeEntry Entry; + Entry.fileName = FileName.str(); + Entry.isAngled = IsAngled; + Entry.includerPath = std::move(IncluderPath); + Entry.isFromMainFile = IncluderID == SM.getMainFileID(); +#if LLVM_VERSION_MAJOR > 6 + Entry.isSystem = FileType != clang::SrcMgr::C_User; +#endif - if (file) { + if (File) { #if LLVM_VERSION_MAJOR < 15 - entry.resolvedPath = file->tryGetRealPathName().str(); - if (entry.resolvedPath.empty()) - entry.resolvedPath = file->getName().str(); + Entry.resolvedPath = File->tryGetRealPathName().str(); + if (Entry.resolvedPath.empty()) + Entry.resolvedPath = File->getName().str(); #else - entry.resolvedPath = file->getFileEntry().tryGetRealPathName().str(); - if (entry.resolvedPath.empty()) - entry.resolvedPath = file->getName().str(); + Entry.resolvedPath = File->getFileEntry().tryGetRealPathName().str(); + if (Entry.resolvedPath.empty()) + Entry.resolvedPath = File->getName().str(); #endif } - Entries.push_back(std::move(entry)); + Entries.push_back(std::move(Entry)); } }; -class IncludeCollectorAction : public clang::PreprocessorFrontendAction { +// PreprocessOnlyAction supplies the lexing loop and IgnorePragmas(). +class IncludeCollectorAction : public clang::PreprocessOnlyAction { std::vector &Entries; public: - explicit IncludeCollectorAction(std::vector &entries) - : Entries(entries) {} + explicit IncludeCollectorAction(std::vector &Entries) + : Entries(Entries) {} - void ExecuteAction() override { - clang::CompilerInstance &CI = getCompilerInstance(); - clang::Preprocessor &PP = CI.getPreprocessor(); - PP.addPPCallbacks(std::make_unique( - CI.getSourceManager(), Entries)); - - PP.EnterMainSourceFile(); - clang::Token Tok; - do { - PP.Lex(Tok); - } while (Tok.isNot(clang::tok::eof)); +protected: +#if LLVM_VERSION_MAJOR < 5 + bool BeginSourceFileAction(clang::CompilerInstance &CI, StringRef) override { +#else + bool BeginSourceFileAction(clang::CompilerInstance &CI) override { +#endif + CI.getPreprocessor().addPPCallbacks( + std::make_unique(CI.getSourceManager(), + Entries)); + return true; } }; @@ -137,6 +157,49 @@ class IncludeCollectorActionFactory : public FrontendActionFactory { #endif }; +// Returns entries.size() when headerPath was never included. +size_t findFirstInclusion(const std::vector &entries, + StringRef headerPath) { + size_t i = 0; + for (; i < entries.size(); ++i) + if (entries[i].resolvedPath == headerPath) + break; + return i; +} + +// Files to `-include` in front of headerPath so that a header which is not +// self-contained still sees what its ancestors included before it. +std::vector +buildIncludeContext(const std::vector &entries, + StringRef headerPath) { + const size_t pos = findFirstInclusion(entries, headerPath); + if (pos == entries.size()) + return std::vector(); + + StringSet<> ancestors; + std::string file = entries[pos].includerPath; + while (!file.empty() && ancestors.insert(file).second) { + const size_t idx = findFirstInclusion(entries, file); + file = idx == entries.size() ? std::string() : entries[idx].includerPath; + } + + std::vector context; + StringSet<> seen; + for (size_t i = 0; i < pos; ++i) { + const IncludeEntry &e = entries[i]; + // An ancestor would re-include headerPath, whose guard then empties the + // copy being hipified. + if (!ancestors.count(e.includerPath) || ancestors.count(e.resolvedPath)) + continue; + // System headers are injected as spelled, to be found via the search paths. + std::string arg = e.isSystem ? e.fileName : e.resolvedPath; + if (arg.empty() || !seen.insert(arg).second) + continue; + context.push_back(std::move(arg)); + } + return context; +} + } // namespace bool appendArgumentsAdjusters(ct::RefactoringTool &Tool, @@ -147,134 +210,74 @@ bool collectIncludeTree(const std::string &srcPath, const ct::CompilationDatabase *compDB, ct::CommonOptionsParser *OptionsParserPtr, const char *hipify_exe, - const std::string &mainContextPath, std::vector &outEntries) { outEntries.clear(); - const ct::CompilationDatabase &baseDB = - compDB ? *compDB : OptionsParserPtr->getCompilations(); - - // If srcPath has no entry in the compilation database, fall back to a - // FixedCompilationDatabase rooted at the mainContextPath's directory so that - // the tool doesn't skip the file. - std::vector cmds = baseDB.getCompileCommands(srcPath); - std::unique_ptr fallbackDB; - if (cmds.empty()) { - std::string dir = sys::path::parent_path(mainContextPath).str(); - fallbackDB = std::make_unique( - dir, std::vector()); - } - - ct::RefactoringTool Tool(fallbackDB ? *fallbackDB : baseDB, {srcPath}); + ct::RefactoringTool Tool( + compDB ? *compDB : OptionsParserPtr->getCompilations(), {srcPath}); - if (!appendArgumentsAdjusters(Tool, mainContextPath, hipify_exe)) { + if (!appendArgumentsAdjusters(Tool, srcPath, hipify_exe)) { return false; } IncludeCollectorActionFactory factory(outEntries); - Tool.run(&factory); - return true; + return Tool.run(&factory) == 0; } -bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, - const ct::CompilationDatabase *compDB, - ct::CommonOptionsParser *OptionsParserPtr, - const char *hipify_exe, - std::vector &outHeaders) { +bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe, bool recursive) { std::vector entries; if (!collectIncludeTree(mainSourceAbsPath, compDB, OptionsParserPtr, - hipify_exe, mainSourceAbsPath, entries)) { + hipify_exe, entries)) { errs() << "\n" << sHipify << sError << "Failed to collect includes from: " << mainSourceAbsPath << "\n"; return false; } - std::set uniq; - for (const auto &e : entries) { - if (!e.isAngled && !e.resolvedPath.empty()) - uniq.insert(e.resolvedPath); - } - outHeaders.assign(uniq.begin(), uniq.end()); - return true; -} - -bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, - const ct::CompilationDatabase *compDB, - ct::CommonOptionsParser *OptionsParserPtr, - const char *hipify_exe, bool recursive) { - - std::vector initial; - if (!collectLocalQuotedIncludes(mainSourceAbsPath, compDB, OptionsParserPtr, - hipify_exe, initial)) { - return false; + // The single run already walked the whole tree; recursion is just the + // unfiltered result. Include guards make each path appear at most once. + std::vector headers; + StringSet<> seen; + for (const IncludeEntry &e : entries) { + if (e.isAngled || e.isSystem || e.resolvedPath.empty()) + continue; + if (!recursive && !e.isFromMainFile) + continue; + if (seen.insert(e.resolvedPath).second) + headers.push_back(e.resolvedPath); } - if (initial.empty()) { + if (headers.empty()) { outs() << "\n" << sHipify << "No local headers detected in " << sys::path::filename(mainSourceAbsPath) << "\n"; return true; } - outs() << "\n" << sHipify << "Local headers found: " << initial.size() + outs() << "\n" << sHipify << "Local headers found: " << headers.size() << " in " << sys::path::filename(mainSourceAbsPath) << "\n"; - for (size_t i = 0; i < initial.size(); ++i) { - outs() << (i + 1) << "/" << initial.size() - << ": " << sys::path::filename(initial[i]) << "\n"; + for (size_t i = 0; i < headers.size(); ++i) { + outs() << (i + 1) << "/" << headers.size() << ": " + << sys::path::filename(headers[i]) << "\n"; } - std::vector work(initial.begin(), initial.end()); - std::set processed; - std::set queued(initial.begin(), initial.end()); - size_t total = initial.size(); - size_t current = 0; - - while (!work.empty()) { - std::string hdr = work.back(); - work.pop_back(); - if (processed.count(hdr)) { - continue; - } - processed.insert(hdr); - ++current; - + for (size_t i = 0; i < headers.size(); ++i) { + const std::string &hdr = headers[i]; std::string hipOut = hdr + ".hip"; - bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, - hipify_exe, mainSourceAbsPath, false); - - if (!ok) { - errs() << "\n" << sHipify << sError - << "Hipify failed for header [" << current << "/" << total - << "]: " << sys::path::filename(hdr) << "\n"; + if (!hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, hipify_exe, + mainSourceAbsPath, false, + buildIncludeContext(entries, hdr))) { + errs() << "\n" << sHipify << sError << "Hipify failed for header [" + << (i + 1) << "/" << headers.size() << "]: " + << sys::path::filename(hdr) << "\n"; return false; } outs() << sHipify << "Successfully hipified header file" << "\n"; - - if (recursive) { - std::vector childEntries; - if (collectIncludeTree(hdr, compDB, OptionsParserPtr, hipify_exe, - mainSourceAbsPath, childEntries)) { - std::vector newHeaders; - for (const auto &e : childEntries) { - if (!e.isAngled && !e.resolvedPath.empty() && - !processed.count(e.resolvedPath) && - !queued.count(e.resolvedPath)) { - newHeaders.push_back(e.resolvedPath); - work.push_back(e.resolvedPath); - queued.insert(e.resolvedPath); - } - } - if (!newHeaders.empty()) { - total += newHeaders.size(); - outs() << sHipify << " Recursive: found " << newHeaders.size() - << " additional local header(s) in " - << sys::path::filename(hdr) << "\n"; - } - } - } } outs() << "\n" << sHipify << "Local header hipification complete: " - << processed.size() << " header(s) processed.\n"; + << headers.size() << " header(s) processed.\n"; return true; } diff --git a/src/LocalHeader.h b/src/LocalHeader.h index caaba0970..bf2af038f 100644 --- a/src/LocalHeader.h +++ b/src/LocalHeader.h @@ -22,8 +22,6 @@ THE SOFTWARE. #pragma once -#include "clang/Frontend/FrontendActions.h" -#include "clang/Lex/PPCallbacks.h" #include "clang/Tooling/CommonOptionsParser.h" #include "clang/Tooling/Refactoring.h" #include @@ -31,10 +29,14 @@ THE SOFTWARE. namespace ct = clang::tooling; +// A single `#include` directive, recorded in preprocessing order. struct IncludeEntry { std::string fileName; std::string resolvedPath; - bool isAngled; + std::string includerPath; + bool isAngled = false; + bool isSystem = false; + bool isFromMainFile = false; }; extern bool hipifySingleSource(const std::string &srcPath, @@ -52,15 +54,9 @@ bool hipifyLocalHeaders(const std::string &srcPath, const char *hipify_exe, bool recursive = false); +// Preprocesses srcPath and records its whole include tree in outEntries. bool collectIncludeTree(const std::string &srcPath, const ct::CompilationDatabase *compDB, ct::CommonOptionsParser *OptionsParserPtr, const char *hipify_exe, - const std::string &mainContextPath, std::vector &outEntries); - -bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, - const ct::CompilationDatabase *compDB, - ct::CommonOptionsParser *OptionsParserPtr, - const char *hipify_exe, - std::vector &outHeaders); diff --git a/src/main.cpp b/src/main.cpp index 572872284..2742810b0 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -278,13 +278,25 @@ bool hipifySingleSource(const std::string &srcPath, return false; } - for (auto it = additionalIncludes.rbegin(); it != additionalIncludes.rend(); ++it) { - Tool.appendArgumentsAdjuster( - ct::getInsertArgumentAdjuster(it->c_str(), - ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster( - ct::getInsertArgumentAdjuster("-include", - ct::ArgumentInsertPosition::BEGIN)); + // The copy is preprocessed from a temporary directory, so includes quoted + // relative to the original one need an explicit search path. + StringRef srcDir = sys::path::parent_path(srcPath); + if (!srcDir.empty()) { + std::string sSrcDir = "-I" + srcDir.str(); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster( + sSrcDir.c_str(), ct::ArgumentInsertPosition::BEGIN)); + } + + // Appended at the end to follow the implicit CUDA headers, inserted at the + // beginning: a header may use CUDA declarations without including any. + if (!additionalIncludes.empty()) { + ct::CommandLineArguments includeArgs; + for (const std::string &include : additionalIncludes) { + includeArgs.push_back("-include"); + includeArgs.push_back(include); + } + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster( + includeArgs, ct::ArgumentInsertPosition::END)); } // Hipify _all_ the things! diff --git a/tests/lit.cfg b/tests/lit.cfg index 332a64637..f896836e3 100644 --- a/tests/lit.cfg +++ b/tests/lit.cfg @@ -24,6 +24,8 @@ config.excludes.append('transitive_child.h') config.excludes.append('diamond_left.h') config.excludes.append('diamond_right.h') config.excludes.append('diamond_shared.h') +# Not self-contained by design; only hipifiable through non_self_contained.cu. +config.excludes.append('nsc_dependent.h') delimiter = "==============================================================="; print(delimiter) diff --git a/tests/unit_tests/headers/local_headers/non_self_contained.cu b/tests/unit_tests/headers/local_headers/non_self_contained.cu new file mode 100644 index 000000000..ed787c5a7 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/non_self_contained.cu @@ -0,0 +1,23 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --local-headers %clang_args + +// nsc_dependent.h uses a type from nsc_context.h without including it, so it +// hipifies only if the includes preceding it here are replayed in front of it. + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "nsc_context.h" +// CHECK: #include "nsc_dependent.h" +#include +#include "nsc_context.h" +#include "nsc_dependent.h" + +__global__ void nscKernel(NscVec* vectors) { + int idx = threadIdx.x; + vectors[idx].data = make_float3(0.0f, 0.0f, 0.0f); +} + +int main() { + NscVec vec; + nsc_fill(&vec, 1); + return 0; +} diff --git a/tests/unit_tests/headers/local_headers/nsc_context.h b/tests/unit_tests/headers/local_headers/nsc_context.h new file mode 100644 index 000000000..8a9691f56 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/nsc_context.h @@ -0,0 +1,15 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +#ifndef NSC_CONTEXT_H +#define NSC_CONTEXT_H + +// CHECK: #include +// CHECK-NOT: #include +#include + +struct NscVec { + float3 data; + cudaError_t status; +}; + +#endif diff --git a/tests/unit_tests/headers/local_headers/nsc_dependent.h b/tests/unit_tests/headers/local_headers/nsc_dependent.h new file mode 100644 index 000000000..4553c07eb --- /dev/null +++ b/tests/unit_tests/headers/local_headers/nsc_dependent.h @@ -0,0 +1,12 @@ +#ifndef NSC_DEPENDENT_H +#define NSC_DEPENDENT_H + +// Deliberately not self-contained: NscVec comes from the includer. Driven by +// non_self_contained.cu, and excluded from standalone testing in tests/lit.cfg. + +inline cudaError_t nsc_fill(NscVec *v, int count) { + v->status = cudaMemset(&v->data, 0, count * sizeof(float3)); + return v->status; +} + +#endif