[SYCLomatic] Emit warning for dynamic parallelism (#2251)

Signed-off-by: Jiang, Zhiwei <[email protected]>

Author: zhiweij1

clang/lib/DPCT/ASTTraversal.cpp

@@ -8407,6 +8407,12 @@ void KernelCallRule::runRule(
     auto FD = getAssistNodeAsType<FunctionDecl>(Result, "callContext");
     if (!FD)
       return;
+    if (FD->hasAttr<CUDAGlobalAttr>()) {
+      report(KCall->getBeginLoc(), Diagnostics::NOT_SUPPORT_DYN_PARALLEL,
+             false);
+      return;
+    }
+
     const auto &SM = (*Result.Context).getSourceManager();
 
     if (SM.isMacroArgExpansion(KCall->getCallee()->getBeginLoc())) {

clang/lib/DPCT/Diagnostics.inc

@@ -286,6 +286,8 @@ DEF_WARNING(NOT_DEVICE_COPYABLE, 1128, MEDIUM_LEVEL, "The type \"%0\" is not dev
 DEF_COMMENT(NOT_DEVICE_COPYABLE, 1128, MEDIUM_LEVEL, "The type \"{0}\" is not device copyable for {1}. It is used in the SYCL kernel, please rewrite the code.")
 DEF_WARNING(NOT_DEVICE_COPYABLE_ADD_SPECIALIZATION, 1129, MEDIUM_LEVEL, "The type \"%0\" is used in the SYCL kernel, but it is not device copyable. The sycl::is_device_copyable specialization has been added for this type. Please review the code.")
 DEF_COMMENT(NOT_DEVICE_COPYABLE_ADD_SPECIALIZATION, 1129, MEDIUM_LEVEL, "The type \"{0}\" is used in the SYCL kernel, but it is not device copyable. The sycl::is_device_copyable specialization has been added for this type. Please review the code.")
+DEF_WARNING(NOT_SUPPORT_DYN_PARALLEL, 1130, HIGH_LEVEL, "SYCL 2020 standard does not support dynamic parallelism (launching kernel in device code). Please rewrite the code.")
+DEF_COMMENT(NOT_SUPPORT_DYN_PARALLEL, 1130, HIGH_LEVEL, "SYCL 2020 standard does not support dynamic parallelism (launching kernel in device code). Please rewrite the code.")
 // clang-format on
 
 #undef DEF_COMMENT

clang/test/dpct/nestedqueue.cu

@@ -74,3 +74,19 @@ __host__ void foo4(){
   cublasIsamax(handle, n, x_S, incx, result);
 }
 
+__global__ void childKernel() {}
+__global__ void parentKernel() {
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1130:{{[0-9]+}}: SYCL 2020 standard does not support dynamic parallelism (launching kernel in device code). Please rewrite the code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: childKernel<<<1, 1>>>();
+  childKernel<<<1, 1>>>();
+}
+void foo5() {
+  // CHECK: dpct::get_out_of_order_queue().parallel_for(
+  // CHECK-NEXT:   sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
+  // CHECK-NEXT:   [=](sycl::nd_item<3> item_ct1) {
+  // CHECK-NEXT:     parentKernel();
+  // CHECK-NEXT:   });
+  parentKernel<<<1, 1>>>();
+}