[SYCLomatic] Emit warning message for replacing macro/constant variab…

…le by literal in local_accessor range (#315)

Signed-off-by: Tang, Jiajun [email protected]

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);

clang/lib/DPCT/AnalysisInfo.h

@@ -2282,6 +2282,7 @@ class CtTypeInfo {
   }
   std::set<HelperFeatureEnum> getHelperFeatureSet() { return HelperFeatureSet; }
   inline bool containSizeofType() { return ContainSizeofType; }
+  inline std::vector<std::string> 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<TemplateDependentStringInfo> TDSI;
   std::set<HelperFeatureEnum> HelperFeatureSet;
   bool ContainSizeofType = false;
+  std::vector<std::string> ArraySizeOriginExprs{};
 };
 
 // variable info includes name, type and location.

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.")

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<int, 1> counter_acc_ct1(sycl::range<1>(32/*WARP_SIZE*/), cgh);
   //CHECK-NEXT:    dpct::access_wrapper<int *> dOut_acc_ct0(dOut, cgh);
   //CHECK-EMPTY:

clang/test/dpct/macro_test.cu

@@ -1153,6 +1153,15 @@ template<class T1, class T2, int N> __global__ void foo31();
 #define FOO31(DIMS) foo31<unsigned int, float, DIMS><<<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<double, 2> red_acc_acc_ct1(
 //CHECK-NEXT:         sycl::range<2>(8 /*8*/, 8 /*BLOCK_PAIR / SIMD_SIZE*/), cgh);
 

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: */

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: */

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<T, 2> 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<T, 3> s3_acc_ct1(sycl::range<3>(64/*size * 2*/, 128/*size * 4*/, 32/*size*/), cgh);
   // CHECK-NEXT:     dpct::access_wrapper<T *> 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<int, 0> 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<int, 1> 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<int, 2> 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<int, 3> 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: