Skip to content

Commit 3e1326e

Browse files
zhiweij1zhiweij1
authored
[SYCLomatic] Add 12.9 header file migration support (#2822)
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent e744c15 commit 3e1326e

39 files changed

Lines changed: 136 additions & 212 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: 19 additions & 6 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 or later
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,16 @@ 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",
320-
"CUjit_option", "CUfunction", "cudaMemcpyKind", "cudaComputeMode",
321-
"__nv_bfloat16", "cooperative_groups::__v1::thread_group",
331+
thrustNamespace() + "pair", "CUdeviceptr", "cudaDeviceAttr",
332+
"CUmodule", "CUjit_option", "CUfunction", "cudaMemcpyKind",
333+
"cudaComputeMode", "__nv_bfloat16",
334+
"cooperative_groups::__v1::thread_group",
322335
"cooperative_groups::__v1::thread_block", "libraryPropertyType_t",
323336
"libraryPropertyType", "cudaDataType_t", "cudaDataType",
324337
"cublasComputeType_t", "cublasAtomicsMode_t", "cublasMath_t",
325338
"CUmem_advise_enum", "CUmem_advise", "CUmemorytype",
326-
"CUmemorytype_enum", "thrust::tuple_element",
327-
"thrust::tuple_size", "thrust::zip_iterator",
339+
"CUmemorytype_enum", thrustNamespace() + "tuple_element",
340+
thrustNamespace() + "tuple_size", "thrust::zip_iterator",
328341
"cudaPointerAttributes", "CUpointer_attribute",
329342
"cusolverEigRange_t", "cudaUUID_t", "cusolverDnFunction_t",
330343
"cusolverAlgMode_t", "cusparseIndexType_t", "cusparseFormat_t",

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

Lines changed: 8 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) {
@@ -854,9 +848,9 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
854848
Obj->getType().getCanonicalType());
855849
if (OpType == "cub::Sum" || OpType == "cuda::std::plus<void>") {
856850
OpRepl = MapNames::getClNamespace() + "plus<>()";
857-
} else if (OpType == "cub::Max") {
851+
} else if (OpType == "cub::Max" || OpType == "cuda::maximum<void>") {
858852
OpRepl = MapNames::getClNamespace() + "maximum<>()";
859-
} else if (OpType == "cub::Min") {
853+
} else if (OpType == "cub::Min" || OpType == "cuda::minimum<void>") {
860854
OpRepl = MapNames::getClNamespace() + "minimum<>()";
861855
}
862856
};
@@ -873,7 +867,9 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
873867
std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
874868
D->getType().getCanonicalType());
875869
if (OpType == "cub::Sum" || OpType == "cub::Max" ||
876-
OpType == "cub::Min" || OpType == "cuda::std::plus<void>") {
870+
OpType == "cub::Min" || OpType == "cuda::std::plus<void>" ||
871+
OpType == "cuda::maximum<void>" ||
872+
OpType == "cuda::minimum<void>") {
877873
ExprAnalysis EA(Operator);
878874
OpRepl = EA.getReplacedString();
879875
}

‎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: 35 additions & 107 deletions
Original file line numberDiff line numberDiff line change
@@ -22,116 +22,44 @@ void ThrustAPIRule::registerMatcher(ast_matchers::MatchFinder &MF) {
2222
// API register
2323
auto functionName = [&]() { return hasAnyName("on"); };
2424

25-
// THRUST_200302___CUDA_ARCH_LIST___NS is newly imported inline
26-
// namespace by thrust library in CUDA header file 12.4.
27-
auto thrustFuncNameCuda124 = [&]() {
28-
return hasAnyName("THRUST_200302___CUDA_ARCH_LIST___NS",
29-
"THRUST_200302___CUDA_ARCH_LIST___NS::detail",
30-
"THRUST_200302___CUDA_ARCH_LIST___NS::system");
31-
};
32-
33-
// THRUST_200400___CUDA_ARCH_LIST___NS is newly imported inline
34-
// namespace by thrust library in CUDA header file 12.5.
35-
auto thrustFuncNameCuda125 = [&]() {
36-
return hasAnyName("THRUST_200400___CUDA_ARCH_LIST___NS",
37-
"THRUST_200400___CUDA_ARCH_LIST___NS::detail",
38-
"THRUST_200400___CUDA_ARCH_LIST___NS::system");
39-
};
40-
41-
// THRUST_200500___CUDA_ARCH_LIST___NS is newly imported inline
42-
// namespace by thrust library in CUDA header file 12.6.
43-
auto thrustFuncNameCuda126 = [&]() {
44-
return hasAnyName("THRUST_200500___CUDA_ARCH_LIST___NS",
45-
"THRUST_200500___CUDA_ARCH_LIST___NS::detail",
46-
"THRUST_200500___CUDA_ARCH_LIST___NS::system");
47-
};
48-
49-
// THRUST_200700___CUDA_ARCH_LIST___NS is newly imported inline
50-
// namespace by thrust library in CUDA header file 12.6.
51-
auto thrustFuncNameCuda128 = [&]() {
52-
return hasAnyName("THRUST_200700___CUDA_ARCH_LIST___NS",
53-
"THRUST_200700___CUDA_ARCH_LIST___NS::detail",
54-
"THRUST_200700___CUDA_ARCH_LIST___NS::system");
55-
};
56-
57-
auto thrustFuncNameCudaCommon = [&]() {
58-
return hasAnyName("thrust", "thrust::detail", "thrust::system", "__4");
59-
};
60-
6125
int ThrustMajorVersion = ThrustVersion / 100000;
6226
int ThrustMinorVersion = ThrustVersion / 100 % 1000;
6327

64-
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 3) {
65-
// For CUDA-12.4
66-
MF.addMatcher(
67-
callExpr(
68-
anyOf(callee(functionDecl(anyOf(
69-
hasDeclContext(namespaceDecl(thrustFuncNameCuda124())),
70-
hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
71-
functionName()))),
72-
callee(unresolvedLookupExpr(
73-
hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
74-
anyOf(thrustFuncNameCuda124(),
75-
thrustFuncNameCudaCommon())))))))))
76-
.bind("thrustFuncCall"),
77-
this);
78-
79-
} else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 4) {
80-
// For CUDA-12.5
81-
MF.addMatcher(
82-
callExpr(
83-
anyOf(callee(functionDecl(anyOf(
84-
hasDeclContext(namespaceDecl(thrustFuncNameCuda125())),
85-
hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
86-
functionName()))),
87-
callee(unresolvedLookupExpr(
88-
hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
89-
anyOf(thrustFuncNameCuda125(),
90-
thrustFuncNameCudaCommon())))))))))
91-
.bind("thrustFuncCall"),
92-
this);
93-
} else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 5) {
94-
// For CUDA-12.6
95-
MF.addMatcher(
96-
callExpr(
97-
anyOf(callee(functionDecl(anyOf(
98-
hasDeclContext(namespaceDecl(thrustFuncNameCuda126())),
99-
hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
100-
functionName()))),
101-
callee(unresolvedLookupExpr(
102-
hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
103-
anyOf(thrustFuncNameCuda126(),
104-
thrustFuncNameCudaCommon())))))))))
105-
.bind("thrustFuncCall"),
106-
this);
107-
} else if (ThrustMajorVersion == 2 && ThrustMinorVersion == 7) {
108-
// For CUDA-12.8
109-
MF.addMatcher(
110-
callExpr(
111-
anyOf(callee(functionDecl(anyOf(
112-
hasDeclContext(namespaceDecl(thrustFuncNameCuda128())),
113-
hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
114-
functionName()))),
115-
callee(unresolvedLookupExpr(
116-
hasAnyDeclaration(namedDecl(hasDeclContext(namespaceDecl(
117-
anyOf(thrustFuncNameCuda128(),
118-
thrustFuncNameCudaCommon())))))))))
119-
.bind("thrustFuncCall"),
120-
this);
121-
} else {
122-
// For CUDA SDK versions before CUDA-12.4
123-
MF.addMatcher(
124-
callExpr(
125-
anyOf(callee(functionDecl(anyOf(
126-
hasDeclContext(namespaceDecl(thrustFuncNameCudaCommon())),
127-
128-
functionName()))),
129-
callee(unresolvedLookupExpr(
130-
hasAnyDeclaration(namedDecl(hasDeclContext(
131-
namespaceDecl(thrustFuncNameCudaCommon()))))))))
132-
.bind("thrustFuncCall"),
133-
this);
134-
}
28+
auto thrustFuncName = [&]() {
29+
#define COMMON "thrust", "thrust::detail", "thrust::system", "__4"
30+
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 3)
31+
return hasAnyName("THRUST_200302___CUDA_ARCH_LIST___NS",
32+
"THRUST_200302___CUDA_ARCH_LIST___NS::detail",
33+
"THRUST_200302___CUDA_ARCH_LIST___NS::system", COMMON);
34+
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 4)
35+
return hasAnyName("THRUST_200400___CUDA_ARCH_LIST___NS",
36+
"THRUST_200400___CUDA_ARCH_LIST___NS::detail",
37+
"THRUST_200400___CUDA_ARCH_LIST___NS::system", COMMON);
38+
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 5)
39+
return hasAnyName("THRUST_200500___CUDA_ARCH_LIST___NS",
40+
"THRUST_200500___CUDA_ARCH_LIST___NS::detail",
41+
"THRUST_200500___CUDA_ARCH_LIST___NS::system", COMMON);
42+
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 7)
43+
return hasAnyName("THRUST_200700___CUDA_ARCH_LIST___NS",
44+
"THRUST_200700___CUDA_ARCH_LIST___NS::detail",
45+
"THRUST_200700___CUDA_ARCH_LIST___NS::system", COMMON);
46+
if (ThrustMajorVersion == 2 && ThrustMinorVersion == 8)
47+
return hasAnyName("THRUST_200802_SM___CUDA_ARCH_LIST___NS",
48+
"THRUST_200802_SM___CUDA_ARCH_LIST___NS::detail",
49+
"THRUST_200802_SM___CUDA_ARCH_LIST___NS::system",
50+
COMMON);
51+
return hasAnyName(COMMON);
52+
#undef COMMON
53+
};
54+
55+
MF.addMatcher(
56+
callExpr(anyOf(callee(functionDecl(
57+
anyOf(hasDeclContext(namespaceDecl(thrustFuncName())),
58+
functionName()))),
59+
callee(unresolvedLookupExpr(hasAnyDeclaration(namedDecl(
60+
hasDeclContext(namespaceDecl(thrustFuncName()))))))))
61+
.bind("thrustFuncCall"),
62+
this);
13563

