From dcdce3db87676997a214ecd525d7414d19d23138 Mon Sep 17 00:00:00 2001 From: tangjj11 Date: Fri, 11 Nov 2022 17:17:02 +0800 Subject: [PATCH] [SYCLomatic] Emit warning message for replacing macro/constant variable by literal in local_accessor range (#315) Signed-off-by: Tang, Jiajun jiajun.tang@intel.com --- clang/lib/DPCT/AnalysisInfo.cpp | 8 ++++++ clang/lib/DPCT/AnalysisInfo.h | 5 +++- clang/lib/DPCT/Diagnostics.inc | 6 +++++ clang/test/dpct/curand-device.cu | 3 +++ clang/test/dpct/macro_test.cu | 9 +++++++ clang/test/dpct/sharedmem-dim-noUSM.cu | 3 +++ clang/test/dpct/sharedmem-dim.cu | 3 +++ clang/test/dpct/sharedmem_var_static.cu | 33 +++++++++++++++++++++++++ 8 files changed, 69 insertions(+), 1 deletion(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index c8fde4a63e23..b7d184a36002 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -3160,6 +3160,14 @@ void MemVarInfo::appendAccessorOrPointerDecl(const std::string &ExternMemSize, OS << "cgh)"; OS << ";"; StmtWithWarning AccDecl(OS.str()); + for (const auto &OriginExpr : getType()->getArraySizeOriginExprs()) { + DiagnosticsUtils::report(getFilePath(), getOffset(), + Diagnostics::MACRO_EXPR_REPLACED, false, false, + OriginExpr); + AccDecl.Warnings.push_back( + DiagnosticsUtils::getWarningTextAndUpdateUniqueID( + Diagnostics::MACRO_EXPR_REPLACED, OriginExpr)); + } if ((isExtern() && ExternEmitWarning) || getType()->containSizeofType()) { DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::SIZEOF_WARNING, false, false); diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 68e440eca2ae..a67258ddd5c2 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -2282,6 +2282,7 @@ class CtTypeInfo { } std::set getHelperFeatureSet() { return HelperFeatureSet; } inline bool containSizeofType() { return ContainSizeofType; } + inline std::vector getArraySizeOriginExprs() { return ArraySizeOriginExprs; } private: // For ConstantArrayType, size in generated code is folded as an integer. @@ -2301,8 +2302,9 @@ class CtTypeInfo { if (TL.getSizeExpr()->getStmtClass() == Stmt::IntegerLiteralClass && TL.getSizeExpr()->getBeginLoc().isFileID()) return toString(TL.getTypePtr()->getSize(), 10, false, false); + ArraySizeOriginExprs.push_back(getStmtSpelling(TL.getSizeExpr())); return buildString(toString(TL.getTypePtr()->getSize(), 10, false, false), - "/*", getStmtSpelling(TL.getSizeExpr()), "*/"); + "/*", ArraySizeOriginExprs.back(), "*/"); } // Get original array size expression. @@ -2349,6 +2351,7 @@ class CtTypeInfo { std::shared_ptr TDSI; std::set HelperFeatureSet; bool ContainSizeofType = false; + std::vector ArraySizeOriginExprs{}; }; // variable info includes name, type and location. diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index 8810f6d58d2c..1598d26b2474 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -211,3 +211,9 @@ DEF_WARNING(PRIMITIVE_WORKSPACE, 1097, "The function \"%0\" may require the work DEF_COMMENT(PRIMITIVE_WORKSPACE, 1097, "The function \"{0}\" may require the workspace used to save intermediate results from function \"{1}\". By default, a workspace from engine_ext is selected according to the source data pointer, but this may be incorrect and cause a workspace data race. You may need to rewrite this code.") DEF_WARNING(MATH_EMULATION_EXPRESSION, 1098, "The %1 expression is used instead of the %0 call. These two expressions do not provide the exact same functionality. Check the generated code for potential precision and/or performance issues.") DEF_COMMENT(MATH_EMULATION_EXPRESSION, 1098, "The {1} expression is used instead of the {0} call. These two expressions do not provide the exact same functionality. Check the generated code for potential precision and/or performance issues.") +DEF_WARNING(FFT_DIRECTION_AND_PLACEMENT, 1099, "Verify if the default value of the direction and placement used in the function \"%0\" is correct.") +DEF_COMMENT(FFT_DIRECTION_AND_PLACEMENT, 1099, "Verify if the default value of the direction and placement used in the function \"{0}\" is correct.") +DEF_WARNING(FFT_EXTERNAL_WORKSPACE, 1100, "Currently the DFT external workspace feature in the Intel(R) oneAPI Math Kernel Library (oneMKL) is only supported for GPU devices. Use the internal workspace if your code should run on non-GPU devices.") +DEF_COMMENT(FFT_EXTERNAL_WORKSPACE, 1100, "Currently the DFT external workspace feature in the Intel(R) oneAPI Math Kernel Library (oneMKL) is only supported for GPU devices. Use the internal workspace if your code should run on non-GPU devices.") +DEF_WARNING(MACRO_EXPR_REPLACED, 1101, "'%0' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct.") +DEF_COMMENT(MACRO_EXPR_REPLACED, 1101, "'{0}' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct.") diff --git a/clang/test/dpct/curand-device.cu b/clang/test/dpct/curand-device.cu index a471f3b8dbb0..cae05d85c3ee 100644 --- a/clang/test/dpct/curand-device.cu +++ b/clang/test/dpct/curand-device.cu @@ -76,6 +76,9 @@ int main(int argc, char **argv) { int *dOut; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { + //CHECK-NEXT: /* + //CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'WARP_SIZE' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + //CHECK-NEXT: */ //CHECK-NEXT: sycl::local_accessor counter_acc_ct1(sycl::range<1>(32/*WARP_SIZE*/), cgh); //CHECK-NEXT: dpct::access_wrapper dOut_acc_ct0(dOut, cgh); //CHECK-EMPTY: diff --git a/clang/test/dpct/macro_test.cu b/clang/test/dpct/macro_test.cu index 3600618ac468..c76608a90096 100644 --- a/clang/test/dpct/macro_test.cu +++ b/clang/test/dpct/macro_test.cu @@ -1153,6 +1153,15 @@ template __global__ void foo31(); #define FOO31(DIMS) foo31<<<1,1>>>(); //CHECK: q_ct1.submit([&](sycl::handler &cgh) { +//CHECK-NEXT: /* +//CHECK-NEXT: DPCT1101:{{[0-9]+}}: '8' expression was replaced with a value. Modify the code to +//CHECK-NEXT: use original expression, provided in comments, if it is correct. +//CHECK-NEXT: */ +//CHECK-NEXT: /* +//CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'BLOCK_PAIR / SIMD_SIZE' expression was replaced with a +//CHECK-NEXT: value. Modify the code to use original expression, provided in comments, +//CHECK-NEXT: if it is correct. +//CHECK-NEXT: */ //CHECK-NEXT: sycl::local_accessor red_acc_acc_ct1( //CHECK-NEXT: sycl::range<2>(8 /*8*/, 8 /*BLOCK_PAIR / SIMD_SIZE*/), cgh); diff --git a/clang/test/dpct/sharedmem-dim-noUSM.cu b/clang/test/dpct/sharedmem-dim-noUSM.cu index c38f3f22e2ba..6d3a522c1141 100644 --- a/clang/test/dpct/sharedmem-dim-noUSM.cu +++ b/clang/test/dpct/sharedmem-dim-noUSM.cu @@ -83,6 +83,9 @@ __global__ void kernel3() { } void foo3() { + // CHECK: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'sizeof(float3) * 3' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK: /* // CHECK-NEXT: DPCT1083:{{[0-9]+}}: The size of local memory in the migrated code may be different from the original code. Check that the allocated memory size in the migrated code is correct. // CHECK-NEXT: */ diff --git a/clang/test/dpct/sharedmem-dim.cu b/clang/test/dpct/sharedmem-dim.cu index f35821eea5f2..2b518cf08cb1 100644 --- a/clang/test/dpct/sharedmem-dim.cu +++ b/clang/test/dpct/sharedmem-dim.cu @@ -85,6 +85,9 @@ __global__ void kernel3() { } void foo3() { + // CHECK: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'sizeof(float3) * 3' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK: /* // CHECK-NEXT: DPCT1083:{{[0-9]+}}: The size of local memory in the migrated code may be different from the original code. Check that the allocated memory size in the migrated code is correct. // CHECK-NEXT: */ diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index 6330650f232c..d6d80933d076 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -82,7 +82,22 @@ void testTemplate() { // CHECK: dpct::get_default_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 2' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 4' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK-NEXT: sycl::local_accessor s_acc_ct1(sycl::range<2>(64/*size * 2*/, 128/*size * 4*/), cgh); + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 2' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 4' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK-NEXT: sycl::local_accessor s3_acc_ct1(sycl::range<3>(64/*size * 2*/, 128/*size * 4*/, 32/*size*/), cgh); // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); // CHECK-EMPTY: @@ -119,6 +134,9 @@ int main(void) { // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: sycl::local_accessor a0_acc_ct1(cgh); + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK-NEXT: sycl::local_accessor s_acc_ct1(sycl::range<1>(64/*size*/), cgh); // CHECK-NEXT: auto d_d_acc_ct0 = dpct::get_access(d_d, cgh); // CHECK-EMPTY: @@ -133,7 +151,22 @@ int main(void) { // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 2' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 4' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK-NEXT: sycl::local_accessor s_acc_ct1(sycl::range<2>(64/*size * 2*/, 128/*size * 4*/), cgh); + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 2' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size * 4' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size' expression was replaced with a value. Modify the code to use original expression, provided in comments, if it is correct. + // CHECK-NEXT: */ // CHECK-NEXT: sycl::local_accessor s3_acc_ct1(sycl::range<3>(64/*size * 2*/, 128/*size * 4*/, 32/*size*/), cgh); // CHECK-NEXT: auto d_d_acc_ct0 = dpct::get_access(d_d, cgh); // CHECK-EMPTY: