Skip to content

Commit 3e7f712

Browse files
zhiweij1zhiweij1
committed
Merge remote-tracking branch 'origin/SYCLomatic' into fix_pd
2 parents 54bdeba + 9b620bf commit 3e7f712

27 files changed

Lines changed: 639 additions & 479 deletions

‎clang/lib/DPCT/AnalysisInfo.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6667,8 +6667,19 @@ KernelPrinter &KernelCallExpr::SubmitStmtsList::print(KernelPrinter &Printer) {
66676667
Printer.line("cgh.depends_on(dpct::get_current_device().get_in_order_"
66686668
"queues_last_events());");
66696669
} else {
6670+
Printer.line("#ifdef __INTEL_LLVM_COMPILER");
6671+
Printer.newLine();
66706672
Printer.line("cgh.depends_on(dpct::get_default_queue().ext_oneapi_get_"
66716673
"last_event());");
6674+
Printer.newLine();
6675+
Printer.line("#else");
6676+
Printer.newLine();
6677+
Printer.line("auto e_opt = dpct::get_default_queue().ext_oneapi_get_last_"
6678+
"event();");
6679+
Printer.newLine();
6680+
Printer.line("if (e_opt) cgh.depends_on(*e_opt);");
6681+
Printer.newLine();
6682+
Printer.line("#endif");
66726683
}
66736684
Printer.newLine();
66746685
}

‎clang/lib/DPCT/RuleInfra/MapNames.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -338,6 +338,9 @@ void MapNames::setExplicitNamespaceMap(
338338
{"cudaStream_t",
339339
std::make_shared<TypeNameRule>(getDpctNamespace() + "queue_ptr",
340340
HelperFeatureEnum::device_ext)},
341+
{"cudaHostFn_t",
342+
std::make_shared<TypeNameRule>(getDpctNamespace() + "host_func",
343+
HelperFeatureEnum::device_ext)},
341344
{"CUstream",
342345
std::make_shared<TypeNameRule>(getDpctNamespace() + "queue_ptr",
343346
HelperFeatureEnum::device_ext)},

‎clang/lib/DPCT/RulesLang/RulesLang.cpp‎

Lines changed: 82 additions & 42 deletions
Original file line numberDiff line numberDiff line change
@@ -346,7 +346,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
346346
"cublasLtMatmulHeuristicResult_t", "CUjit_target",
347347
"cublasLtMatrixTransformDesc_t", "cudaGraphicsMapFlags",
348348
"cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
349-
"CUstreamCallback"))))))
349+
"CUstreamCallback", "cudaHostFn_t"))))))
350350
.bind("cudaTypeDef"),
351351
this);
352352

@@ -4367,11 +4367,12 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) {
43674367
StreamName = "{{NEEDREPLACEQ" + std::to_string(Index) + "}}.";
43684368
ReplStr = StreamName + "ext_oneapi_empty()";
43694369
} else {
4370-
StreamName = getStmtSpelling(StreamArg);
4370+
ExprAnalysis EA(StreamArg);
4371+
ReplStr = EA.getReplacedString();
43714372
if (needExtraParensInMemberExpr(StreamArg)) {
4372-
StreamName = "(" + StreamName + ")";
4373+
ReplStr = "(" + ReplStr + ")";
43734374
}
4374-
ReplStr = StreamName + "->" + "ext_oneapi_empty()";
4375+
ReplStr = ReplStr + "->" + "ext_oneapi_empty()";
43754376
}
43764377
if (IsAssigned) {
43774378
ReplStr = MapNames::getCheckErrorMacroName() + "((" + ReplStr + "))";
@@ -4414,7 +4415,12 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) {
44144415

44154416
StmtStr0 = "{{NEEDREPLACEQ" + std::to_string(Index) + "}}.";
44164417
} else {
4417-
StmtStr0 = getStmtSpelling(CE->getArg(0)) + "->";
4418+
ExprAnalysis StreamArgEA(StreamArg);
4419+
StmtStr0 = StreamArgEA.getReplacedString();
4420+
if (needExtraParensInMemberExpr(StreamArg)) {
4421+
StmtStr0 = "(" + StmtStr0 + ")";
4422+
}
4423+
StmtStr0 += "->";
44184424
}
44194425
ReplStr = StmtStr0 + "ext_oneapi_submit_barrier({" +
44204426
StmtStr1 + "})";
@@ -4622,7 +4628,8 @@ void KernelCallRule::registerMatcher(ast_matchers::MatchFinder &MF) {
46224628
this);
46234629

