diff --git a/src/HeaderInjection.cpp b/src/HeaderInjection.cpp new file mode 100644 index 000000000..16d32b940 --- /dev/null +++ b/src/HeaderInjection.cpp @@ -0,0 +1,303 @@ +/* +Copyright (c) 2015 - 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 "HeaderInjection.h" +#include "LocalHeader.h" +#include "LLVMCompat.h" + +#include +#include +#include + +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/Path.h" +#include "llvm/Support/MemoryBuffer.h" +#include "llvm/Support/raw_ostream.h" + +using namespace llvm; +using namespace std; + +namespace { + +// Matches system/library includes +static const regex SystemIncludeRe( + R"(^\s*#\s*include\s*<([^>\n]+)>)", regex::ECMAScript); + +// Matches local (quoted) includes +static const regex LocalIncludeRe( + R"(^\s*#\s*include\s*\"([^\"\n]+)\")", regex::ECMAScript); + +// Matches #ifndef guard +static const regex IfndefGuardRe( + R"(^\s*#\s*ifndef\s+(\w+)\s*$)", regex::ECMAScript); + +// Matches #define for guard +static const regex DefineGuardRe( + R"(^\s*#\s*define\s+(\w+)\s*$)", regex::ECMAScript); + +static const regex PragmaOnceRe( + R"(^\s*#\s*pragma\s+once\s*$)", regex::ECMAScript); + +bool readFileContent(const string &path, string &out) { + auto MBOrErr = MemoryBuffer::getFile(path); + if (!MBOrErr) return false; + out = MBOrErr->get()->getBuffer().str(); + return true; +} + +string extractIncludePath(const string &line) { + smatch m; + if (regex_search(line, m, SystemIncludeRe)) { + return m[1].str(); + } + return ""; +} + +} + +bool collectPrecedingSystemIncludes(const string &mainSourceAbsPath, + const string &targetHeaderAbsPath, + vector &outSystemIncludes) { + string content; + if (!readFileContent(mainSourceAbsPath, content)) { + errs() << sHipify << sError << "Cannot read source file: " << mainSourceAbsPath << "\n"; + return false; + } + + string targetFileName = string(sys::path::filename(targetHeaderAbsPath)); + + istringstream iss(content); + string line; + smatch sysMatch, localMatch; + + while (getline(iss, line)) { + if (regex_search(line, localMatch, LocalIncludeRe)) { + string localInc = localMatch[1].str(); + string localFileName = string(sys::path::filename(localInc)); + if (localFileName == targetFileName) { + break; + } + continue; + } + + if (regex_search(line, sysMatch, SystemIncludeRe)) { + outSystemIncludes.push_back(line); + } + } + + return true; +} + +void detectIncludeGuard(const string &headerContent, + size_t &guardEndLine, + string &guardType) { + guardEndLine = 0; + guardType = "none"; + + istringstream iss(headerContent); + string line; + size_t lineNum = 0; + string ifndefSymbol; + + while (getline(iss, line)) { + smatch m; + + if (regex_match(line, PragmaOnceRe)) { + guardType = "pragma_once"; + guardEndLine = lineNum; + return; + } + + if (regex_match(line, m, IfndefGuardRe)) { + ifndefSymbol = m[1].str(); + for (int i = 0; i < 5 && getline(iss, line); ++i) { + lineNum++; + if (regex_match(line, m, DefineGuardRe)) { + if (m[1].str() == ifndefSymbol) { + guardType = "ifndef"; + guardEndLine = lineNum; + return; + } + } + if (line.empty() || line.find("//") == 0 || line.find("/*") == 0) { + continue; + } + break; + } + } + + lineNum++; + } +} + +void getExistingIncludes(const string &headerContent, + set &existingIncludes) { + istringstream iss(headerContent); + string line; + smatch m; + + while (getline(iss, line)) { + if (regex_search(line, m, SystemIncludeRe)) { + existingIncludes.insert(m[1].str()); + } + } +} + +bool createInjectedHeader(const string &mainSourceAbsPath, + const string &targetHeaderAbsPath, + const string &injectedFilePath) { + string headerContent; + if (!readFileContent(targetHeaderAbsPath, headerContent)) { + errs() << sHipify << sError << "Cannot read target header: " << targetHeaderAbsPath << "\n"; + return false; + } + + vector systemIncludes; + if (!collectPrecedingSystemIncludes(mainSourceAbsPath, targetHeaderAbsPath, + systemIncludes)) { + } + + set existingIncludes; + getExistingIncludes(headerContent, existingIncludes); + + vector uniqueIncludes; + for (const auto &inc : systemIncludes) { + string path = extractIncludePath(inc); + if (!path.empty() && existingIncludes.find(path) == existingIncludes.end()) { + uniqueIncludes.push_back(inc); + existingIncludes.insert(path); + } + } + + if (uniqueIncludes.empty()) { + ofstream out(injectedFilePath); + if (!out.is_open()) { + errs() << sHipify << sError << "Cannot create injected file: " << injectedFilePath << "\n"; + return false; + } + out << headerContent; + out.close(); + return true; + } + + size_t guardEndLine; + string guardType; + detectIncludeGuard(headerContent, guardEndLine, guardType); + + string mainFileName = string(sys::path::filename(mainSourceAbsPath)); + ostringstream injection; + injection << "// --- HIPIFY: Injected dependencies from " << mainFileName << " ---\n"; + for (const auto &inc : uniqueIncludes) { + injection << inc << "\n"; + } + injection << "// --- End injected dependencies ---\n"; + injection << "\n"; + + ofstream out(injectedFilePath); + if (!out.is_open()) { + errs() << sHipify << sError << "Cannot create injected file: " << injectedFilePath << "\n"; + return false; + } + + istringstream iss(headerContent); + string line; + size_t lineNum = 0; + bool injected = false; + + while (getline(iss, line)) { + out << line << "\n"; + + if (!injected && lineNum == guardEndLine && guardType != "none") { + out << injection.str(); + injected = true; + } + + lineNum++; + } + + if (!injected && guardType == "none") { + out.close(); + ofstream outNew(injectedFilePath); + if (!outNew.is_open()) { + errs() << sHipify << sError << "Cannot create injected file: " << injectedFilePath << "\n"; + return false; + } + outNew << injection.str(); + outNew << headerContent; + outNew.close(); + } else { + out.close(); + } + + return true; +} + +bool hipifyHeaderWithInjection(const string &headerAbsPath, + const string &outputPath, + const string &mainSourceAbsPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe) { + string headerStem = string(sys::path::stem(headerAbsPath)); + string headerExt = string(sys::path::extension(headerAbsPath)); + + if (!headerExt.empty() && headerExt[0] == '.') { + headerExt = headerExt.substr(1); + } + if (headerExt.empty()) { + headerExt = "h"; + } + + string tempPrefix = "inject_" + headerStem; + + SmallString<256> injectedPath; + error_code EC = sys::fs::createTemporaryFile(tempPrefix, headerExt, injectedPath); + if (EC) { + errs() << sHipify << sError << "Cannot create temporary file: " << EC.message() << "\n"; + return false; + } + + if (!createInjectedHeader(mainSourceAbsPath, headerAbsPath, string(injectedPath.str()))) { + sys::fs::remove(injectedPath); + return false; + } + + bool hipifyOk = hipifySingleSource( + string(injectedPath.str()), + outputPath, + compDB, + OptionsParserPtr, + hipify_exe, + mainSourceAbsPath, + false + ); + + sys::fs::remove(injectedPath); + + if (!hipifyOk) { + errs() << sHipify << sError << "Failed to hipify (injection): " << headerAbsPath << "\n"; + return false; + } + + return true; +} + diff --git a/src/HeaderInjection.h b/src/HeaderInjection.h new file mode 100644 index 000000000..7394a484e --- /dev/null +++ b/src/HeaderInjection.h @@ -0,0 +1,53 @@ +/* +Copyright (c) 2015 - 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 +#include +#include +#include "clang/Tooling/CommonOptionsParser.h" + +namespace ct = clang::tooling; + +bool collectPrecedingSystemIncludes(const std::string &mainSourceAbsPath, + const std::string &targetHeaderAbsPath, + std::vector &outSystemIncludes); + +void detectIncludeGuard(const std::string &headerContent, + size_t &guardEndLine, + std::string &guardType); + +void getExistingIncludes(const std::string &headerContent, + std::set &existingIncludes); + +bool createInjectedHeader(const std::string &mainSourceAbsPath, + const std::string &targetHeaderAbsPath, + const std::string &injectedFilePath); + +bool hipifyHeaderWithInjection(const std::string &headerAbsPath, + const std::string &outputPath, + const std::string &mainSourceAbsPath, + const ct::CompilationDatabase *compDB, + ct::CommonOptionsParser *OptionsParserPtr, + const char *hipify_exe); + diff --git a/src/LocalHeader.cpp b/src/LocalHeader.cpp index 7d1d129bd..bdf3cdf62 100644 --- a/src/LocalHeader.cpp +++ b/src/LocalHeader.cpp @@ -1,101 +1,119 @@ +/* +Copyright (c) 2015 - 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 "HeaderInjection.h" +#include "StderrCapture.h" +#include "LLVMCompat.h" #include #include #include #include +#include #include +#include #include "llvm/Support/FileSystem.h" #include "llvm/Support/Path.h" #include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/raw_ostream.h" -#include "clang/Tooling/Refactoring.h" -#include "clang/Tooling/Tooling.h" - -#include "CUDA2HIP.h" -#include "LLVMCompat.h" -#include "HipifyAction.h" - -using namespace clang; -using namespace clang::tooling; using namespace llvm; using namespace std; +using hipify::StderrCapture; -static std::string normalizeSmallStringPath(SmallString<256> &p) { - llvm::sys::path::remove_dots(p, true); +// Matches local (quoted) includes +static const regex LocalIncludeRe( + R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?$)", regex::ECMAScript); + +static string normalizeSmallStringPath(SmallString<256> &p) { + sys::path::remove_dots(p, true); SmallString<256> realBuf; - std::error_code ec = llvm::sys::fs::real_path(p, realBuf); + error_code ec = sys::fs::real_path(p, realBuf); if (!ec) { - return std::string(realBuf.str()); + return string(realBuf.str()); } - return std::string(p.str()); + return string(p.str()); } -static bool pathExists(const std::string &p) { +static bool pathExists(const string &p) { SmallString<256> in(p.begin(), p.end()); SmallString<256> realBuf; - std::error_code ec = llvm::sys::fs::real_path(in, realBuf); + error_code ec = 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); + sys::path::remove_dots(norm, true); + return sys::fs::exists(norm); } -namespace { - static const std::regex LocalIncludeRe( - R"(^\s*#\s*include\s*\"([^\"\n]+)\"\s*(?://.*)?$)", std::regex::ECMAScript); +bool readFile(const string &path, string &out) { + auto MBOrErr = MemoryBuffer::getFile(path); + if (!MBOrErr) return false; + out = MBOrErr->get()->getBuffer().str(); + return true; +} - bool readFile(const std::string &path, std::string &out) { - auto MBOrErr = llvm::MemoryBuffer::getFile(path); - if (!MBOrErr) return false; - out = MBOrErr->get()->getBuffer().str(); +bool resolveLocalIncludeInternal(const string &mainSourceAbsPath, + const string &includeTok, + 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(string(candidate.str()))) { + outAbs = normalizeSmallStringPath(candidate); return true; } + return false; +} - 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; - } - return false; - } -} - -bool resolveLocalInclude(const std::string &mainSourceAbsPath, - const std::string &includeToken, - std::string &outAbsPath) { +bool resolveLocalInclude(const string &mainSourceAbsPath, + const string &includeToken, + string &outAbsPath) { return resolveLocalIncludeInternal(mainSourceAbsPath, includeToken, outAbsPath); } -bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, - std::vector &outHeaders) { - std::string content; +bool collectLocalQuotedIncludes(const string &mainSourceAbsPath, + vector &outHeaders) { + string content; if (!readFile(mainSourceAbsPath, content)) { errs() << "\n" << sHipify << sError << "Cannot read source file: " << 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; + set uniq; + smatch m; + istringstream iss(content); + string line; + while (getline(iss, line)) { + if (regex_match(line, m, LocalIncludeRe)) { + string rel = m[1].str(); + string abs; if (resolveLocalIncludeInternal(mainSourceAbsPath, rel, abs)){ uniq.insert(abs); } else { @@ -109,13 +127,13 @@ bool collectLocalQuotedIncludes(const std::string &mainSourceAbsPath, return true; } -bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, +bool hipifyLocalHeaders(const string &mainSourceAbsPath, const ct::CompilationDatabase *compDB, ct::CommonOptionsParser *OptionsParserPtr, const char *hipify_exe, bool recursive) { - std::vector initial; + vector initial; if (!collectLocalQuotedIncludes(mainSourceAbsPath, initial)) { return false; } @@ -125,41 +143,110 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, return true; } - std::vector work(initial.begin(), initial.end()); - std::set processed; + outs() << "\n"; + outs() << sHipify << "Found " << initial.size() << " local header(s) to process\n"; + outs() << sHipify << "Note: Compilation errors during direct attempts may be safely ignored\n"; + outs() << sHipify << " if the injection fallback succeeds.\n"; + outs() << "\n"; + outs().flush(); + + vector work(initial.begin(), initial.end()); + set processed; + + vector directSuccess; + vector injectionSuccess; + vector failed; + + // Store captured error output for failed files + map capturedErrors; while (!work.empty()) { - std::string hdr = work.back(); + 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; + string original; if (!readFile(hdr, original)) { errs() << sHipify << sError << "Cannot read header: " << hdr << "\n"; + failed.push_back(hdr); continue; } - std::string hipOut = hdr + ".hip"; - bool ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, - hipify_exe, mainSourceAbsPath, false); + string hipOut = hdr + ".hip"; + string hdrFileName = string(sys::path::filename(hdr)); + bool ok = false; - if (!ok) { - errs() << sHipify << sError << "Hipify failed for header: " << hdr << "\n"; - return false; + // HYBRID APPROACH: + // Step 1: Try direct hipification first (works for self-contained headers) + // Capture stderr - if both approaches fail, we'll show the errors + outs() << sHipify << "[" << (directSuccess.size() + injectionSuccess.size() + 1) + << "/" << initial.size() << "] Hipifying source: " << hdr << "\n"; + outs().flush(); + + string directErrors; + { + // Capture stderr during direct attempt + StderrCapture capture; + ok = hipifySingleSource(hdr, hipOut, compDB, OptionsParserPtr, + hipify_exe, mainSourceAbsPath, false, true); + if (!ok) { + directErrors = capture.getCapturedOutput(); + } } + if (ok) { + outs() << sHipify << " -> OK (direct)\n"; + directSuccess.push_back(hdrFileName); + } else { + // Step 2: If direct fails, inject preceding includes + outs() << sHipify << " -> Trying injection approach...\n"; + outs().flush(); + + string injectionErrors; + { + // Capture stderr during injection attempt + StderrCapture capture; + ok = hipifyHeaderWithInjection(hdr, hipOut, mainSourceAbsPath, + compDB, OptionsParserPtr, hipify_exe); + if (!ok) { + injectionErrors = capture.getCapturedOutput(); + } + } + + if (ok) { + outs() << sHipify << " -> OK (injection)\n"; + injectionSuccess.push_back(hdrFileName); + } else { + outs() << sHipify << " -> FAILED\n"; + failed.push_back(hdrFileName); + + // Store errors for this file - combine both attempts' errors + string combinedErrors; + if (!directErrors.empty()) { + combinedErrors += "=== Direct approach errors ===\n" + directErrors; + } + if (!injectionErrors.empty()) { + if (!combinedErrors.empty()) combinedErrors += "\n"; + combinedErrors += "=== Injection approach errors ===\n" + injectionErrors; + } + if (!combinedErrors.empty()) { + capturedErrors[hdrFileName] = combinedErrors; + } + } + } + + // If recursive, find and queue nested local headers 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; + smatch m; + istringstream iss(original); + string line; + while (getline(iss, line)) { + if (regex_match(line, m, LocalIncludeRe)) { + string rel = m[1].str(); + string abs; if (resolveLocalIncludeInternal(hdr, rel, abs) && !processed.count(abs)) work.push_back(abs); @@ -168,5 +255,54 @@ bool hipifyLocalHeaders(const std::string &mainSourceAbsPath, } } + outs() << "\n"; + outs() << sHipify << "Local Header Hipification Summary\n"; + + size_t total = directSuccess.size() + injectionSuccess.size() + failed.size(); + size_t success = directSuccess.size() + injectionSuccess.size(); + + if (!directSuccess.empty()) { + outs() << sHipify << " Direct: " << directSuccess.size() << " header(s)\n"; + } + if (!injectionSuccess.empty()) { + outs() << sHipify << " Injection: " << injectionSuccess.size() << " header(s) (needed deps from main source)\n"; + } + if (!failed.empty()) { + outs() << sHipify << " Failed: " << failed.size() << " header(s)\n"; + for (const auto &f : failed) { + outs() << sHipify << " - " << f << "\n"; + } + } + + outs() << sHipify << " Total: " << success << "/" << total << " succeeded\n"; + outs() << "\n"; + + // If there were failures, show the captured error details + if (!failed.empty()) { + errs() << sHipify << sError << "The following headers failed to hipify:\n"; + errs() << "\n"; + + for (const auto &f : failed) { + errs() << "━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━\n"; + errs() << " Failed: " << f << "\n"; + errs() << "━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━\n"; + + auto it = capturedErrors.find(f); + if (it != capturedErrors.end() && !it->second.empty()) { + errs() << it->second << "\n"; + } else { + errs() << " (No detailed error output captured)\n\n"; + } + } + + errs() << sHipify << "Hint: Check if the headers have:\n"; + errs() << sHipify << " - Missing #include dependencies\n"; + errs() << sHipify << " - Syntax errors\n"; + errs() << sHipify << " - Types/operators not available in HIP\n"; + errs() << "\n"; + + return false; + } + return true; } diff --git a/src/LocalHeader.h b/src/LocalHeader.h index 31763c697..086dec1f6 100644 --- a/src/LocalHeader.h +++ b/src/LocalHeader.h @@ -1,3 +1,25 @@ +/* +Copyright (c) 2015 - 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, + bool silent = false); bool hipifyLocalHeaders(const std::string &srcPath, const ct::CompilationDatabase *compDB, diff --git a/src/StderrCapture.h b/src/StderrCapture.h new file mode 100644 index 000000000..936902595 --- /dev/null +++ b/src/StderrCapture.h @@ -0,0 +1,173 @@ +/* +Copyright (c) 2015 - 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 +#include +#include +#include +#include + +#ifdef _WIN32 + #ifndef NOMINMAX + #define NOMINMAX + #endif + #include + #include + #include + #include + #define STDERR_FD _fileno(stderr) + #define DUP(fd) _dup(fd) + #define DUP2(fd1, fd2) _dup2(fd1, fd2) + #define CLOSE(fd) _close(fd) +#else + #include + #include + #define STDERR_FD STDERR_FILENO + #define DUP(fd) dup(fd) + #define DUP2(fd1, fd2) dup2(fd1, fd2) + #define CLOSE(fd) close(fd) +#endif + +namespace hipify { + +// Capture stderr output to a temp file. +class StderrCapture { +public: + StderrCapture() : lock_(getMutex()), saved_stderr_(-1), temp_fd_(-1), active_(false) { +#ifdef _WIN32 + char tempDir[MAX_PATH]; + DWORD tempDirLen = GetTempPathA(MAX_PATH, tempDir); + if (tempDirLen == 0 || tempDirLen > MAX_PATH) return; + + char tempFile[MAX_PATH]; + if (GetTempFileNameA(tempDir, "hip", 0, tempFile) == 0) return; + + tempFilePath_ = tempFile; + + temp_fd_ = _open(tempFile, _O_RDWR | _O_CREAT | _O_TRUNC, _S_IREAD | _S_IWRITE); + if (temp_fd_ == -1) { + DeleteFileA(tempFile); + tempFilePath_.clear(); + return; + } +#else + char tmpPath[] = "/tmp/hipify_stderr_XXXXXX"; + temp_fd_ = mkstemp(tmpPath); + if (temp_fd_ == -1) return; + + tempFilePath_ = tmpPath; +#endif + + saved_stderr_ = DUP(STDERR_FD); + if (saved_stderr_ == -1) { + CLOSE(temp_fd_); + removeTempFile(); + temp_fd_ = -1; + return; + } + + if (DUP2(temp_fd_, STDERR_FD) != -1) { + active_ = true; + } + } + + // Restores stderr and deletes temp file. + ~StderrCapture() { + restore(); + cleanup(); + } + + StderrCapture(const StderrCapture&) = delete; + StderrCapture& operator=(const StderrCapture&) = delete; + + // Restores stderr to original state. + void restore() { + if (active_ && saved_stderr_ != -1) { + fflush(stderr); + DUP2(saved_stderr_, STDERR_FD); + active_ = false; + } + if (saved_stderr_ != -1) { + CLOSE(saved_stderr_); + saved_stderr_ = -1; + } + } + + // Returns captured stderr content and restores stderr. + std::string getCapturedOutput() { + std::string content; + restore(); + + if (temp_fd_ != -1 && !tempFilePath_.empty()) { + std::ifstream file(tempFilePath_); + if (file.is_open()) { + std::stringstream buffer; + buffer << file.rdbuf(); + content = buffer.str(); + file.close(); + } + } + return content; + } + + bool isActive() const { return active_; } + +private: + static std::mutex& getMutex() { + static std::mutex mtx; + return mtx; + } + + void removeTempFile() { + if (!tempFilePath_.empty()) { +#ifdef _WIN32 + DeleteFileA(tempFilePath_.c_str()); +#else + unlink(tempFilePath_.c_str()); +#endif + tempFilePath_.clear(); + } + } + + void cleanup() { + if (temp_fd_ != -1) { + CLOSE(temp_fd_); + temp_fd_ = -1; + } + removeTempFile(); + } + + std::unique_lock lock_; + int saved_stderr_; + int temp_fd_; + bool active_; + std::string tempFilePath_; +}; + +} // namespace hipify + +#undef STDERR_FD +#undef DUP +#undef DUP2 +#undef CLOSE diff --git a/src/main.cpp b/src/main.cpp index b00dbcf2e..df08f9c1f 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -206,6 +206,7 @@ bool appendArgumentsAdjusters(ct::RefactoringTool &Tool, const std::string &sSou Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN)); } } + // Standard c++ to use in hipification by default llcompat::setStdCPP(Tool); std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str(); @@ -239,7 +240,13 @@ bool hipifySingleSource(const std::string &srcPath, ct::CommonOptionsParser *OptionsParserPtr, const char *hipify_exe_path, const std::string &mainContextPath, - bool preserveTemp) { + bool preserveTemp, + bool silent) { + // Print info message for each file being hipified (unless silent) + if (!silent) { + llvm::outs() << sHipify << "Hipifying source: " << srcPath << "\n"; + } + std::error_code EC; SmallString<128> tmpFile; StringRef srcFileName = sys::path::filename(srcPath); 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