AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Merge 'master' into 'amd-master'

Sfoglia il codice sorgente 
Change-Id: I999da34da4bbfba663144dd778d9b32fb67a1a70
This commit is contained in:
Jenkins
2018-04-30 04:09:42 -05:00
parent 1964bb1115 b36001bbbf
commit c509cebcde
16 ha cambiato i file con 2469 aggiunte e 3996 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
hipify-clang/src/ArgParse.cpp
+27 -15
Vedi File
@@ -2,27 +2,39 @@
cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options");
cl::opt<std::string> OutputFilename("o", cl::desc("Output filename"), cl::value_desc("filename"),
cl::cat(ToolTemplateCategory));
cl::opt<std::string> OutputFilename("o",
cl::desc("Output filename"),
cl::value_desc("filename"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> Inplace("inplace",
cl::desc("Modify input file inplace, replacing input with hipified output, "
"save backup in .prehip file"),
cl::value_desc("inplace"), cl::cat(ToolTemplateCategory));
cl::desc("Modify input file inplace, replacing input with hipified output, save backup in .prehip file"),
cl::value_desc("inplace"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> NoBackup("no-backup", cl::desc("Don't create a backup file for the hipified source"),
cl::value_desc("no-backup"), cl::cat(ToolTemplateCategory));
cl::opt<bool> NoBackup("no-backup",
cl::desc("Don't create a backup file for the hipified source"),
cl::value_desc("no-backup"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> NoOutput("no-output", cl::desc("Don't write any translated output to stdout"),
cl::value_desc("no-output"), cl::cat(ToolTemplateCategory));
cl::opt<bool> NoOutput("no-output",
cl::desc("Don't write any translated output to stdout"),
cl::value_desc("no-output"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> PrintStats("print-stats", cl::desc("Print translation statistics"),
cl::value_desc("print-stats"), cl::cat(ToolTemplateCategory));
cl::opt<bool> PrintStats("print-stats",
cl::desc("Print translation statistics"),
cl::value_desc("print-stats"),
cl::cat(ToolTemplateCategory));
cl::opt<std::string> OutputStatsFilename("o-stats", cl::desc("Output filename for statistics"),
cl::value_desc("filename"), cl::cat(ToolTemplateCategory));
cl::opt<std::string> OutputStatsFilename("o-stats",
cl::desc("Output filename for statistics"),
cl::value_desc("filename"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> Examine("examine", cl::desc("Combines -no-output and -print-stats options"),
cl::value_desc("examine"), cl::cat(ToolTemplateCategory));
cl::opt<bool> Examine("examine",
cl::desc("Combines -no-output and -print-stats options"),
cl::value_desc("examine"),
cl::cat(ToolTemplateCategory));
cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage);
hipify-clang/src/CUDA2HipMap.cpp
+1993 -3678
Vedi File
File diff soppresso perché troppo grande Carica Diff
hipify-clang/src/HipifyAction.cpp
+55 -79
Vedi File
@@ -25,8 +25,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
const auto found = CUDA_RENAMES_MAP().find(name);
if (found != CUDA_RENAMES_MAP().end()) {
StringRef repName = found->second.hipName;
hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL,
ApiTypes::API_RUNTIME, found->second.unsupported};
hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported};
Statistics::current().incrementCounter(counter, name.str());
if (!counter.unsupported) {
@@ -78,8 +77,7 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
if (found->second.unsupported) {
// An unsupported identifier? Curses! Warn the user.
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning,
"CUDA identifier unsupported in hip");
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier unsupported in hip");
DE.Report(sl, ID);
return;
}
@@ -96,10 +94,8 @@ clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRan
clang::SourceLocation begin = exprRange.getBegin();
clang::SourceLocation end = exprRange.getEnd();
bool beginSafe = !SM.isMacroBodyExpansion(begin) ||
clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{});
bool endSafe = !SM.isMacroBodyExpansion(end) ||
clang::Lexer::isAtEndOfMacroExpansion(end, SM, clang::LangOptions{});
bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{});
bool endSafe = !SM.isMacroBodyExpansion(end) || clang::Lexer::isAtEndOfMacroExpansion(end, SM, clang::LangOptions{});
if (beginSafe && endSafe) {
return {SM.getFileLoc(begin), SM.getFileLoc(end)};
@@ -124,9 +120,7 @@ clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRa
StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
return clang::Lexer::getSourceText(
clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM,
clang::LangOptions(), nullptr);
return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr);
}
/**
@@ -141,36 +135,28 @@ std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Exp
}
}
} // anonymous namespace
} // anonymous namespace
bool HipifyAction::Exclude(const hipCounter& hipToken) {
bool HipifyAction::Exclude(const hipCounter & hipToken) {
switch (hipToken.type) {
case CONV_INCLUDE_CUDA_MAIN_H:
switch (hipToken.apiType) {
case API_DRIVER:
case API_RUNTIME:
if (insertedRuntimeHeader) {
return true;
}
if (insertedRuntimeHeader) { return true; }
insertedRuntimeHeader = true;
return false;
case API_BLAS:
if (insertedBLASHeader) {
return true;
}
if (insertedBLASHeader) { return true; }
insertedBLASHeader = true;
return false;
case API_RAND:
if (hipToken.hipName == "hiprand_kernel.h") {
if (insertedRAND_kernelHeader) {
return true;
}
if (insertedRAND_kernelHeader) { return true; }
insertedRAND_kernelHeader = true;
return false;
} else if (hipToken.hipName == "hiprand.h") {
if (insertedRANDHeader) {
return true;
}
if (insertedRANDHeader) { return true; }
insertedRANDHeader = true;
return false;
}
@@ -181,9 +167,7 @@ bool HipifyAction::Exclude(const hipCounter& hipToken) {
case CONV_INCLUDE:
switch (hipToken.apiType) {
case API_RAND:
if (insertedRAND_kernelHeader) {
return true;
}
if (insertedRAND_kernelHeader) { return true; }
insertedRAND_kernelHeader = true;
return false;
default:
@@ -196,11 +180,13 @@ bool HipifyAction::Exclude(const hipCounter& hipToken) {
return false;
}
void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clang::Token&,
StringRef file_name, bool is_angled,
void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
const clang::Token&,
StringRef file_name,
bool is_angled,
clang::CharSourceRange filename_range,
const clang::FileEntry*, StringRef, StringRef,
const clang::Module*) {
const clang::FileEntry*, StringRef,
StringRef, const clang::Module*) {
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
if (!SM.isWrittenInMainFile(hash_loc)) {
return;
@@ -222,8 +208,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clan
clang::SourceLocation sl = filename_range.getBegin();
if (found->second.unsupported) {
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
DE.Report(sl,
DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header"));
DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header"));
return;
}
@@ -235,22 +220,19 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clan
if (is_angled) {
newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer);
} else {
newInclude =
llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer);
newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer);
}
} else {
// hashLoc is location of the '#', thus replacing the whole include directive by empty
// newInclude starting with '#'.
// hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'.
sl = hash_loc;
}
const char* B = SM.getCharacterData(sl);
const char* E = SM.getCharacterData(filename_range.getEnd());
const char *B = SM.getCharacterData(sl);
const char *E = SM.getCharacterData(filename_range.getEnd());
ct::Replacement Rep(SM, sl, E - B, newInclude);
insertReplacement(Rep, clang::FullSourceLoc{sl, SM});
}
void HipifyAction::PragmaDirective(clang::SourceLocation Loc,
clang::PragmaIntroducerKind Introducer) {
void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) {
if (pragmaOnce) {
return;
}
@@ -284,8 +266,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
const clang::Expr& calleeExpr = *(launchKernel->getCallee());
OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", ";
// Next up are the four kernel configuration parameters, the last two of which are optional and
// default to zero.
// Next up are the four kernel configuration parameters, the last two of which are optional and default to zero.
const clang::CallExpr& config = *(launchKernel->getConfig());
// Copy the two dimensional arguments verbatim.
@@ -312,14 +293,11 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
OS << ")";
clang::SourceRange replacementRange =
getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()});
clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()});
clang::SourceLocation launchStart = replacementRange.getBegin();
clang::SourceLocation launchEnd = replacementRange.getEnd();
size_t length = SM->getCharacterData(
clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) -
SM->getCharacterData(launchStart);
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart);
ct::Replacement Rep(*SM, launchStart, length, OS.str());
clang::FullSourceLoc fullSL(launchStart, *SM);
@@ -330,8 +308,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
return true;
}
bool HipifyAction::cudaSharedIncompleteArrayVar(
const clang::ast_matchers::MatchFinder::MatchResult& Result) {
bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
StringRef refName = "cudaSharedIncompleteArrayVar";
auto* sharedVar = Result.Nodes.getNodeAs<clang::VarDecl>(refName);
if (!sharedVar) {
@@ -379,8 +356,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(
return true;
}
void HipifyAction::insertReplacement(const ct::Replacement& rep,
const clang::FullSourceLoc& fullSL) {
void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) {
llcompat::insertReplacement(*replacements, rep);
if (PrintStats) {
rep.getLength();
@@ -389,19 +365,22 @@ void HipifyAction::insertReplacement(const ct::Replacement& rep,
}
}
std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI,
llvm::StringRef) {
std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) {
Finder.reset(new clang::ast_matchers::MatchFinder);
// Replace the <<<...>>> language extension with a hip kernel launch
Finder->addMatcher(
mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
Finder->addMatcher(mat::varDecl(mat::isExpansionInMainFile(),
mat::allOf(mat::hasAttr(clang::attr::CUDAShared),
mat::hasType(mat::incompleteArrayType())))
.bind("cudaSharedIncompleteArrayVar"),
this);
Finder->addMatcher(
mat::varDecl(
mat::isExpansionInMainFile(),
mat::allOf(
mat::hasAttr(clang::attr::CUDAShared),
mat::hasType(mat::incompleteArrayType())
)
).bind("cudaSharedIncompleteArrayVar"),
this
);
// Ownership is transferred to the caller...
return Finder->newASTConsumer();
@@ -410,9 +389,9 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
void HipifyAction::EndSourceFileAction() {
// Insert the hip header, if we didn't already do it by accident during substitution.
if (!insertedRuntimeHeader) {
// It's not sufficient to just replace CUDA headers with hip ones, because numerous CUDA
// headers are implicitly included by the compiler. Instead, we _delete_ CUDA headers, and
// unconditionally insert one copy of the hip include into every file.
// It's not sufficient to just replace CUDA headers with hip ones, because numerous CUDA headers are
// implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert
// one copy of the hip include into every file.
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
clang::SourceLocation sl;
if (pragmaOnce) {
@@ -439,25 +418,22 @@ namespace {
class PPCallbackProxy : public clang::PPCallbacks {
HipifyAction& hipifyAction;
public:
explicit PPCallbackProxy(HipifyAction& action) : hipifyAction(action) {}
public:
explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {}
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,
StringRef file_name, bool is_angled, clang::CharSourceRange filename_range,
const clang::FileEntry* file, StringRef search_path, StringRef relative_path,
const clang::Module* imported) override {
hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled,
filename_range, file, search_path, relative_path, imported);
hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported);
}
void PragmaDirective(clang::SourceLocation Loc,
clang::PragmaIntroducerKind Introducer) override {
void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override {
hipifyAction.PragmaDirective(Loc, Introducer);
}
};
} // namespace
}
void HipifyAction::ExecuteAction() {
clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
@@ -468,10 +444,10 @@ void HipifyAction::ExecuteAction() {
clang::Lexer RawLex(SM.getMainFileID(), FromFile, SM, PP.getLangOpts());
RawLex.SetKeepWhitespaceMode(true);
// Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us
// enough information to tell the difference between identifiers, string literals, and "other
// stuff". It also ignores preprocessor directives, so this transformation will operate inside
// preprocessor-deleted code.
// Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us enough
// information to tell the difference between identifiers, string literals, and "other stuff". It also
// ignores preprocessor directives, so this transformation will operate inside preprocessor-deleted
// code.
clang::Token RawTok;
RawLex.LexFromRawLexer(RawTok);
while (RawTok.isNot(clang::tok::eof)) {
hipify-clang/src/HipifyAction.h
+27 -25
Vedi File
@@ -15,15 +15,14 @@ namespace ct = clang::tooling;
*/
class HipifyAction : public clang::ASTFrontendAction,
public clang::ast_matchers::MatchFinder::MatchCallback {
private:
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.
/// 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 insertedBLASHeader = false;
bool insertedRANDHeader = false;
@@ -41,11 +40,12 @@ class HipifyAction : public clang::ASTFrontendAction,
/**
* Replace a CUDA identifier with the corresponding hip identifier, if applicable.
*/
void RewriteToken(const clang::Token& t);
void RewriteToken(const clang::Token &t);
public:
explicit HipifyAction(ct::Replacements* replacements)
: clang::ASTFrontendAction(), replacements(replacements) {}
public:
explicit HipifyAction(ct::Replacements *replacements):
clang::ASTFrontendAction(),
replacements(replacements) {}
// MatchCallback listeners
bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result);
@@ -55,21 +55,24 @@ class HipifyAction : public clang::ASTFrontendAction,
/**
* 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);
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.
*/
* Called by the preprocessor for each pragma directive during the non-raw lexing pass.
*/
void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer);
protected:
protected:
/**
* Add a Replacement for the current file. These will all be applied after executing the
* FrontendAction.
* 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);
@@ -84,13 +87,12 @@ class HipifyAction : public clang::ASTFrontendAction,
void EndSourceFileAction() override;
/**
* MatchCallback API entry point. Called by the AST visitor while searching the AST for things
* we registered an interest for.
* 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;
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
bool Exclude(const hipCounter& hipToken);
bool Exclude(const hipCounter & hipToken);
};
hipify-clang/src/LLVMCompat.cpp
+2 -3
Vedi File
@@ -32,8 +32,7 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re
#endif
}
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token* start, size_t len,
bool DisableMacroExpansion) {
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
_pp.EnterTokenStream(start, len, false, DisableMacroExpansion);
#else
@@ -41,4 +40,4 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token*
#endif
}
} // namespace llcompat
} // namespace llcompat
hipify-clang/src/LLVMCompat.h
+6 -4
Vedi File
@@ -18,9 +18,9 @@ namespace llcompat {
* remain unchanged, so let's be slightly ugly about it here. :D
*/
#if LLVM_VERSION_MAJOR > 4
#define GET_NUM_ARGS() getNumParams()
#define GET_NUM_ARGS() getNumParams()
#else
#define GET_NUM_ARGS() getNumArgs()
#define GET_NUM_ARGS() getNumArgs()
#endif
void PrintStackTraceOnErrorSignal();
@@ -41,7 +41,9 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re
/**
* Version-agnostic version of Preprocessor::EnterTokenStream().
*/
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token* start, size_t len,
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
const clang::Token *start,
size_t len,
bool DisableMacroExpansion);
} // namespace llcompat
} // namespace llcompat
hipify-clang/src/ReplacementsFrontendActionFactory.h
+7 -4
Vedi File
@@ -17,9 +17,12 @@ template <typename T>
class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory {
ct::Replacements* replacements;
public:
explicit ReplacementsFrontendActionFactory(ct::Replacements* r)
: ct::FrontendActionFactory(), replacements(r) {}
public:
explicit ReplacementsFrontendActionFactory(ct::Replacements* r):
ct::FrontendActionFactory(),
replacements(r) {}
clang::FrontendAction* create() override { return new T(replacements); }
clang::FrontendAction* create() override {
return new T(replacements);
}
};
hipify-clang/src/Statistics.cpp
+69 -72
Vedi File
@@ -4,26 +4,26 @@
#include <iomanip>
const char* counterNames[NUM_CONV_TYPES] = {"version", "init", "device",
"mem", "kern", "coord_func",
"math_func", "device_func", "special_func",
"stream", "event", "occupancy",
"ctx", "peer", "module",
"cache", "exec", "err",
"def", "tex", "gl",
"graphics", "surface", "jit",
"d3d9", "d3d10", "d3d11",
"vdpau", "egl", "thread",
"other", "include", "include_cuda_main_header",
"type", "literal", "numeric_literal"};
const char *counterNames[NUM_CONV_TYPES] = {
"version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func",
"special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
"cache", "exec", "err", "def", "tex", "gl", "graphics",
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
"thread", "other", "include", "include_cuda_main_header", "type", "literal",
"numeric_literal"
};
const char* apiNames[NUM_API_TYPES] = {"CUDA Driver API", "CUDA RT API", "CUBLAS API",
"CURAND API"};
const char *apiNames[NUM_API_TYPES] = {
"CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API"
};
namespace {
template <typename ST, typename ST2>
void conditionalPrint(ST* stream1, ST2* stream2, const std::string& s1, const std::string& s2) {
template<typename ST, typename ST2>
void conditionalPrint(ST *stream1,
ST2* stream2,
const std::string& s1,
const std::string& s2) {
if (stream1) {
*stream1 << s1;
}
@@ -37,8 +37,8 @@ void conditionalPrint(ST* stream1, ST2* stream2, const std::string& s1, const st
/**
* Print a named stat value to both the terminal and the CSV file.
*/
template <typename T>
void printStat(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& name, T value) {
template<typename T>
void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string &name, T value) {
if (printOut) {
*printOut << " " << name << ": " << value << "\n";
}
@@ -49,12 +49,12 @@ void printStat(std::ostream* csv, llvm::raw_ostream* printOut, const std::string
}
} // Anonymous namespace
} // Anonymous namespace
void StatCounter::incrementCounter(const hipCounter& counter, std::string name) {
counters[name]++;
apiCounters[(int)counter.apiType]++;
convTypeCounters[(int)counter.type]++;
apiCounters[(int) counter.apiType]++;
convTypeCounters[(int) counter.type]++;
}
void StatCounter::add(const StatCounter& other) {
@@ -81,36 +81,32 @@ int StatCounter::getConvSum() {
}
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) {
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n",
"[HIPIFY] info: " + prefix + " refs by type:\n");
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (convTypeCounters[i] > 0) {
printStat(csv, printOut, counterNames[i], convTypeCounters[i]);
}
}
conditionalPrint(csv, printOut, "\nCUDA API;Count\n",
"[HIPIFY] info: " + prefix + " refs by API:\n");
conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n");
for (int i = 0; i < NUM_API_TYPES; i++) {
printStat(csv, printOut, apiNames[i], apiCounters[i]);
}
conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n",
"[HIPIFY] info: " + prefix + " refs by names:\n");
for (const auto& it : counters) {
conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n", "[HIPIFY] info: " + prefix + " refs by names:\n");
for (const auto &it : counters) {
printStat(csv, printOut, it.first, it.second);
}
}
Statistics::Statistics(std::string name) : fileName(name) {
Statistics::Statistics(std::string name): fileName(name) {
// Compute the total bytes/lines in the input file.
std::ifstream src_file(name, std::ios::binary | std::ios::ate);
src_file.clear();
src_file.seekg(0);
totalLines = (int)std::count(std::istreambuf_iterator<char>(src_file),
std::istreambuf_iterator<char>(), '\n');
totalBytes = (int)src_file.tellg();
totalLines = (int) std::count(std::istreambuf_iterator<char>(src_file), std::istreambuf_iterator<char>(), '\n');
totalBytes = (int) src_file.tellg();
// Mark the start time...
startTime = chr::steady_clock::now();
@@ -119,7 +115,7 @@ Statistics::Statistics(std::string name) : fileName(name) {
///////// Counter update routines //////////
void Statistics::incrementCounter(const hipCounter& counter, std::string name) {
void Statistics::incrementCounter(const hipCounter &counter, std::string name) {
if (counter.unsupported) {
unsupported.incrementCounter(counter, name);
} else {
@@ -127,7 +123,7 @@ void Statistics::incrementCounter(const hipCounter& counter, std::string name) {
}
}
void Statistics::add(const Statistics& other) {
void Statistics::add(const Statistics &other) {
supported.add(other.supported);
unsupported.add(other.unsupported);
totalBytes += other.totalBytes;
@@ -135,60 +131,61 @@ void Statistics::add(const Statistics& other) {
touchedBytes += other.touchedBytes;
}
void Statistics::lineTouched(int lineNumber) { touchedLines.insert(lineNumber); }
void Statistics::bytesChanged(int bytes) { touchedBytes += bytes; }
void Statistics::markCompletion() { completionTime = chr::steady_clock::now(); }
void Statistics::lineTouched(int lineNumber) {
touchedLines.insert(lineNumber);
}
void Statistics::bytesChanged(int bytes) {
touchedBytes += bytes;
}
void Statistics::markCompletion() {
completionTime = chr::steady_clock::now();
}
///////// Output functions //////////
void Statistics::print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader) {
if (!skipHeader) {
std::string str = "file \'" + fileName + "\' statistics:\n";
conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str);
}
if (!skipHeader) {
std::string str = "file \'" + fileName + "\' statistics:\n";
conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str);
}
size_t changedLines = touchedLines.size();
size_t changedLines = touchedLines.size();
// Total number of (un)supported refs that were converted.
int supportedSum = supported.getConvSum();
int unsupportedSum = unsupported.getConvSum();
// Total number of (un)supported refs that were converted.
int supportedSum = supported.getConvSum();
int unsupportedSum = unsupported.getConvSum();
printStat(csv, printOut, "CONVERTED refs count", supportedSum);
printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum);
printStat(
csv, printOut, "CONVERSION %",
100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum)));
printStat(csv, printOut, "REPLACED bytes", touchedBytes);
printStat(csv, printOut, "TOTAL bytes", totalBytes);
printStat(csv, printOut, "CHANGED lines of code", changedLines);
printStat(csv, printOut, "TOTAL lines of code", totalLines);
printStat(csv, printOut, "CONVERTED refs count", supportedSum);
printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum);
printStat(csv, printOut, "CONVERSION %", 100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum)));
printStat(csv, printOut, "REPLACED bytes", touchedBytes);
printStat(csv, printOut, "TOTAL bytes", totalBytes);
printStat(csv, printOut, "CHANGED lines of code", changedLines);
printStat(csv, printOut, "TOTAL lines of code", totalLines);
if (totalBytes > 0) {
printStat(csv, printOut, "CODE CHANGED (in bytes) %",
std::lround(double(touchedBytes * 100) / double(totalBytes)));
}
if (totalBytes > 0) {
printStat(csv, printOut, "CODE CHANGED (in bytes) %", std::lround(double(touchedBytes * 100) / double(totalBytes)));
}
if (totalLines > 0) {
printStat(csv, printOut, "CODE CHANGED (in lines) %",
std::lround(double(changedLines * 100) / double(totalLines)));
}
if (totalLines > 0) {
printStat(csv, printOut, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
}
typedef std::chrono::duration<double, std::milli> duration;
duration elapsed = completionTime - startTime;
std::stringstream stream;
stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
printStat(csv, printOut, "TIME ELAPSED s", stream.str());
typedef std::chrono::duration<double, std::milli> duration;
duration elapsed = completionTime - startTime;
std::stringstream stream;
stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
printStat(csv, printOut, "TIME ELAPSED s", stream.str());
supported.print(csv, printOut, "CONVERTED");
unsupported.print(csv, printOut, "UNCONVERTED");
supported.print(csv, printOut, "CONVERTED");
unsupported.print(csv, printOut, "UNCONVERTED");
}
void Statistics::printAggregate(std::ostream* csv, llvm::raw_ostream* printOut) {
void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) {
Statistics globalStats = getAggregate();
conditionalPrint(csv, printOut, "\nTOTAL statistics:\n",
"\n[HIPIFY] info: TOTAL statistics:\n");
conditionalPrint(csv, printOut, "\nTOTAL statistics:\n", "\n[HIPIFY] info: TOTAL statistics:\n");
// A file is considered "converted" if we made any changes to it.
int convertedFiles = 0;
hipify-clang/src/Statistics.h
+26 -20
Vedi File
@@ -49,14 +49,20 @@ enum ConvTypes {
CONV_NUMERIC_LITERAL,
CONV_LAST
};
constexpr int NUM_CONV_TYPES = (int)ConvTypes::CONV_LAST;
constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;
enum ApiTypes { API_DRIVER = 0, API_RUNTIME, API_BLAS, API_RAND, API_LAST };
constexpr int NUM_API_TYPES = (int)ApiTypes::API_LAST;
enum ApiTypes {
API_DRIVER = 0,
API_RUNTIME,
API_BLAS,
API_RAND,
API_LAST
};
constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
// The names of various fields in in the statistics reports.
extern const char* counterNames[NUM_CONV_TYPES];
extern const char* apiNames[NUM_API_TYPES];
extern const char *counterNames[NUM_CONV_TYPES];
extern const char *apiNames[NUM_API_TYPES];
struct hipCounter {
@@ -71,14 +77,14 @@ struct hipCounter {
* Tracks a set of named counters, as well as counters for each of the type enums defined above.
*/
class StatCounter {
private:
private:
// Each thing we track is either "supported" or "unsupported"...
std::map<std::string, int> counters;
int apiCounters[NUM_API_TYPES] = {};
int convTypeCounters[NUM_CONV_TYPES] = {};
public:
public:
void incrementCounter(const hipCounter& counter, std::string name);
/**
@@ -109,15 +115,15 @@ class Statistics {
chr::steady_clock::time_point startTime;
chr::steady_clock::time_point completionTime;
public:
public:
Statistics(std::string name);
void incrementCounter(const hipCounter& counter, std::string name);
void incrementCounter(const hipCounter &counter, std::string name);
/**
* Add the counters from `other` onto the counters of this object.
*/
void add(const Statistics& other);
void add(const Statistics &other);
void lineTouched(int lineNumber);
void bytesChanged(int bytes);
@@ -129,18 +135,18 @@ class Statistics {
/////// Output functions ///////
public:
/**
public:
/**
* Pretty-print the statistics stored in this object.
*
* @param csv Pointer to an output stream for the CSV to write. If null, no CSV is written
* @param printOut Pointer to an output stream to print human-readable textual stats to. If
* null, no such stats are produced.
* @param printOut Pointer to an output stream to print human-readable textual stats to. If null, no
* such stats are produced.
*/
void print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader = false);
/// Print aggregated statistics for all registered counters.
static void printAggregate(std::ostream* csv, llvm::raw_ostream* printOut);
static void printAggregate(std::ostream *csv, llvm::raw_ostream* printOut);
/////// Static nonsense ///////
@@ -156,15 +162,15 @@ class Statistics {
static Statistics getAggregate();
/**
* Convenient global entry point for updating the "active" Statistics. Since we operate
* single-threadedly processing one file at a time, this allows us to simply expose the stats
* for the current file globally, simplifying things.
* Convenient global entry point for updating the "active" Statistics. Since we operate single-threadedly
* processing one file at a time, this allows us to simply expose the stats for the current file globally,
* simplifying things.
*/
static Statistics& current();
/**
* Set the active Statistics object to the named one, creating it if necessary, and write the
* completion timestamp into the currently active one.
* Set the active Statistics object to the named one, creating it if necessary, and write the completion
* timestamp into the currently active one.
*/
static void setActive(std::string name);
};
hipify-clang/src/StringUtils.cpp
+1 -1
Vedi File
@@ -8,7 +8,7 @@ llvm::StringRef unquoteStr(llvm::StringRef s) {
return s;
}
void removePrefixIfPresent(std::string& s, std::string prefix) {
void removePrefixIfPresent(std::string &s, std::string prefix) {
if (s.find(prefix) != 0) {
return;
}
hipify-clang/src/StringUtils.h
+1 -1
Vedi File
@@ -11,4 +11,4 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
/**
* If `s` starts with `prefix`, remove it. Otherwise, does nothing.
*/
void removePrefixIfPresent(std::string& s, std::string prefix);
void removePrefixIfPresent(std::string &s, std::string prefix);
hipify-clang/src/main.cpp
+88 -93
Vedi File
@@ -46,116 +46,111 @@ namespace ct = clang::tooling;
namespace {
void copyFile(const std::string& src, const std::string& dst) {
std::ifstream source(src, std::ios::binary);
std::ofstream dest(dst, std::ios::binary);
dest << source.rdbuf();
std::ifstream source(src, std::ios::binary);
std::ofstream dest(dst, std::ios::binary);
dest << source.rdbuf();
}
} // anonymous namespace
} // anonymous namespace
int main(int argc, const char** argv) {
llcompat::PrintStackTraceOnErrorSignal();
int main(int argc, const char **argv) {
llcompat::PrintStackTraceOnErrorSignal();
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
std::string dst = OutputFilename;
if (!dst.empty() && fileSources.size() > 1) {
llvm::errs() << "[HIPIFY] conflict: -o and multiple source files are specified.\n";
return 1;
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
std::string dst = OutputFilename;
if (!dst.empty() && fileSources.size() > 1) {
llvm::errs() << "[HIPIFY] conflict: -o and multiple source files are specified.\n";
return 1;
}
if (NoOutput) {
if (Inplace) {
llvm::errs() << "[HIPIFY] conflict: both -no-output and -inplace options are specified.\n";
return 1;
}
if (!dst.empty()) {
llvm::errs() << "[HIPIFY] conflict: both -no-output and -o options are specified.\n";
return 1;
}
}
if (Examine) {
NoOutput = PrintStats = true;
}
int Result = 0;
// Arguments for the Statistics print routines.
std::unique_ptr<std::ostream> csv = nullptr;
llvm::raw_ostream* statPrint = nullptr;
if (!OutputStatsFilename.empty()) {
csv = std::unique_ptr<std::ostream>(new std::ofstream(OutputStatsFilename, std::ios_base::trunc));
}
if (PrintStats) {
statPrint = &llvm::errs();
}
for (const auto & src : fileSources) {
if (dst.empty()) {
if (Inplace) {
dst = src;
} else {
dst = src + ".hip";
}
} else if (Inplace) {
llvm::errs() << "[HIPIFY] conflict: both -o and -inplace options are specified.\n";
return 1;
}
if (NoOutput) {
if (Inplace) {
llvm::errs()
<< "[HIPIFY] conflict: both -no-output and -inplace options are specified.\n";
return 1;
}
if (!dst.empty()) {
llvm::errs() << "[HIPIFY] conflict: both -no-output and -o options are specified.\n";
return 1;
}
}
std::string tmpFile = src + ".hipify-tmp";
if (Examine) {
NoOutput = PrintStats = true;
}
// Create a copy of the file to work on. When we're done, we'll move this onto the
// output (which may mean overwriting the input, if we're in-place).
// Should we fail for some reason, we'll just leak this file and not corrupt the input.
copyFile(src, tmpFile);
int Result = 0;
// Initialise the statistics counters for this file.
Statistics::setActive(src);
// Arguments for the Statistics print routines.
std::unique_ptr<std::ostream> csv = nullptr;
llvm::raw_ostream* statPrint = nullptr;
if (!OutputStatsFilename.empty()) {
csv = std::unique_ptr<std::ostream>(
new std::ofstream(OutputStatsFilename, std::ios_base::trunc));
}
if (PrintStats) {
statPrint = &llvm::errs();
}
// RefactoringTool operates on the file in-place. Giving it the output path is no good,
// because that'll break relative includes, and we don't want to overwrite the input file.
// So what we do is operate on a copy, which we then move to the output.
ct::RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile);
ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile);
for (const auto& src : fileSources) {
if (dst.empty()) {
if (Inplace) {
dst = src;
} else {
dst = src + ".hip";
}
} else if (Inplace) {
llvm::errs() << "[HIPIFY] conflict: both -o and -inplace options are specified.\n";
return 1;
}
ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
std::string tmpFile = src + ".hipify-tmp";
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
// Create a copy of the file to work on. When we're done, we'll move this onto the
// output (which may mean overwriting the input, if we're in-place).
// Should we fail for some reason, we'll just leak this file and not corrupt the input.
copyFile(src, tmpFile);
// Initialise the statistics counters for this file.
Statistics::setActive(src);
// RefactoringTool operates on the file in-place. Giving it the output path is no good,
// because that'll break relative includes, and we don't want to overwrite the input file.
// So what we do is operate on a copy, which we then move to the output.
ct::RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile);
ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile);
ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
Tool.appendArgumentsAdjuster(
ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
// Ensure at least c++11 is used.
Tool.appendArgumentsAdjuster(
ct::getInsertArgumentAdjuster("-std=c++11", ct::ArgumentInsertPosition::BEGIN));
// Ensure at least c++11 is used.
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-std=c++11", ct::ArgumentInsertPosition::BEGIN));
#if defined(HIPIFY_CLANG_RES)
Tool.appendArgumentsAdjuster(
ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
#endif
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
// Hipify _all_ the things!
if (Tool.runAndSave(&actionFactory)) {
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
}
// Either move the tmpfile to the output, or remove it.
if (!NoOutput) {
rename(tmpFile.c_str(), dst.c_str());
} else {
remove(tmpFile.c_str());
}
Statistics::current().markCompletion();
Statistics::current().print(csv.get(), statPrint);
dst.clear();
// Hipify _all_ the things!
if (Tool.runAndSave(&actionFactory)) {
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
}
if (fileSources.size() > 1) {
Statistics::printAggregate(csv.get(), statPrint);
// Either move the tmpfile to the output, or remove it.
if (!NoOutput) {
rename(tmpFile.c_str(), dst.c_str());
} else {
remove(tmpFile.c_str());
}
return Result;
Statistics::current().markCompletion();
Statistics::current().print(csv.get(), statPrint);
dst.clear();
}
if (fileSources.size() > 1) {
Statistics::printAggregate(csv.get(), statPrint);
}
return Result;
}
include/hip/hcc_detail/hip_runtime_api.h
+21
Vedi File
@@ -50,6 +50,13 @@ THE SOFTWARE.
#define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*)0x02)
#define HIP_LAUNCH_PARAM_END ((void*)0x03)
#ifdef __cplusplus
#define __dparm(x) \
= x
#else
#define __dparm(x)
#endif
// Structure definitions:
#ifdef __cplusplus
extern "C" {
@@ -1436,6 +1443,20 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height);
/**
* @brief Fills asynchronously the memory area pointed to by dst with the constant value.
*
* @param[in] dst Pointer to device memory
* @param[in] pitch - data size in bytes
* @param[in] value - constant value to be set
* @param[in] width
* @param[in] height
* @param[in] stream
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree
*/
hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height,hipStream_t stream __dparm(0));
/**
* @brief Query memory info.
* Return snapshot of free memory, and total allocatable memory on the device.
include/hip/nvcc_detail/hip_runtime_api.h
+8
Vedi File
@@ -625,6 +625,14 @@ inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, s
return hipCUResultTohipError(cuMemsetD8(dest, value, sizeBytes));
}
inline static hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
return hipCUDAErrorTohipError(cudaMemset2D(dst, pitch, value, width, height));
}
inline static hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream __dparm(0)) {
return hipCUDAErrorTohipError(cudaMemset2DAsync(dst, pitch, value, width, height, stream));
}
inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int device) {
struct cudaDeviceProp cdprop;
cudaError_t cerror;
src/hip_memory.cpp
+19 -1
Vedi File
@@ -1552,7 +1552,7 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea
{
hipError_t e = hipSuccess;
if (stream) {
if (stream && (dst != NULL)) {
if(copyDataType == ihipMemsetDataTypeChar){
if ((sizeBytes & 0x3) == 0) {
// use a faster dword-per-workitem copy:
@@ -1647,6 +1647,24 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
return ihipLogStatus(e);
}
hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream )
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height, stream);
hipError_t e = hipSuccess;
stream = ihipSyncAndResolveStream(stream);
if (stream) {
size_t sizeBytes = pitch * height;
e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
} else {
e = hipErrorInvalidValue;
}
return ihipLogStatus(e);
};
hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
tests/src/runtimeApi/memory/hipMemset2D.cpp
+119
Vedi File
@@ -0,0 +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.
*/
// Simple test for memset.
// Also serves as a template for other tests.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
#include "test_common.h"
bool testhipMemset2D(int memsetval,int p_gpuDevice)
{
size_t numH = 256;
size_t numW = 256;
size_t pitch_A;
size_t width = numW * sizeof(char);
size_t sizeElements = width * numH;
size_t elements = numW* numH;
printf ("testhipMemset2D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
char *A_d;
char *A_h;
bool testResult = true;
HIPCHECK ( hipMallocPitch((void**)&A_d, &pitch_A, width , numH) );
A_h = (char*)malloc(sizeElements);
HIPASSERT(A_h != NULL);
for (size_t i=0; i<elements; i++) {
A_h[i] = 1;
}
HIPCHECK ( hipMemset2D(A_d, pitch_A, memsetval, numW, numH) );
HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
for (int i=0; i<elements; i++) {
if (A_h[i] != memsetval) {
testResult = false;
printf("testhipMemset2D mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
break;
}
}
hipFree(A_d);
free(A_h);
return testResult;
}
bool testhipMemset2DAsync(int memsetval,int p_gpuDevice)
{
size_t numH = 256;
size_t numW = 256;
size_t pitch_A;
size_t width = numW * sizeof(char);
size_t sizeElements = width * numH;
size_t elements = numW* numH;
printf ("testhipMemset2DAsync memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
char *A_d;
char *A_h;
bool testResult = true;
HIPCHECK ( hipMallocPitch((void**)&A_d, &pitch_A, width , numH) );
A_h = (char*)malloc(sizeElements);
HIPASSERT(A_h != NULL);
for (size_t i=0; i<elements; i++) {
A_h[i] = 1;
}
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK ( hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
for (int i=0; i<elements; i++) {
if (A_h[i] != memsetval) {
testResult = false;
printf("testhipMemset2DAsync mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
break;
}
}
hipFree(A_d);
HIPCHECK(hipStreamDestroy(stream));
free(A_h);
return testResult;
}
int main(int argc, char *argv[])
{
HipTest::parseStandardArguments(argc, argv, true);
bool testResult = false;
HIPCHECK(hipSetDevice(p_gpuDevice));
testResult = testhipMemset2D(memsetval, p_gpuDevice);
testResult = testhipMemset2DAsync(memsetval, p_gpuDevice);
passed();
}