46244630
auto launchAPIName = [&]() {
4625-
return hasAnyName("cudaLaunchKernel", "cudaLaunchCooperativeKernel");
4631+
return hasAnyName("cudaLaunchKernel", "cudaLaunchCooperativeKernel",
4632+
"cudaLaunchHostFunc");
46264633
};
46274634
MF.addMatcher(
46284635
callExpr(allOf(callee(functionDecl(launchAPIName())), parentStmt()))
@@ -4837,56 +4844,89 @@ void KernelCallRule::runRule(
48374844
LaunchKernelCall = getNodeAsType<CallExpr>(Result, "launchUsed");
48384845
IsAssigned = true;
48394846
}
4840-
if (!LaunchKernelCall)
4847+
auto FD = LaunchKernelCall->getDirectCallee();
4848+
if (!LaunchKernelCall || !FD)
48414849
return;
4842-
const Expr *CalleeDRE = LaunchKernelCall->getArg(0);
4843-
bool IsFuncTypeErased = true;
4844-
auto QT = CalleeDRE->getType();
4845-
4846-
if (QT->isPointerType()) {
4847-
QT = QT->getPointeeType();
4848-
}
4849-
if (QT->isFunctionType()) {
4850-
IsFuncTypeErased = false;
4851-
}
4852-
4853-
if (!getAddressedRef(CalleeDRE)) {
4854-
if (IsFuncTypeErased) {
4855-
DpctGlobalInfo::setCVersionCUDALaunchUsed();
4850+
std::string FuncName = FD->getNameAsString();
4851+
std::cout << FuncName << std::endl;
4852+
if (FuncName == "cudaLaunchHostFunc") {
4853+
if (DpctGlobalInfo::getUsmLevel() != UsmLevel::UL_Restricted) {
4854+
report(LaunchKernelCall->getBeginLoc(), Diagnostics::API_NOT_MIGRATED,
4855+
false, "cudaLaunchHostFunc");
4856+
return;
48564857
}
48574858
std::string ReplStr;
48584859
llvm::raw_string_ostream OS(ReplStr);
4860+
std::string IndentStr = getIndent(LaunchKernelCall->getBeginLoc(),
4861+
DpctGlobalInfo::getSourceManager())
4862+
.str();
48594863
if (IsAssigned) {
48604864
OS << MapNames::getCheckErrorMacroName() << "(";
48614865
}
4862-
OS << MapNames::getDpctNamespace() << "kernel_launcher::launch(";
4863-
size_t ArgsNum = LaunchKernelCall->getNumArgs();
4864-
for (size_t i = 0; i < ArgsNum; i++) {
4865-
if (auto Arg = LaunchKernelCall->getArg(i)) {
4866-
if (i == 0) {
4867-
if (auto E = getAddressedRef(CalleeDRE, false, nullptr)) {
4868-
OS << ExprAnalysis::ref(E);
4866+
OS << ExprAnalysis::ref(LaunchKernelCall->getArg(0))
4867+
<< "->submit([&](sycl::handler &cgh) {" << getNL() << IndentStr
4868+
<< " cgh.host_task([=](){" << getNL() << IndentStr << " "
4869+
<< ExprAnalysis::ref(LaunchKernelCall->getArg(1)) << "("
4870+
<< ExprAnalysis::ref(LaunchKernelCall->getArg(2)) << ");" << getNL()
4871+
<< IndentStr << " });" << getNL() << IndentStr << "})";
4872+
if (IsAssigned) {
4873+
OS << ")";
4874+
}
4875+
auto Repl = new ReplaceStmt(LaunchKernelCall, OS.str());
4876+
Repl->setBlockLevelFormatFlag();
4877+
emplaceTransformation(Repl);
4878+
return;
4879+
} else {
4880+
const Expr *CalleeDRE = LaunchKernelCall->getArg(0);
4881+
bool IsFuncTypeErased = true;
4882+
auto QT = CalleeDRE->getType();
4883+
4884+
if (QT->isPointerType()) {
4885+
QT = QT->getPointeeType();
4886+
}
4887+
if (QT->isFunctionType()) {
4888+
IsFuncTypeErased = false;
4889+
}
4890+
4891+
if (!getAddressedRef(CalleeDRE)) {
4892+
if (IsFuncTypeErased) {
4893+
DpctGlobalInfo::setCVersionCUDALaunchUsed();
4894+
}
4895+
std::string ReplStr;
4896+
llvm::raw_string_ostream OS(ReplStr);
4897+
if (IsAssigned) {
4898+
OS << MapNames::getCheckErrorMacroName() << "(";
4899+
}
4900+
OS << MapNames::getDpctNamespace() << "kernel_launcher::launch(";
4901+
size_t ArgsNum = LaunchKernelCall->getNumArgs();
4902+
for (size_t i = 0; i < ArgsNum; i++) {
4903+
if (auto Arg = LaunchKernelCall->getArg(i)) {
4904+
if (i == 0) {
4905+
if (auto E = getAddressedRef(CalleeDRE, false, nullptr)) {
4906+
OS << ExprAnalysis::ref(E);
4907+
} else {
4908+
OS << ExprAnalysis::ref(Arg);
4909+
}
48694910
} else {
4870-
OS << ExprAnalysis::ref(Arg);
4911+
OS << ", " << ExprAnalysis::ref(Arg);
48714912
}
4872-
} else {
4873-
OS << ", " << ExprAnalysis::ref(Arg);
48744913
}
48754914
}
4876-
}
4877-
OS << ")";
4878-
if (IsAssigned) {
48794915
OS << ")";
4916+
if (IsAssigned) {
4917+
OS << ")";
4918+
}
4919+
emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str()));
4920+
return;
48804921
}
4881-
emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str()));
4882-
return;
4883-
}
48844922

4885-
if (!IsAssigned)
4886-
findAndRemoveTrailingSemicolon(LaunchKernelCall, Result);
4887-
if (DpctGlobalInfo::getInstance().buildLaunchKernelInfo(LaunchKernelCall,
4888-
IsAssigned)) {
4889-
emplaceTransformation(new ReplaceStmt(LaunchKernelCall, true, false, ""));
4923+
if (!IsAssigned)
4924+
findAndRemoveTrailingSemicolon(LaunchKernelCall, Result);
4925+
if (DpctGlobalInfo::getInstance().buildLaunchKernelInfo(LaunchKernelCall,
4926+
IsAssigned)) {
4927+
emplaceTransformation(
4928+
new ReplaceStmt(LaunchKernelCall, true, false, ""));
4929+
}
48904930
}
48914931
}
48924932
}

‎clang/lib/DPCT/SrcAPI/APINames.inc‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,7 @@ ENTRY(cudaGridDependencySynchronize, cudaGridDependencySynchronize, false, NO_FL
148148
ENTRY(cudaLaunchCooperativeKernel, cudaLaunchCooperativeKernel, true, NO_FLAG, P4, "Partial: DPCT1007")
149149
ENTRY(cudaLaunchCooperativeKernelMultiDevice, cudaLaunchCooperativeKernelMultiDevice, false, NO_FLAG, P4, "comment")
150150
ENTRY(cudaLaunchDevice, cudaLaunchDevice, false, NO_FLAG, P4, "comment")
151-
ENTRY(cudaLaunchHostFunc, cudaLaunchHostFunc, false, NO_FLAG, P4, "comment")
151+
ENTRY(cudaLaunchHostFunc, cudaLaunchHostFunc, true, NO_FLAG, P4, "comment")
152152
ENTRY(cudaLaunchKernel, cudaLaunchKernel, true, NO_FLAG, P0, "Partial: DPCT1007, success only when directly using of kernel function name")
153153
ENTRY(cudaLaunchKernelExC, cudaLaunchKernelExC, false, NO_FLAG, P4, "comment")
154154
ENTRY(cudaSetDoubleForDevice, cudaSetDoubleForDevice, false, NO_FLAG, P0, "comment")

‎clang/lib/DPCT/UserDefinedRules/UserDefinedRules.cpp‎

Lines changed: 138 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,7 @@ namespace clang {
2626
namespace dpct {
2727
std::vector<clang::tooling::UnifiedPath> MetaRuleObject::RuleFiles;
2828
std::vector<std::shared_ptr<MetaRuleObject>> MetaRules;
29+
llvm::DenseSet<llvm::StringRef> ProcessedYamlFiles;
2930

3031
OutputBuilder::~OutputBuilder() {}
3132

@@ -344,17 +345,147 @@ MetaRuleObject::PatternRewriter::PatternRewriter(
344345
Subrules = S;
345346
}
346347

348+
// Read a YAML file recursively and substitute any "!include <filename>"
349+
// directive with the contents of the referenced file.
350+
std::unique_ptr<llvm::MemoryBuffer>
351+
readYAMLFile(const llvm::StringRef &RuleFilePath) {
352+
// Check if the rule file has already been processed
353+
// to avoid infinite recursion
354+
if (!ProcessedYamlFiles.insert(RuleFilePath).second) {
355+
return llvm::MemoryBuffer::getMemBufferCopy("");
356+
}
357+
358+
// Load the rule file into a MemoryBuffer
359+
llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> Buffer =
360+
llvm::MemoryBuffer::getFile(RuleFilePath);
361+
if (!Buffer) {
362+
llvm::errs() << "Error: failed to read " << RuleFilePath << ": "
363+
<< Buffer.getError().message() << "\n";
364+
clang::dpct::ShowStatus(MigrationErrorInvalidRuleFilePath);
365+
dpctExit(MigrationErrorInvalidRuleFilePath);
366+
}
367+
368+
// Get the directory path of the rule file.
369+
llvm::SmallString<128> DirectoryPath(RuleFilePath);
370+
llvm::sys::path::remove_filename(DirectoryPath);
371+
372+
// Iterate over the input line by line.
373+
std::stringstream Output, IncRuleFilePath;
374+
const std::string IncDirective = "!include";
375+
376+
size_t Idx = 0;
377+
bool IncDirectiveFound = false;
378+
bool SkipWhiteSpaces = false;
379+
bool IsLineBeginning = true;
380+
381+
llvm::StringRef BuffContent = std::move(*Buffer)->getBuffer();
382+
const auto BuffSize = BuffContent.size();
383+
while (Idx < BuffSize) {
384+
unsigned char Ch = BuffContent[Idx];
385+
386+
// Skip white spaces at the beginning of the line if it contains a directive
387+
if (IsLineBeginning) {
388+
auto i = Idx;
389+
auto c = Ch;
390+
391+
// lookahead for "!" directive after white spaces
392+
for (; i < BuffContent.size(); i++) {
393+
c = BuffContent[i];
394+
395+
// Stop at new line or first non-white space character
396+
if (c == '\n' || !std::isspace(c)) {
397+
break;
398+
}
399+
}
400+
401+
// Check if the line starts with a directive
402+
if (c == '!') {
403+
// Move Idx to the beginning of directive
404+
Idx = i;
405+
406+
// Check if the directive is "!include"
407+
if (!IncDirectiveFound && Idx + IncDirective.length() <= BuffSize &&
408+
BuffContent.substr(Idx, IncDirective.length()) == IncDirective) {
409+
// Move Idx to the end of directive
410+
Idx += IncDirective.length();
411+
IncDirectiveFound = true;
412+
SkipWhiteSpaces = true;
413+
}
414+
415+
// Update current character
416+
Ch = BuffContent[Idx];
417+
}
418+
419+
IsLineBeginning = false;
420+
}
421+
422+
// Skip return carriage character
423+
if (Ch == '\r') {
424+
Idx++;
425+
continue;
426+
}
427+
428+
// Process IncRuleFilePath at end of the line
429+
if (Ch == '\n') {
430+
if (IncDirectiveFound) {
431+
auto IncRuleFilePathStr = IncRuleFilePath.str();
432+
433+
if (!IncRuleFilePathStr.empty()) {
434+
// Find the absolute path for the included rule file path
435+
llvm::SmallString<128> IncRuleFileAbsPath = DirectoryPath;
436+
llvm::sys::path::append(IncRuleFileAbsPath, IncRuleFilePathStr);
437+
438+
// Recursively process the included file
439+
if (llvm::sys::fs::exists(IncRuleFileAbsPath)) {
440+
Output << readYAMLFile(IncRuleFileAbsPath.str())->getBuffer().str();
441+
} else {
442+
Output << readYAMLFile(IncRuleFilePathStr)->getBuffer().str();
443+
}
444+
445+
// Clear the contents of include rule file path
446+
IncRuleFilePath.str("");
447+
}
448+
}
449+
450+
// Reset include directive info for each new line
451+
IncDirectiveFound = false;
452+
SkipWhiteSpaces = false;
453+
IsLineBeginning = true;
454+
}
455+
456+
if (IncDirectiveFound) {
457+
// Skip adding quotes to the include rule file path
458+
if (Ch == '"' || Ch == '\'') {
459+
Idx++;
460+
// Flip white space skip flag at the boundaries of quotes
461+
SkipWhiteSpaces = !SkipWhiteSpaces;
462+
continue;
463+
}
464+
465+
// Skip white space characters
466+
if (SkipWhiteSpaces && std::isspace(Ch)) {
467+
Idx++;
468+
continue;
469+
}
470+
471+
// Append the character to the include rule file path for !include line
472+
IncRuleFilePath << Ch;
473+
} else {
474+
// Append the character to the output buffer
475+
Output << Ch;
476+
}
477+
478+
Idx++;
479+
}
480+
481+
return llvm::MemoryBuffer::getMemBufferCopy(Output.str(), RuleFilePath);
482+
}
483+
347484
void importRules(std::vector<clang::tooling::UnifiedPath> &RuleFiles) {
348485
for (auto &RuleFile : RuleFiles) {
349486
// open the yaml file
350487
llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> Buffer =
351-
llvm::MemoryBuffer::getFile(RuleFile.getCanonicalPath());
352-
if (!Buffer) {
353-
llvm::errs() << "Error: failed to read " << RuleFile << ": "
354-
<< Buffer.getError().message() << "\n";
355-
clang::dpct::ShowStatus(MigrationErrorInvalidRuleFilePath);
356-
dpctExit(MigrationErrorInvalidRuleFilePath);
357-
}
488+
readYAMLFile(RuleFile.getCanonicalPath());
358489

359490
// load rules
360491
std::vector<std::shared_ptr<MetaRuleObject>> CurrentRules;

‎clang/runtime/dpct-rt/include/dpct/device.hpp‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -592,7 +592,14 @@ class device_ext : public sycl::device {
592592
lock.unlock();
593593
for (const auto &q : current_queues) {
594594
if (q->is_in_order()) {
595+
#ifdef __INTEL_LLVM_COMPILER
595596
last_events.push_back(q->ext_oneapi_get_last_event());
597+
#else
598+
auto last_event = q->ext_oneapi_get_last_event();
599+
if (last_event) {
600+
last_events.push_back(*last_event);
601+
}
602+
#endif
596603
}
597604
}
598605
// Guard the destruct of current_queues to make sure the ref count is safe.

0 commit comments

Comments
 (0)