[SYCLomatic] Emit warning for shared local memory if the data type defined by using or typedef

clang/lib/DPCT/ASTTraversal.cpp

@@ -9592,6 +9592,33 @@ void MemVarAnalysisRule::processTypeDeclaredLocal(
   }
 }
 
+#define TYPE_CAST(Target) dyn_cast<Target>(T)
+SourceLocation getTypedefOrUsingLoc(QualType QT) {
+  const Type *T = QT.getTypePtr();
+  switch (T->getTypeClass()) {
+  case Type::TypeClass::IncompleteArray:
+    return getTypedefOrUsingLoc(
+        TYPE_CAST(IncompleteArrayType)->getElementType());
+  case Type::TypeClass::ConstantArray:
+    return getTypedefOrUsingLoc(TYPE_CAST(ConstantArrayType)->getElementType());
+  case Type::TypeClass::Pointer:
+    return getTypedefOrUsingLoc(TYPE_CAST(PointerType)->getPointeeType());
+  case Type::TypeClass::Elaborated:
+    return getTypedefOrUsingLoc(TYPE_CAST(ElaboratedType)->desugar());
+  case Type::TypeClass::Typedef:
+    return TYPE_CAST(TypedefType)
+        ->getDecl()
+        ->getTypeSourceInfo()
+        ->getTypeLoc()
+        .getBeginLoc();
+  case Type::TypeClass::Using:
+    return TYPE_CAST(clang::UsingType)->getFoundDecl()->getBeginLoc();
+  default:
+    return SourceLocation();
+  }
+}
+#undef TYPE_CAST
+
 void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) {
   std::string CanonicalType;
   if (auto MemVar = getAssistNodeAsType<VarDecl>(Result, "var")) {
@@ -9621,6 +9648,12 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) {
       emplaceTransformation(ReplaceVarDecl::getVarDeclReplacement(
           MemVar, Info->getDeclarationReplacement(MemVar)));
     }
+    if (MemVar->hasAttr<CUDASharedAttr>()) {
+      SourceLocation SL = getTypedefOrUsingLoc(MemVar->getType());
+      if (SL.isValid()) {
+        report(SL, Diagnostics::MOVE_TYPE_DEFINITION, true);
+      }
+    }
     return;
   }
   auto MemVarRef = getNodeAsType<DeclRefExpr>(Result, "used");
@@ -9649,6 +9682,12 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) {
         }
       }
     }
+    if (VD->hasAttr<CUDASharedAttr>()) {
+      SourceLocation SL = getTypedefOrUsingLoc(VD->getType());
+      if (SL.isValid()) {
+        report(SL, Diagnostics::MOVE_TYPE_DEFINITION, true);
+      }
+    }
   }
 }
 

clang/lib/DPCT/Diagnostics.inc

@@ -274,6 +274,8 @@ DEF_WARNING(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in
 DEF_COMMENT(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in template declare, it may cause template function or class redefinition, please adjust the code.")
 DEF_WARNING(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.")
 DEF_COMMENT(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.")
+DEF_WARNING(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used in other places. You may need adjust the type definition location.")
+DEF_COMMENT(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used in other places. You may need adjust the type definition location.")
 // clang-format on
 
 #undef DEF_COMMENT

clang/test/dpct/kernel_without_name.cu

@@ -1,6 +1,6 @@
 // RUN: dpct --format-range=none --usm-level=none -out-root %T/kernel_without_name %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
 // RUN: FileCheck --input-file %T/kernel_without_name/kernel_without_name.dp.cpp --match-full-lines %s
-// RUN: %if build_lit %{icpx -c -fsycl %T/kernel_without_name/kernel_without_name.dp.cpp -o %T/kernel_without_name/kernel_without_name.dp.o %}
+// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/kernel_without_name/kernel_without_name.dp.cpp -o %T/kernel_without_name/kernel_without_name.dp.o %}
 
 __global__ void testKernel(int L, int M, int N);
 
@@ -277,3 +277,18 @@ void run_foo7(T *a, const T *b, const T *c, const T *d, const T *e, const int f,
     foo_kernel6<<<grid, block, 0, stream>>>(a, b, c, d, e, f, g, h);
   }
 }
+
+#ifndef BUILD_TEST
+template <typename T> struct kernel_type_t {
+  using Type = T;
+};
+
+template <typename T> __global__ void foo_kernel7() {
+  //CHECK:/*
+  //CHECK-NEXT:DPCT1124:{{[0-9]+}}: This type is used in other places. You may need adjust the type definition location.
+  //CHECK-NEXT:*/
+  //CHECK-NEXT:using Tk = typename kernel_type_t<T>::Type;
+  using Tk = typename kernel_type_t<T>::Type;
+  __shared__ Tk mem[256];
+}
+#endif