Skip to content

Commit b12a054

Browse files
zhiweij1zhiweij1
committed
[SYCLomatic] Add 12.9 header file migration support
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent 971e72b commit b12a054

37 files changed

Lines changed: 116 additions & 78 deletions

‎README.md‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ in daily releases. None of the branches in the project are stable or rigorously
2727
tested for production quality control, so the quality of these releases is
2828
expected to be similar to the daily releases.
2929

30-
SYCLomatic supports migrating programs implemented with CUDA versions 8.0, 9.x, 10.x, 11.x, 12.0-12.8. The list of supported languages and versions may be extended in the future.
30+
SYCLomatic supports migrating programs implemented with CUDA versions 8.0, 9.x, 10.x, 11.x, 12.0-12.9. The list of supported languages and versions may be extended in the future.
3131

3232
## Build from source code
3333
### Prerequisites

‎clang/include/clang/Basic/Cuda.h‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -50,13 +50,16 @@ enum class CudaVersion {
5050
CUDA_126,
5151
CUDA_128,
5252
#ifdef SYCLomatic_CUSTOMIZATION
53-
FULLY_SUPPORTED = CUDA_126,
53+
CUDA_129,
54+
FULLY_SUPPORTED = CUDA_123,
55+
PARTIALLY_SUPPORTED =
56+
CUDA_129, // Partially supported. Proceed with a warning.
5457
#else
5558
FULLY_SUPPORTED = CUDA_123,
56-
#endif
5759
PARTIALLY_SUPPORTED =
5860
CUDA_128, // Partially supported. Proceed with a warning.
59-
NEW = 10000, // Too new. Issue a warning, but allow using it.
61+
#endif
62+
NEW = 10000, // Too new. Issue a warning, but allow using it.
6063
};
6164
const char *CudaVersionToString(CudaVersion V);
6265
#ifdef SYCLomatic_CUSTOMIZATION

‎clang/lib/Basic/Cuda.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,9 @@ static const CudaVersionMapEntry CudaNameVersionMap[] = {
4545
CUDA_ENTRY(12, 5),
4646
CUDA_ENTRY(12, 6),
4747
CUDA_ENTRY(12, 8),
48+
#ifdef SYCLomatic_CUSTOMIZATION
49+
CUDA_ENTRY(12, 9),
50+
#endif
4851
{"", CudaVersion::NEW, llvm::VersionTuple(std::numeric_limits<int>::max())},
4952
{"unknown", CudaVersion::UNKNOWN, {}} // End of list tombstone.
5053
};

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

Lines changed: 16 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,7 @@ using namespace clang::tooling;
5353
extern clang::tooling::UnifiedPath DpctInstallPath; // Installation directory for this tool
5454
extern DpctOption<opt, bool> ProcessAll;
5555
extern DpctOption<opt, bool> AsyncHandler;
56+
extern int ThrustVersion;
5657

