Skip to content

Commit 6af2cd1

Browse files
LU-JOHNLU-JOHN
authored
[SYCLomatic] Migrate CU_JIT* options (#641)
* [SYCLomatic] Migrate CU_JIT* options Signed-off-by: Lu, John <[email protected]> * CU_JIT test not supported with CUDA 8.0 Signed-off-by: Lu, John <[email protected]> * Remove CU_JIT options not available in CUDA 9.2 Signed-off-by: Lu, John <[email protected]> * Migrate CU_JIT_* options to 0 and give a warning. Signed-off-by: Lu, John <[email protected]> * Formatting changes Signed-off-by: Lu, John <[email protected]> * Remove unnecessary code Signed-off-by: Lu, John <[email protected]> --------- Signed-off-by: Lu, John <[email protected]>
1 parent b453e06 commit 6af2cd1

File tree

4 files changed

+155
-1
lines changed

4 files changed

+155
-1
lines changed

‎clang/lib/DPCT/ASTTraversal.cpp‎

Lines changed: 25 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2008,7 +2008,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
20082008
"cufftReal", "cufftDoubleReal", "cufftComplex",
20092009
"cufftDoubleComplex", "cufftResult_t", "cufftResult",
20102010
"cufftType_t", "cufftType", "thrust::pair", "CUdeviceptr",
2011-
"cudaDeviceAttr", "CUmodule", "CUfunction", "cudaMemcpyKind",
2011+
"cudaDeviceAttr", "CUmodule", "CUjit_option",
2012+
"CUfunction", "cudaMemcpyKind",
20122013
"cudaComputeMode", "__nv_bfloat16",
20132014
"cooperative_groups::__v1::thread_block_tile",
20142015
"cooperative_groups::__v1::thread_block",
@@ -3885,6 +3886,29 @@ void FFTEnumsRule::runRule(const MatchFinder::MatchResult &Result) {
38853886

38863887
REGISTER_RULE(FFTEnumsRule, PassKind::PK_Migration)
38873888

3889+
// Rule for CU_JIT enums.
3890+
void CU_JITEnumsRule::registerMatcher(MatchFinder &MF) {
3891+
MF.addMatcher(
3892+
declRefExpr(
3893+
to(enumConstantDecl(matchesName(
3894+
"(CU_JIT_*)"))))
3895+
.bind("CU_JITConstants"),
3896+
this);
3897+
}
3898+
3899+
void CU_JITEnumsRule::runRule(const MatchFinder::MatchResult &Result) {
3900+
if (const DeclRefExpr *DE =
3901+
getNodeAsType<DeclRefExpr>(Result, "CU_JITConstants")) {
3902+
emplaceTransformation(new ReplaceStmt(DE, "0"));
3903+
3904+
report(DE->getBeginLoc(),
3905+
Diagnostics::HOSTALLOCMACRO_NO_MEANING,
3906+
true, DE->getDecl()->getNameAsString());
3907+
}
3908+
}
3909+
3910+
REGISTER_RULE(CU_JITEnumsRule, PassKind::PK_Migration)
3911+
38883912
// Rule for BLAS enums.
38893913
// Migrate BLAS status values to corresponding int values
38903914
// Other BLAS named values are migrated to corresponding named values

‎clang/lib/DPCT/ASTTraversal.h‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -656,6 +656,13 @@ class FFTEnumsRule : public NamedMigrationRule<FFTEnumsRule> {
656656
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
657657
};
658658

659+
/// Migration rule for CU_JIT enums.
660+
class CU_JITEnumsRule : public NamedMigrationRule<CU_JITEnumsRule> {
661+
public:
662+
void registerMatcher(ast_matchers::MatchFinder &MF) override;
663+
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
664+
};
665+
659666
/// Migration rule for BLAS enums.
660667
class BLASEnumsRule : public NamedMigrationRule<BLASEnumsRule> {
661668
public:

‎clang/lib/DPCT/MapNames.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -99,6 +99,7 @@ void MapNames::setExplicitNamespaceMap() {
9999
{"cudaError",
100100
std::make_shared<TypeNameRule>(getDpctNamespace() + "err0",
101101
HelperFeatureEnum::Util_err_types)},
102+
{"CUjit_option", std::make_shared<TypeNameRule>("int")},
102103
{"CUresult", std::make_shared<TypeNameRule>("int")},
103104
{"CUcontext", std::make_shared<TypeNameRule>("int")},
104105
{"CUmodule", std::make_shared<TypeNameRule>(getDpctNamespace() + "kernel_library",

‎clang/test/dpct/cu_jit.cu‎

Lines changed: 122 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,122 @@
1+
// UNSUPPORTED: cuda-8.0
2+
// UNSUPPORTED: v8.0
3+
// RUN: dpct --format-range=none --assume-nd-range-dim=1 -out-root %T/cu_jit %s --cuda-include-path="%cuda-path/include"
4+
// RUN: FileCheck --input-file %T/cu_jit/cu_jit.dp.cpp --match-full-lines %s
5+
6+
#define CU_JIT_NOT_A_CUDA_OPTION 1241
7+
8+
int main() {
9+
int a[40];
10+
int CUvar;
11+
12+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_MAX_REGISTERS is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
13+
//CHECK:a[0] = 0;
14+
a[0] = CU_JIT_MAX_REGISTERS;
15+
16+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_THREADS_PER_BLOCK is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
17+
//CHECK:a[1] = 0;
18+
a[1] = CU_JIT_THREADS_PER_BLOCK;
19+
20+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_WALL_TIME is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
21+
//CHECK:a[2] = 0;
22+
a[2] = CU_JIT_WALL_TIME;
23+
24+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INFO_LOG_BUFFER is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
25+
//CHECK:a[3] = 0;
26+
a[3] = CU_JIT_INFO_LOG_BUFFER;
27+
28+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
29+
//CHECK:a[4] = 0;
30+
a[4] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
31+
32+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_ERROR_LOG_BUFFER is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
33+
//CHECK:a[5] = 0;
34+
a[5] = CU_JIT_ERROR_LOG_BUFFER;
35+
36+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
37+
//CHECK:a[6] = 0;
38+
a[6] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
39+
40+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_OPTIMIZATION_LEVEL is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
41+
//CHECK:a[7] = 0;
42+
a[7] = CU_JIT_OPTIMIZATION_LEVEL;
43+
44+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_TARGET_FROM_CUCONTEXT is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
45+
//CHECK:a[8] = 0;
46+
a[8] = CU_JIT_TARGET_FROM_CUCONTEXT;
47+
48+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_TARGET is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
49+
//CHECK:a[9] = 0;
50+
a[9] = CU_JIT_TARGET;
51+
52+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_FALLBACK_STRATEGY is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
53+
//CHECK:a[10] = 0;
54+
a[10] = CU_JIT_FALLBACK_STRATEGY;
55+
56+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_GENERATE_DEBUG_INFO is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
57+
//CHECK:a[11] = 0;
58+
a[11] = CU_JIT_GENERATE_DEBUG_INFO;
59+
60+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_LOG_VERBOSE is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
61+
//CHECK:a[12] = 0;
62+
a[12] = CU_JIT_LOG_VERBOSE;
63+
64+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_GENERATE_LINE_INFO is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
65+
//CHECK:a[13] = 0;
66+
a[13] = CU_JIT_GENERATE_LINE_INFO;
67+
68+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_CACHE_MODE is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
69+
//CHECK:a[14] = 0;
70+
a[14] = CU_JIT_CACHE_MODE;
71+
72+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_NEW_SM3X_OPT is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
73+
//CHECK:a[15] = 0;
74+
a[15] = CU_JIT_NEW_SM3X_OPT;
75+
76+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_FAST_COMPILE is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
77+
//CHECK:a[16] = 0;
78+
a[16] = CU_JIT_FAST_COMPILE;
79+
80+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_NUM_OPTIONS is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
81+
//CHECK:a[17] = 0;
82+
a[17] = CU_JIT_NUM_OPTIONS;
83+
84+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_CACHE_OPTION_NONE is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
85+
//CHECK:a[18] = 0;
86+
a[18] = CU_JIT_CACHE_OPTION_NONE;
87+
88+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_CACHE_OPTION_CG is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
89+
//CHECK:a[19] = 0;
90+
a[19] = CU_JIT_CACHE_OPTION_CG;
91+
92+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_CACHE_OPTION_CA is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
93+
//CHECK:a[20] = 0;
94+
a[20] = CU_JIT_CACHE_OPTION_CA;
95+
96+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INPUT_CUBIN is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
97+
//CHECK:a[21] = 0;
98+
a[21] = CU_JIT_INPUT_CUBIN;
99+
100+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INPUT_PTX is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
101+
//CHECK:a[22] = 0;
102+
a[22] = CU_JIT_INPUT_PTX;
103+
104+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INPUT_FATBINARY is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
105+
//CHECK:a[23] = 0;
106+
a[23] = CU_JIT_INPUT_FATBINARY;
107+
108+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INPUT_OBJECT is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
109+
//CHECK:a[24] = 0;
110+
a[24] = CU_JIT_INPUT_OBJECT;
111+
112+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_INPUT_LIBRARY is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
113+
//CHECK:a[25] = 0;
114+
a[25] = CU_JIT_INPUT_LIBRARY;
115+
116+
//CHECK:DPCT1048:{{[0-9]+}}: The original value CU_JIT_NUM_INPUT_TYPES is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
117+
//CHECK:a[26] = 0;
118+
a[26] = CU_JIT_NUM_INPUT_TYPES;
119+
120+
//CHECK:a[27] = CU_JIT_NOT_A_CUDA_OPTION;
121+
a[27] = CU_JIT_NOT_A_CUDA_OPTION;
122+
}

0 commit comments

Comments
 (0)