Skip to content

Commit

Permalink
Merge branch 'SYCLomatic' into fpointer
Browse files Browse the repository at this point in the history
  • Loading branch information
@intwanghao
intwanghao committed Dec 18, 2024
2 parents 829c3d5 + cfe128c commit 2b73dcb
Show file tree
Hide file tree
Showing 20 changed files with 543 additions and 348 deletions.
5 changes: 3 additions & 2 deletions clang/lib/DPCT/AnalysisInfo.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -246,7 +246,8 @@ std::shared_ptr<TargetType> makeTextureObjectInfo(const ValueDecl *D,
if (auto VD = dyn_cast<VarDecl>(D)) {
return std::make_shared<TargetType>(VD);
}
} else if (auto PVD = dyn_cast<ParmVarDecl>(D)) {
} else if (const auto *PVD = dyn_cast<ParmVarDecl>(D);
PVD && PVD->getTypeSourceInfo()) {
return std::make_shared<TargetType>(PVD);
}
return std::shared_ptr<TargetType>();
Expand Down Expand Up @@ -5043,7 +5044,7 @@ void DeviceFunctionDecl::buildTextureObjectParamsInfo(
std::string ParamName =
DpctGlobalInfo::getUnqualifiedTypeName(Param->getType());
if (ParamName == "cudaTextureObject_t" ||
ParamName == "cudaSurfaceObject_t") {
ParamName == "cudaSurfaceObject_t" || ParamName == "CUsurfObject") {
TextureObjectList[Idx] = std::make_shared<TextureObjectInfo>(Param);
}
}
Expand Down
7 changes: 7 additions & 0 deletions clang/lib/DPCT/RuleInfra/MapNames.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -524,6 +524,13 @@ void MapNames::setExplicitNamespaceMap(
"ext::oneapi::experimental::unsampled_image_handle"
: getDpctNamespace() + "image_wrapper_base_p",
HelperFeatureEnum::device_ext)},
{"CUsurfObject",
std::make_shared<TypeNameRule>(
DpctGlobalInfo::useExtBindlessImages()
? getClNamespace() +
"ext::oneapi::experimental::unsampled_image_handle"
: getDpctNamespace() + "image_wrapper_base_p",
HelperFeatureEnum::device_ext)},
{"textureReference",
std::make_shared<TypeNameRule>(getDpctNamespace() + "image_wrapper_base",
HelperFeatureEnum::device_ext)},
Expand Down
3 changes: 2 additions & 1 deletion clang/lib/DPCT/RulesInclude/InclusionHeaders.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -200,7 +200,8 @@ void IncludesCallbacks::InclusionDirective(
DpctGlobalInfo::getIncludeMapSet().push_back({IncludedFile, Repl});
}
}
return;
if (Global.isInRoot(IncludedFile))
return;
}

if (!Global.isInAnalysisScope(LocInfo.first) &&
Expand Down
3 changes: 1 addition & 2 deletions clang/lib/DPCT/RulesInclude/InclusionHeaders.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -50,8 +50,7 @@ REGIST_INCLUSION("curand_kernel.h", FullMatch, Rng, Replace, false,
REGIST_INCLUSION("cusparse.h", FullMatch, Sparse, Replace, false,
HeaderType::HT_DPCT_SPBLAS_Utils)
REGIST_INCLUSION("cusparse_v2.h", FullMatch, Sparse, Replace, false,
HeaderType::HT_DPCT_SPBLAS_Utils,
HeaderType::HT_DPCT_BLAS_Utils)
HeaderType::HT_DPCT_SPBLAS_Utils)

REGIST_INCLUSION("cufft.h", FullMatch, FFT, Replace, false,
HeaderType::HT_DPCT_FFT_Utils)
Expand Down
29 changes: 29 additions & 0 deletions clang/lib/DPCT/RulesLang/APINamesTexture.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,14 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
ARG(0), QUEUESTR)),
DELETER_FACTORY_ENTRY("cudaDestroySurfaceObject", ARG(0))))

ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("cuSurfObjectDestroy",
CALL(MapNames::getDpctNamespace() +
"experimental::destroy_bindless_image",
ARG(0), QUEUESTR)),
DELETER_FACTORY_ENTRY("cuSurfObjectDestroy", ARG(0))))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
Expand Down Expand Up @@ -120,6 +128,17 @@ CONDITIONAL_FACTORY_ENTRY(
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
"cudaGetSurfaceObjectResourceDesc", DEREF(0),
MEMBER_CALL(ARG(1), true, "get_data")))))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
"cuSurfObjectGetResourceDesc", DEREF(0),
CALL(MapNames::getDpctNamespace() + "experimental::get_data", ARG(1)))),
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
"cuSurfObjectGetResourceDesc", DEREF(0),
MEMBER_CALL(ARG(1), true, "get_data")))))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
FEATURE_REQUEST_FACTORY(
Expand Down Expand Up @@ -263,6 +282,16 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
"create_image_wrapper",
DEREF(1)))))

FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
"cuSurfObjectCreate", DEREF(0),
CALL(DpctGlobalInfo::useExtBindlessImages()
? MapNames::getDpctNamespace() +
"experimental::create_bindless_image"
: MapNames::getDpctNamespace() +
"create_image_wrapper",
DEREF(1)))))

ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc",
Diagnostics::API_NOT_MIGRATED)

Expand Down
7 changes: 7 additions & 0 deletions clang/lib/DPCT/RulesLang/RulesLang.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1082,6 +1082,13 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
}
Str = Itr->second;
}
} else if (llvm::StringRef(TypeStr).starts_with("cublas")) {
// In most cases, we do not need to insert blas_utils.hpp manually since
// the cublas_v2.h will be migrated. However, when the include directive
// of cublas_v2.h is not in the in-root, the migrated code cannot be
// built successfully.
DpctGlobalInfo::getInstance().insertHeader(
TL->getBeginLoc(), HeaderType::HT_DPCT_BLAS_Utils);
}

// Add '#include <complex>' directive to the file only once
Expand Down
28 changes: 15 additions & 13 deletions clang/lib/DPCT/RulesLang/RulesLangTexture.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -496,12 +496,11 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
)
.bind("tex"),
this);
MF.addMatcher(
typeLoc(
loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
"cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
.bind("texObj"),
this);
MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
"cudaTextureObject_t", "cudaSurfaceObject_t",
"CUsurfObject", "CUtexObject"))))))
.bind("texObj"),
this);
MF.addMatcher(
memberExpr(
hasObjectExpression(hasType(type(hasUnqualifiedDesugaredType(
Expand All @@ -516,13 +515,13 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
loc(qualType(hasDeclaration(namedDecl(hasAnyName(
"cudaChannelFormatDesc", "cudaChannelFormatKind",
"cudaTextureDesc", "cudaResourceDesc", "cudaResourceType",
"cudaTextureAddressMode", "cudaTextureFilterMode", "cudaArray",
"cudaArray_t", "CUarray_st", "CUarray", "CUarray_format",
"CUarray_format_enum", "CUresourcetype", "CUresourcetype_enum",
"CUaddress_mode", "CUaddress_mode_enum", "CUfilter_mode",
"CUfilter_mode_enum", "CUDA_RESOURCE_DESC", "CUDA_TEXTURE_DESC",
"CUtexref", "textureReference", "cudaMipmappedArray",
"cudaMipmappedArray_t"))))))
"CUDA_RESOURCE_DESC", "cudaTextureAddressMode",
"cudaTextureFilterMode", "cudaArray", "cudaArray_t", "CUarray_st",
"CUarray", "CUarray_format", "CUarray_format_enum",
"CUresourcetype", "CUresourcetype_enum", "CUaddress_mode",
"CUaddress_mode_enum", "CUfilter_mode", "CUfilter_mode_enum",
"CUDA_TEXTURE_DESC", "CUtexref", "textureReference",
"cudaMipmappedArray", "cudaMipmappedArray_t"))))))
.bind("texType"),
this);

Expand Down Expand Up @@ -571,6 +570,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
"cudaCreateSurfaceObject",
"cudaDestroySurfaceObject",
"cudaGetSurfaceObjectResourceDesc",
"cuSurfObjectCreate",
"cuSurfObjectGetResourceDesc",
"cuSurfObjectDestroy",
"cuArray3DCreate_v2",
"cuArrayCreate_v2",
"cuArrayDestroy",
Expand Down
3 changes: 3 additions & 0 deletions clang/lib/DPCT/RulesMathLib/MapNamesBlas.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1642,6 +1642,9 @@ void MapNamesBlas::setExplicitNamespaceMap(
{"CUBLASLT_EPILOGUE_GELU_AUX_BIAS",
MapNames::getLibraryHelperNamespace() +
"blas_gemm::experimental::epilogue_t::gelu_aux_bias"},
{"CUBLASLT_EPILOGUE_DGELU",
MapNames::getLibraryHelperNamespace() +
"blas_gemm::experimental::epilogue_t::dgelu"},
{"CUBLASLT_MATRIX_TRANSFORM_DESC_SCALE_TYPE",
MapNames::getLibraryHelperNamespace() +
"blas_gemm::experimental::transform_desc_t::attribute::scale_type"},
Expand Down
6 changes: 3 additions & 3 deletions clang/lib/DPCT/SrcAPI/APINames.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1992,9 +1992,9 @@ ENTRY(cuTexObjectGetResourceViewDesc, cuTexObjectGetResourceViewDesc, false, NO_
ENTRY(cuTexObjectGetTextureDesc, cuTexObjectGetTextureDesc, true, NO_FLAG, P4, "Successful")

// Surface Object Management
ENTRY(cuSurfObjectCreate, cuSurfObjectCreate, false, NO_FLAG, P4, "comment")
ENTRY(cuSurfObjectDestroy, cuSurfObjectDestroy, false, NO_FLAG, P4, "comment")
ENTRY(cuSurfObjectGetResourceDesc, cuSurfObjectGetResourceDesc, false, NO_FLAG, P4, "comment")
ENTRY(cuSurfObjectCreate, cuSurfObjectCreate, true, NO_FLAG, P4, "Successful")
ENTRY(cuSurfObjectDestroy, cuSurfObjectDestroy, true, NO_FLAG, P4, "Successful")
ENTRY(cuSurfObjectGetResourceDesc, cuSurfObjectGetResourceDesc, true, NO_FLAG, P4, "Successfuls")

// Tensor Map Object Managment
ENTRY(cuTensorMapEncodeIm2col, cuTensorMapEncodeIm2col, false, NO_FLAG, P4, "comment")
Expand Down
Loading

0 comments on commit 2b73dcb

Please sign in to comment.