5758
namespace clang{
5859
namespace dpct{
@@ -280,6 +281,17 @@ void MiscAPIRule::runRule(const MatchFinder::MatchResult &Result) {
280281

281282
// Rule for types migration in var declarations and field declarations
282283
void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
284+
int ThrustMajorVersion = ThrustVersion / 100000;
285+
int ThrustMinorVersion = ThrustVersion / 100 % 1000;
286+
287+
auto thrustNamespace = [=]() -> std::string {
288+
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 8) {
289+
// For CUDA-12.9
290+
return "cuda::std::";
291+
}
292+
return "thrust::";
293+
};
294+
283295
MF.addMatcher(
284296
typeLoc(
285297
loc(qualType(hasDeclaration(namedDecl(hasAnyName(
@@ -294,7 +306,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
294306
"thrust::permutation_iterator", "thrust::iterator_difference",
295307
"cusolverDnHandle_t", "cusolverDnParams_t", "gesvdjInfo_t",
296308
"syevjInfo_t", "thrust::device_malloc_allocator",
297-
"thrust::divides", "thrust::tuple", "thrust::maximum",
309+
"thrust::divides", thrustNamespace() + "tuple", "thrust::maximum",
298310
"thrust::multiplies", "thrust::plus", "cudaDataType_t",
299311
"cudaError_t", "CUresult", "CUdevice", "cudaEvent_t",
300312
"cublasStatus_t", "cuComplex", "cuFloatComplex",
@@ -316,15 +328,15 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
316328
"curandRngType_t", "curandOrdering_t", "cufftHandle", "cufftReal",
317329
"cufftDoubleReal", "cufftComplex", "cufftDoubleComplex",
318330
"cufftResult_t", "cufftResult", "cufftType_t", "cufftType",
319-
"thrust::pair", "CUdeviceptr", "cudaDeviceAttr", "CUmodule",
331+
thrustNamespace() + "pair", "CUdeviceptr", "cudaDeviceAttr", "CUmodule",
320332
"CUjit_option", "CUfunction", "cudaMemcpyKind", "cudaComputeMode",
321333
"__nv_bfloat16", "cooperative_groups::__v1::thread_group",
322334
"cooperative_groups::__v1::thread_block", "libraryPropertyType_t",
323335
"libraryPropertyType", "cudaDataType_t", "cudaDataType",
324336
"cublasComputeType_t", "cublasAtomicsMode_t", "cublasMath_t",
325337
"CUmem_advise_enum", "CUmem_advise", "CUmemorytype",
326-
"CUmemorytype_enum", "thrust::tuple_element",
327-
"thrust::tuple_size", "thrust::zip_iterator",
338+
"CUmemorytype_enum", thrustNamespace() + "tuple_element",
339+
thrustNamespace() + "tuple_size", "thrust::zip_iterator",
328340
"cudaPointerAttributes", "CUpointer_attribute",
329341
"cusolverEigRange_t", "cudaUUID_t", "cusolverDnFunction_t",
330342
"cusolverAlgMode_t", "cusparseIndexType_t", "cusparseFormat_t",

‎clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp‎

Lines changed: 10 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -129,15 +129,9 @@ void CubTypeRule::runRule(
129129

130130
bool CubTypeRule::CanMappingToSyclNativeBinaryOp(StringRef OpTypeName) {
131131
return OpTypeName == "cub::Sum" || OpTypeName == "cub::Max" ||
132-
OpTypeName == "cub::Min";
133-
}
134-
135-
bool CubTypeRule::CanMappingToSyclType(StringRef OpTypeName) {
136-
return CanMappingToSyclNativeBinaryOp(OpTypeName) ||
137-
OpTypeName == "cub::Equality" || OpTypeName == "cub::NullType" ||
138-
139-
// Ignore template arguments, .e.g cub::KeyValuePair<int, int>
140-
OpTypeName.starts_with("cub::KeyValuePair");
132+
OpTypeName == "cub::Min" || OpTypeName == "cuda::std::plus<void>" ||
133+
OpTypeName == "cuda::maximum<void>" ||
134+
OpTypeName == "cuda::minimum<void>";
141135
}
142136

143137
void CubDeviceLevelRule::registerMatcher(ast_matchers::MatchFinder &MF) {
@@ -852,11 +846,12 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
852846
auto processOperatorExpr = [&](const Expr *Obj) {
853847
std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
854848
Obj->getType().getCanonicalType());
849+
std::cout << "1OpType:" << OpType << std::endl;
855850
if (OpType == "cub::Sum" || OpType == "cuda::std::plus<void>") {
856851
OpRepl = MapNames::getClNamespace() + "plus<>()";
857-
} else if (OpType == "cub::Max") {
852+
} else if (OpType == "cub::Max" || OpType == "cuda::maximum<void>") {
858853
OpRepl = MapNames::getClNamespace() + "maximum<>()";
859-
} else if (OpType == "cub::Min") {
854+
} else if (OpType == "cub::Min" || OpType == "cuda::minimum<void>") {
860855
OpRepl = MapNames::getClNamespace() + "minimum<>()";
861856
}
862857
};
@@ -872,8 +867,11 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
872867
return OpRepl;
873868
std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
874869
D->getType().getCanonicalType());
870+
std::cout << "2OpType:" << OpType << std::endl;
875871
if (OpType == "cub::Sum" || OpType == "cub::Max" ||
876-
OpType == "cub::Min" || OpType == "cuda::std::plus<void>") {
872+
OpType == "cub::Min" || OpType == "cuda::std::plus<void>" ||
873+
OpType == "cuda::maximum<void>" ||
874+
OpType == "cuda::minimum<void>") {
877875
ExprAnalysis EA(Operator);
878876
OpRepl = EA.getReplacedString();
879877
}

‎clang/lib/DPCT/RulesLangLib/CUBAPIMigration.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,6 @@ class CubTypeRule : public NamedMigrationRule<CubTypeRule> {
2121
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
2222

2323
static bool CanMappingToSyclNativeBinaryOp(StringRef OpTypeName);
24-
static bool CanMappingToSyclType(StringRef OpTypeName);
2524
};
2625

2726
class CubDeviceLevelRule : public NamedMigrationRule<CubDeviceLevelRule> {

‎clang/lib/DPCT/RulesLangLib/ThrustAPIMigration.cpp‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -54,6 +54,12 @@ void ThrustAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) {
5454
"THRUST_200700___CUDA_ARCH_LIST___NS::system");
5555
};
5656

57+
auto thrustFuncNameCuda129 = [&]() {
58+
return hasAnyName("THRUST_200802_SM___CUDA_ARCH_LIST___NS",
59+
"THRUST_200802_SM___CUDA_ARCH_LIST___NS::detail",
60+
"THRUST_200802_SM___CUDA_ARCH_LIST___NS::system");
61+
};
62+
5763
auto thrustFuncNameCudaCommon = [&]() {
5864
return hasAnyName("thrust", "thrust::detail", "thrust::system", "__4");
5965
};
@@ -118,6 +124,20 @@ void ThrustAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) {
118124
thrustFuncNameCudaCommon())))))))))
119125
.bind("thrustFuncCall"),
120126
this);
127+
} else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 8) {
128+
// For CUDA-12.9
129+
MF.addMatcher(
130+
callExpr(
131+
anyOf(callee(functionDecl(anyOf(
132+
hasDeclContext(namespaceDecl(thrustFuncNameCuda129())),
133+
hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
134+
functionName()))),
135+
callee(unresolvedLookupExpr(
136+
hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
137+
anyOf(thrustFuncNameCuda129(),
138+
thrustFuncNameCudaCommon())))))))))
139+
.bind("thrustFuncCall"),
140+
this);
121141
} else {
122142
// For CUDA SDK versions before CUDA-12.4
123143
MF.addMatcher(

‎clang/lib/Driver/ToolChains/Cuda.cpp‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -199,6 +199,8 @@ bool CudaInstallationDetector::ParseCudaVersionFile(const std::string &FilePath)
199199
Version = CudaVersion::CUDA_126;
200200
} else if (Major == 12 && Minor == 8) {
201201
Version = CudaVersion::CUDA_128;
202+
} else if (Major == 12 && Minor == 9) {
203+
Version = CudaVersion::CUDA_129;
202204
}
203205

204206

@@ -264,7 +266,7 @@ CudaVersion getCudaVersion(uint32_t raw_version) {
264266
return CudaVersion::CUDA_124;
265267
if (raw_version < 12060)
266268
return CudaVersion::CUDA_125;
267-
if (raw_version < 12080)
269+
if (raw_version < 12070)
268270
return CudaVersion::CUDA_126;
269271
if (raw_version < 12090)
270272
return CudaVersion::CUDA_128;
@@ -350,7 +352,8 @@ CudaInstallationDetector::CudaInstallationDetector(
350352

351353
// In decreasing order so we prefer newer versions to older versions.
352354
#ifdef SYCLomatic_CUSTOMIZATION
353-
std::initializer_list<const char *> Versions = {"12.8",
355+
std::initializer_list<const char *> Versions = {"12.9",
356+
"12.8",
354357
"12.6",
355358
"12.5",
356359
"12.4",

‎clang/test/dpct/allocator_syclcompat.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// UNSUPPORTED: cuda-8.0, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8
2-
// UNSUPPORTED: v8.0, v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8
1+
// UNSUPPORTED: cuda-8.0, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8, cuda-12.9
2+
// UNSUPPORTED: v8.0, v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8, v12.9
33
// RUN: dpct --format-range=none --use-syclcompat -out-root %T/allocator_syclcompat %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
44
// RUN: FileCheck --match-full-lines --input-file %T/allocator_syclcompat/allocator_syclcompat.dp.cpp %s
55
// RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/allocator_syclcompat/allocator_syclcompat.dp.cpp -o %T/allocator_syclcompat/allocator_syclcompat.dp.o %}

‎clang/test/dpct/comments.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// UNSUPPORTED: cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8
2-
// UNSUPPORTED: v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8
1+
// UNSUPPORTED: cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.8, cuda-12.9
2+
// UNSUPPORTED: v12.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.8, v12.9
33
// RUN: dpct --format-range=none -out-root %T/comments %s --cuda-include-path="%cuda-path/include" --comments -- -std=c++14 -x cuda --cuda-host-only
44
// RUN: FileCheck %s --match-full-lines --input-file %T/comments/comments.dp.cpp
55
// RUN: %if build_lit %{icpx -c -fsycl %T/comments/comments.dp.cpp -o %T/comments/comments.dp.o %}

0 commit comments

Comments
 (0)