13664
// THRUST_STATIC_ASSERT macro register
13765
MF.addMatcher(staticAssertDecl(isExpandedFromMacro("THRUST_STATIC_ASSERT"))

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

Lines changed: 7 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -199,9 +199,10 @@ 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

204-
205206
if (Version != CudaVersion::UNKNOWN) {
206207
IsVersionSupported = true;
207208
return true;
@@ -264,7 +265,7 @@ CudaVersion getCudaVersion(uint32_t raw_version) {
264265
return CudaVersion::CUDA_124;
265266
if (raw_version < 12060)
266267
return CudaVersion::CUDA_125;
267-
if (raw_version < 12080)
268+
if (raw_version < 12070)
268269
return CudaVersion::CUDA_126;
269270
if (raw_version < 12090)
270271
return CudaVersion::CUDA_128;
@@ -350,31 +351,10 @@ CudaInstallationDetector::CudaInstallationDetector(
350351

351352
// In decreasing order so we prefer newer versions to older versions.
352353
#ifdef SYCLomatic_CUSTOMIZATION
353-
std::initializer_list<const char *> Versions = {"12.8",
354-
"12.6",
355-
"12.5",
356-
"12.4",
357-
"12.3",
358-
"12.2",
359-
"12.1",
360-
"12.0",
361-
"11.8",
362-
"11.7",
363-
"11.6",
364-
"11.5",
365-
"11.4",
366-
"11.3",
367-
"11.2",
368-
"11.1",
369-
"10.2",
370-
"10.1",
371-
"10.0",
372-
"9.2",
373-
"9.1",
374-
"9.0",
375-
"8.0",
376-
"7.5",
377-
"7.0"};
354+
std::initializer_list<const char *> Versions = {
355+
"12.9", "12.8", "12.6", "12.5", "12.4", "12.3", "12.2", "12.1", "12.0",
356+
"11.8", "11.7", "11.6", "11.5", "11.4", "11.3", "11.2", "11.1", "10.2",
357+
"10.1", "10.0", "9.2", "9.1", "9.0", "8.0", "7.5", "7.0"};
378358
#else
379359
std::initializer_list<const char *> Versions = {
380360
"11.4", "11.3", "11.2", "11.1", "10.2", "10.1", "10.0",

‎clang/test/dpct/NVTX/NVTX-linux.cu‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
11
// UNSUPPORTED: system-windows
2+
// UNSUPPORTED: cuda-12.9
23
// RUN: dpct --format-range=none -in-root %S -out-root %T %S/NVTX-linux.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
34
// RUN: FileCheck --input-file %T/NVTX-linux.dp.cpp --match-full-lines %s
45

‎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 %}

0 commit comments

Comments
 (0)