Skip to content

Commit

Permalink
Merge remote-tracking branch 'origin/SYCLomatic' into support_more_ge…
Browse files Browse the repository at this point in the history 
…mm_batch
  • Loading branch information
@zhiweij1
zhiweij1 committed Jul 5, 2024
2 parents 3b40caf + 6ea742c commit 258e107
Show file tree
Hide file tree
Showing 21 changed files with 531 additions and 198 deletions.
7 changes: 4 additions & 3 deletions clang/lib/DPCT/APINames.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -385,14 +385,15 @@ ENTRY(cudaRuntimeGetVersion, cudaRuntimeGetVersion, true, NO_FLAG, P0, "DPCT1043
// Graph Management
ENTRY(cudaGraph_t, cudaGraph_t, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaGraphExec_t, cudaGraphExec_t, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaGraphNode_t, cudaGraphNode_t, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaDeviceGetGraphMemAttribute, cudaDeviceGetGraphMemAttribute, false, NO_FLAG, P4, "comment")
ENTRY(cudaDeviceGraphMemTrim, cudaDeviceGraphMemTrim, false, NO_FLAG, P4, "comment")
ENTRY(cudaDeviceSetGraphMemAttribute, cudaDeviceSetGraphMemAttribute, false, NO_FLAG, P4, "comment")
ENTRY(cudaGetCurrentGraphExec, cudaGetCurrentGraphExec, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddChildGraphNode, cudaGraphAddChildGraphNode, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddDependencies, cudaGraphAddDependencies, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddDependencies, cudaGraphAddDependencies, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaGraphAddDependencies_v2, cudaGraphAddDependencies_v2, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddEmptyNode, cudaGraphAddEmptyNode, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddEmptyNode, cudaGraphAddEmptyNode, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaGraphAddEventRecordNode, cudaGraphAddEventRecordNode, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddEventWaitNode, cudaGraphAddEventWaitNode, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphAddExternalSemaphoresSignalNode, cudaGraphAddExternalSemaphoresSignalNode, false, NO_FLAG, P4, "comment")
Expand Down Expand Up @@ -433,7 +434,7 @@ ENTRY(cudaGraphExecMemcpyNodeSetParamsFromSymbol, cudaGraphExecMemcpyNodeSetPara
ENTRY(cudaGraphExecMemcpyNodeSetParamsToSymbol, cudaGraphExecMemcpyNodeSetParamsToSymbol, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphExecMemsetNodeSetParams, cudaGraphExecMemsetNodeSetParams, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphExecNodeSetParams, cudaGraphExecNodeSetParams, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphExecUpdate, cudaGraphExecUpdate, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphExecUpdate, cudaGraphExecUpdate, true, NO_FLAG, P4, "Successful/DPCT1119")
ENTRY(cudaGraphExternalSemaphoresSignalNodeGetParams, cudaGraphExternalSemaphoresSignalNodeGetParams, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphExternalSemaphoresSignalNodeSetParams, cudaGraphExternalSemaphoresSignalNodeSetParams, false, NO_FLAG, P4, "comment")
ENTRY(cudaGraphExternalSemaphoresWaitNodeGetParams, cudaGraphExternalSemaphoresWaitNodeGetParams, false, NO_FLAG, P4, "comment")
Expand Down
31 changes: 31 additions & 0 deletions clang/lib/DPCT/APINamesGraph.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -35,3 +35,34 @@ CONDITIONAL_FACTORY_ENTRY(
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaGraphLaunch"),
ARG("--use-experimental-features=graph")))

CONDITIONAL_FACTORY_ENTRY(
UseExtGraph,
CALL_FACTORY_ENTRY("cudaGraphAddEmptyNode",
CALL(MapNames::getDpctNamespace() +
"experimental::add_empty_node",
ARG(0), ARG(1), ARG(2), ARG(3))),
UNSUPPORT_FACTORY_ENTRY("cudaGraphAddEmptyNode",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaGraphAddEmptyNode"),
ARG("--use-experimental-features=graph")))

CONDITIONAL_FACTORY_ENTRY(
UseExtGraph,
CALL_FACTORY_ENTRY("cudaGraphAddDependencies",
CALL(MapNames::getDpctNamespace() +
"experimental::add_dependencies",
ARG(0), ARG(1), ARG(2), ARG(3))),
UNSUPPORT_FACTORY_ENTRY("cudaGraphAddDependencies",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaGraphAddDependencies"),
ARG("--use-experimental-features=graph")))

CONDITIONAL_FACTORY_ENTRY(
UseExtGraph,
MEMBER_CALL_FACTORY_ENTRY("cudaGraphExecUpdate", ARG(0), true, "update",
DEREF(1)),
UNSUPPORT_FACTORY_ENTRY("cudaGraphExecUpdate",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("cudaGraphExecUpdate"),
ARG("--use-experimental-features=graph")))
9 changes: 9 additions & 0 deletions clang/lib/DPCT/APINamesTemplateType.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -393,3 +393,12 @@ TYPE_REWRITE_ENTRY(
"experimental::command_graph_exec_ptr")),
WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
STR("--use-experimental-features=graph"))))

