diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 2d5e97545..6a36705e4 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -1,167 +1,280 @@ +/* +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" -#include -#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" -#include "LLVMCompat.h" - using namespace clang; using namespace clang::tooling; using namespace llvm; using namespace std; -static std::string normalizeSmallStringPath(SmallString<256> &p) { - llvm::sys::path::remove_dots(p, true); +namespace { + +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) {} - SmallString<256> realBuf; - std::error_code ec = llvm::sys::fs::real_path(p, realBuf); - if (!ec) { - return std::string(realBuf.str()); + 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 + } + + Entries.push_back(std::move(entry)); } +}; - return std::string(p.str()); -} +class IncludeCollectorAction : public clang::PreprocessorFrontendAction { + std::vector &Entries; -static bool pathExists(const std::string &p) { - SmallString<256> in(p.begin(), p.end()); +public: + explicit IncludeCollectorAction(std::vector &entries) + : Entries(entries) {} - SmallString<256> realBuf; - std::error_code ec = llvm::sys::fs::real_path(in, realBuf); - if (!ec) return true; + void ExecuteAction() override { + clang::CompilerInstance &CI = getCompilerInstance(); + clang::Preprocessor &PP = CI.getPreprocessor(); + PP.addPPCallbacks(std::make_unique( + CI.getSourceManager(), Entries)); - SmallString<256> norm = in; - llvm::sys::path::remove_dots(norm, true); - return llvm::sys::fs::exists(norm); -} + PP.EnterMainSourceFile(); + clang::Token Tok; + do { + PP.Lex(Tok); + } while (Tok.isNot(clang::tok::eof)); + } +}; -namespace { - static const std::regex LocalIncludeRe( - R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?$)", std::regex::ECMAScript); +class IncludeCollectorActionFactory : public FrontendActionFactory { + std::vector &Entries; - 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; +public: + explicit IncludeCollectorActionFactory(std::vector &entries) + : Entries(entries) {} + +#if LLVM_VERSION_MAJOR >= 10 + std::unique_ptr create() override { + return std::make_unique(Entries); } +#else + clang::FrontendAction *create() override { + return new IncludeCollectorAction(Entries); + } +#endif +}; - 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; - } +} // 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, + 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}); + + if (!appendArgumentsAdjusters(Tool, mainContextPath, hipify_exe)) { return false; } -} -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_match(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() << 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; + 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)) { - errs() << sHipify << sWarning << "Duplicate local header reference ignored: " << hdr << "\n"; continue; } processed.insert(hdr); - - std::string original; - if (!readFile(hdr, original)) { - errs() << sHipify << sError << "Cannot read header: " << hdr << "\n"; - continue; - } + ++current; std::string hipOut = hdr + ".hip"; bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, - hipify_exe, mainSourceAbsPath, false); + hipify_exe, mainSourceAbsPath, false); 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; - 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)) - work.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"; } } } } + 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..caaba0970 100644 --- a/src/LocalHeader.h +++ b/src/LocalHeader.h @@ -1,18 +1,50 @@ +/* +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 "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; +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, 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, @@ -20,9 +52,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 0675bdfce..572872284 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,26 +527,25 @@ 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) { - if (!hipifyLocalHeaders(sSourceAbsPath, - compilationDatabase.get(), - &OptionsParser, - argv[0], + 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; } } - + std::string outputPath = NoOutput ? "" : dst; if (!hipifySingleSource(src, outputPath, compilationDatabase.get(), &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/lit.cfg b/tests/lit.cfg index aee90a323..12e17d92c 100644 --- a/tests/lit.cfg +++ b/tests/lit.cfg @@ -19,6 +19,11 @@ 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') +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/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/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 3d3e6c692..059d1651a 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 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.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..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_has_cmath.h b/tests/unit_tests/headers/local_headers/injection_has_cmath.h new file mode 100644 index 000000000..fcc8ae96e --- /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 %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..150920889 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_helper.h @@ -0,0 +1,22 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %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..0c213c7fe --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_inner.h @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %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..0283d1047 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_multi.cu @@ -0,0 +1,22 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --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..05747193d --- /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 1 --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..e742b9699 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_outer.h @@ -0,0 +1,16 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %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..86f68bc53 --- /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 %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..4887bd592 --- /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 1 --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..e943933b5 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_recursive.cu @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --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..8bcc93801 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/injection_test.cu @@ -0,0 +1,17 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --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..8ea1b31f2 --- /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 %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/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..79b1bb617 100644 --- a/tests/unit_tests/headers/local_headers/main.cu +++ b/tests/unit_tests/headers/local_headers/main.cu @@ -1,12 +1,25 @@ -// 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 // 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/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.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..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/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/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..9f2c48fde 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 1 --local-headers-recursive %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..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.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_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/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/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..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/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 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..ff86eaeb8 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/transitive_child.h @@ -0,0 +1,12 @@ +#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)); +} + +#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..235e9e8a5 --- /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 1 --local-headers-recursive %clang_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include "transitive_parent.h" +#include +#include +#include "transitive_parent.h" + +int main() { return 0; } 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..ae4eb4a25 --- /dev/null +++ b/tests/unit_tests/headers/local_headers/vector_math.h @@ -0,0 +1,12 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %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