Evgeny Mankov
e90a76a1ef
If in source CUDA file main header (cuda_runtime.h or cuda.h) is not presented, corresponding HIP main header (hip_runtime.h) should be explicitly included in output hipified file. [Algorithm] 1. If #pragma once is presented, HIP main header should be placed just after it; 2. Otherwise if any other (not CUDA main) header is presented, HIP main header should be placed just before it; 3. Otherwise HIP main header should be placed in the beginning of output file. P.S. There might be one more situation when #ifndef #define ... #endif guard for the entire file is presented (make sense for *.h, *.hpp, *.cuh files). In this case HIP main include should be placed just after such #ifdef, or after #pragma once, if it is also presented. This situation will be handled in a separate change.
93 lines
3.5 KiB
C++
93 lines
3.5 KiB
C++
#pragma once
|
|
|
|
#include "clang/Lex/PPCallbacks.h"
|
|
#include "clang/Tooling/Tooling.h"
|
|
#include "clang/Frontend/FrontendAction.h"
|
|
#include "clang/Tooling/Core/Replacement.h"
|
|
#include "clang/ASTMatchers/ASTMatchFinder.h"
|
|
#include "ReplacementsFrontendActionFactory.h"
|
|
|
|
namespace ct = clang::tooling;
|
|
|
|
/**
|
|
* A FrontendAction that hipifies CUDA programs.
|
|
*/
|
|
class HipifyAction : public clang::ASTFrontendAction,
|
|
public clang::ast_matchers::MatchFinder::MatchCallback {
|
|
private:
|
|
ct::Replacements* replacements;
|
|
std::unique_ptr<clang::ast_matchers::MatchFinder> Finder;
|
|
|
|
/// CUDA implicitly adds its runtime header. We rewrite explicitly-provided CUDA includes with equivalent
|
|
// ones, and track - using this flag - if the result led to us including the hip runtime header. If it did
|
|
// not, we insert it at the top of the file when we finish processing it.
|
|
// This approach means we do the best it's possible to do w.r.t preserving the user's include order.
|
|
bool insertedRuntimeHeader = false;
|
|
bool firstNotMainHeader = false;
|
|
bool pragmaOnce = false;
|
|
clang::SourceLocation firstNotMainHeaderLoc;
|
|
clang::SourceLocation pragmaOnceLoc;
|
|
|
|
/**
|
|
* Rewrite a string literal to refer to hip, not CUDA.
|
|
*/
|
|
void RewriteString(StringRef s, clang::SourceLocation start);
|
|
|
|
/**
|
|
* Replace a CUDA identifier with the corresponding hip identifier, if applicable.
|
|
*/
|
|
void RewriteToken(const clang::Token &t);
|
|
|
|
public:
|
|
explicit HipifyAction(ct::Replacements *replacements):
|
|
clang::ASTFrontendAction(),
|
|
replacements(replacements) {}
|
|
|
|
// MatchCallback listeners
|
|
bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
|
bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
|
bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
|
|
|
/**
|
|
* Called by the preprocessor for each include directive during the non-raw lexing pass.
|
|
*/
|
|
void InclusionDirective(clang::SourceLocation hash_loc,
|
|
const clang::Token &include_token,
|
|
StringRef file_name,
|
|
bool is_angled,
|
|
clang::CharSourceRange filename_range,
|
|
const clang::FileEntry *file,
|
|
StringRef search_path,
|
|
StringRef relative_path,
|
|
const clang::Module *imported);
|
|
|
|
/**
|
|
* Called by the preprocessor for each pragma directive during the non-raw lexing pass.
|
|
*/
|
|
void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer);
|
|
|
|
protected:
|
|
/**
|
|
* Add a Replacement for the current file. These will all be applied after executing the FrontendAction.
|
|
*/
|
|
void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL);
|
|
|
|
/**
|
|
* FrontendAction entry point.
|
|
*/
|
|
void ExecuteAction() override;
|
|
|
|
/**
|
|
* Called at the start of each new file to process.
|
|
*/
|
|
void EndSourceFileAction() override;
|
|
|
|
/**
|
|
* MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an
|
|
* interest for.
|
|
*/
|
|
void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override;
|
|
|
|
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
|
|
};
|