TYPE_REWRITE_ENTRY(
"cudaGraphNode_t",
TYPE_CONDITIONAL_FACTORY(
checkEnableGraphForType(),
TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
"experimental::node_ptr")),
WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
STR("--use-experimental-features=graph"))))
122 changes: 75 additions & 47 deletions clang/lib/DPCT/ASTTraversal.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1694,35 +1694,35 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
"thrust::device_ptr", "thrust::device_reference",
"thrust::host_vector", "cublasHandle_t", "CUevent_st", "__half",
"half", "__half2", "half2", "cudaMemoryAdvise", "cudaError_enum",
"cudaDeviceProp", "cudaPitchedPtr", "thrust::counting_iterator",
"thrust::transform_iterator", "thrust::permutation_iterator",
"thrust::iterator_difference", "cusolverDnHandle_t",
"cusolverDnParams_t", "gesvdjInfo_t", "syevjInfo_t",
"thrust::device_malloc_allocator", "thrust::divides",
"thrust::tuple", "thrust::maximum", "thrust::multiplies",
"thrust::plus", "cudaDataType_t", "cudaError_t", "CUresult",
"CUdevice", "cudaEvent_t", "cublasStatus_t", "cuComplex",
"cuFloatComplex", "cuDoubleComplex", "CUevent",
"cublasFillMode_t", "cublasDiagType_t", "cublasSideMode_t",
"cublasOperation_t", "cusolverStatus_t", "cusolverEigType_t",
"cusolverEigMode_t", "curandStatus_t", "cudaStream_t",
"cusparseStatus_t", "cusparseDiagType_t", "cusparseFillMode_t",
"cusparseIndexBase_t", "cusparseMatrixType_t",
"cusparseAlgMode_t", "cusparseOperation_t", "cusparseMatDescr_t",
"cusparseHandle_t", "CUcontext", "cublasPointerMode_t",
"cusparsePointerMode_t", "cublasGemmAlgo_t",
"cusparseSolveAnalysisInfo_t", "cudaDataType", "cublasDataType_t",
"curandState_t", "curandState", "curandStateXORWOW_t",
"curandStateXORWOW", "curandStatePhilox4_32_10_t",
"curandStatePhilox4_32_10", "curandStateMRG32k3a_t",
"curandStateMRG32k3a", "thrust::minus", "thrust::negate",
"thrust::logical_or", "thrust::equal_to", "thrust::less",
"cudaSharedMemConfig", "curandGenerator_t", "curandRngType_t",
"curandOrdering_t", "cufftHandle", "cufftReal", "cufftDoubleReal",
"cufftComplex", "cufftDoubleComplex", "cufftResult_t",
"cufftResult", "cufftType_t", "cufftType", "thrust::pair",
"CUdeviceptr", "cudaDeviceAttr", "CUmodule", "CUjit_option",
"CUfunction", "cudaMemcpyKind", "cudaComputeMode",
"cudaDeviceProp", "cudaGraphExecUpdateResult", "cudaPitchedPtr",
"thrust::counting_iterator", "thrust::transform_iterator",
"thrust::permutation_iterator", "thrust::iterator_difference",
"cusolverDnHandle_t", "cusolverDnParams_t", "gesvdjInfo_t",
"syevjInfo_t", "thrust::device_malloc_allocator",
"thrust::divides", "thrust::tuple", "thrust::maximum",
"thrust::multiplies", "thrust::plus", "cudaDataType_t",
"cudaError_t", "CUresult", "CUdevice", "cudaEvent_t",
"cublasStatus_t", "cuComplex", "cuFloatComplex",
"cuDoubleComplex", "CUevent", "cublasFillMode_t",
"cublasDiagType_t", "cublasSideMode_t", "cublasOperation_t",
"cusolverStatus_t", "cusolverEigType_t", "cusolverEigMode_t",
"curandStatus_t", "cudaStream_t", "cusparseStatus_t",
"cusparseDiagType_t", "cusparseFillMode_t", "cusparseIndexBase_t",
"cusparseMatrixType_t", "cusparseAlgMode_t",
"cusparseOperation_t", "cusparseMatDescr_t", "cusparseHandle_t",
"CUcontext", "cublasPointerMode_t", "cusparsePointerMode_t",
"cublasGemmAlgo_t", "cusparseSolveAnalysisInfo_t", "cudaDataType",
"cublasDataType_t", "curandState_t", "curandState",
"curandStateXORWOW_t", "curandStateXORWOW",
"curandStatePhilox4_32_10_t", "curandStatePhilox4_32_10",
"curandStateMRG32k3a_t", "curandStateMRG32k3a", "thrust::minus",
"thrust::negate", "thrust::logical_or", "thrust::equal_to",
"thrust::less", "cudaSharedMemConfig", "curandGenerator_t",
"curandRngType_t", "curandOrdering_t", "cufftHandle", "cufftReal",
"cufftDoubleReal", "cufftComplex", "cufftDoubleComplex",
"cufftResult_t", "cufftResult", "cufftType_t", "cufftType",
"thrust::pair", "CUdeviceptr", "cudaDeviceAttr", "CUmodule",
"CUjit_option", "CUfunction", "cudaMemcpyKind", "cudaComputeMode",
"__nv_bfloat16", "cooperative_groups::__v1::thread_group",
"cooperative_groups::__v1::thread_block", "libraryPropertyType_t",
"libraryPropertyType", "cudaDataType_t", "cudaDataType",
Expand Down Expand Up @@ -1750,12 +1750,13 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
.bind("cudaTypeDef"),
this);

MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName(
"cooperative_groups::__v1::coalesced_group",
"cooperative_groups::__v1::thread_block_tile",
"cudaGraph_t", "cudaGraphExec_t"))))))
.bind("cudaTypeDefEA"),
this);
MF.addMatcher(
typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName(
"cooperative_groups::__v1::coalesced_group",
"cooperative_groups::__v1::thread_block_tile", "cudaGraph_t",
"cudaGraphExec_t", "cudaGraphNode_t"))))))
.bind("cudaTypeDefEA"),
this);
MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl(
hasAnyTemplateArgument(refersToType(hasDeclaration(
namedDecl(hasName("use_default"))))))))
Expand Down Expand Up @@ -2284,6 +2285,11 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
std::string CanonicalTypeStr = DpctGlobalInfo::getUnqualifiedTypeName(
TL->getType().getCanonicalType());

if (CanonicalTypeStr == "cudaGraphExecUpdateResult") {
report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
CanonicalTypeStr);
return;
}
if (CanonicalTypeStr == "cooperative_groups::__v1::thread_group" ||
CanonicalTypeStr == "cooperative_groups::__v1::thread_block") {
if (auto ETL = TL->getUnqualifiedLoc().getAs<ElaboratedTypeLoc>()) {
Expand Down Expand Up @@ -2883,9 +2889,16 @@ void VectorTypeOperatorRule::registerMatcher(MatchFinder &MF) {
this);

// Matches call of user overloaded operator
MF.addMatcher(cxxOperatorCallExpr(callee(vectorTypeOverLoadedOperator()))
.bind("callOverloadedOperator"),
MF.addMatcher(cxxOperatorCallExpr(callee(vectorTypeOverLoadedOperator()),
hasAncestor(vectorTypeOverLoadedOperator()))
.bind("callOverloadedOperatorInOverloadedOperator"),
this);

MF.addMatcher(
cxxOperatorCallExpr(callee(vectorTypeOverLoadedOperator()),
unless(hasAncestor(vectorTypeOverLoadedOperator())))
.bind("callOverloadedOperatorNotInOverloadedOperator"),
this);
}

const char VectorTypeOperatorRule::NamespaceName[] =
Expand Down Expand Up @@ -2967,10 +2980,15 @@ void VectorTypeOperatorRule::MigrateOverloadedOperatorDecl(
}

void VectorTypeOperatorRule::MigrateOverloadedOperatorCall(
const MatchFinder::MatchResult &Result, const CXXOperatorCallExpr *CE) {
const MatchFinder::MatchResult &Result, const CXXOperatorCallExpr *CE,
bool InOverloadedOperator) {
if (!CE)
return;

if (!InOverloadedOperator &&
(DpctGlobalInfo::findAncestor<FunctionTemplateDecl>(CE) ||
DpctGlobalInfo::findAncestor<ClassTemplateDecl>(CE))) {
return;
}
// Explicitly call user overloaded operator
//
// For non-assignment operator:
Expand Down Expand Up @@ -3004,8 +3022,17 @@ void VectorTypeOperatorRule::runRule(const MatchFinder::MatchResult &Result) {
Result, getNodeAsType<FunctionDecl>(Result, "overloadedOperatorDecl"));

// Explicitly call user overloaded operator
MigrateOverloadedOperatorCall(Result, getNodeAsType<CXXOperatorCallExpr>(
Result, "callOverloadedOperator"));
MigrateOverloadedOperatorCall(
Result,
getNodeAsType<CXXOperatorCallExpr>(
Result, "callOverloadedOperatorInOverloadedOperator"),
true);

MigrateOverloadedOperatorCall(
Result,
getNodeAsType<CXXOperatorCallExpr>(
Result, "callOverloadedOperatorNotInOverloadedOperator"),
false);
}

REGISTER_RULE(VectorTypeOperatorRule, PassKind::PK_Migration)
Expand Down Expand Up @@ -6781,8 +6808,8 @@ EventQueryTraversal::buildCallReplacement(const CallExpr *Call) {
static std::string MemberName = "get_info<" + MapNames::getClNamespace() +
"info::event::command_execution_status>";
std::string ReplStr;
MemberCallPrinter<const Expr *, StringRef> Printer(Call->getArg(0), true,
MemberName);
MemberCallPrinter<const Expr *, StringRef, false> Printer(Call->getArg(0),
true, MemberName);
llvm::raw_string_ostream OS(ReplStr);
Printer.print(OS);
return new ReplaceStmt(Call, std::move(OS.str()));
Expand Down Expand Up @@ -13474,7 +13501,7 @@ void TextureRule::runRule(const MatchFinder::MatchResult &Result) {
std::shared_ptr<CallExprRewriter> Rewriter =
std::make_shared<AssignableRewriter>(
CE, std::make_shared<PrinterRewriter<MemberCallPrinter<
const Expr *, RenameWithSuffix, StringRef>>>(
const Expr *, RenameWithSuffix, false, StringRef>>>(
CE, Name, CE->getArg(0), true,
RenameWithSuffix("set", MethodName), Value));
std::optional<std::string> Result = Rewriter->rewrite();
Expand Down Expand Up @@ -13832,8 +13859,8 @@ bool TextureRule::SettersMerger::applyResult() {

std::string ReplacedText;
llvm::raw_string_ostream OS(ReplacedText);
MemberCallPrinter<StringRef, StringRef, std::vector<std::string>> Printer(
D->getName(), IsArrow, "set", std::move(ArgsList));
MemberCallPrinter<StringRef, StringRef, false, std::vector<std::string>>
Printer(D->getName(), IsArrow, "set", std::move(ArgsList));
Printer.print(OS);

Inserter.success(OS.str());
Expand Down Expand Up @@ -15127,7 +15154,8 @@ REGISTER_RULE(AssertRule, PassKind::PK_Migration)
void GraphRule::registerMatcher(MatchFinder &MF) {
auto functionName = [&]() {
return hasAnyName("cudaGraphInstantiate", "cudaGraphLaunch",
"cudaGraphExecDestroy");
"cudaGraphExecDestroy", "cudaGraphAddEmptyNode",
"cudaGraphAddDependencies", "cudaGraphExecUpdate");
};
MF.addMatcher(
callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"),
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/DPCT/ASTTraversal.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -574,7 +574,7 @@ class VectorTypeOperatorRule
const FunctionDecl *FD);
void MigrateOverloadedOperatorCall(
const ast_matchers::MatchFinder::MatchResult &Result,
const CXXOperatorCallExpr *CE);
const CXXOperatorCallExpr *CE, bool InOverloadedOperator);

private:
static const char NamespaceName[];
Expand Down
Loading

0 comments on commit 258e107

Please sign in to comment.