[SYCL] Diagnose attempt to pass pointer to VLA as kernel arg (#2441)

VLA as well as pointers to VLA require additional AST transformation to make
emission of corresponding type in LLVM IR possible. Without this transformation
Codegen just crashes.
Implementing this transformation is not reasonable whereas SYCL standard
doesn't allow VLA in device code, so emit an error in case if pointer to
VLA is passed as kernel argument.

Author: Fznamznon

clang/lib/Sema/SemaSYCL.cpp

@@ -1298,6 +1298,18 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     return isValid();
   }
 
+  bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
+    while (FieldTy->isAnyPointerType()) {
+      FieldTy = QualType{FieldTy->getPointeeOrArrayElementType(), 0};
+      if (FieldTy->isVariableArrayType()) {
+        Diag.Report(FD->getLocation(), diag::err_vla_unsupported);
+        IsInvalid = true;
+        break;
+      }
+    }
+    return isValid();
+  }
+
   bool handleOtherType(FieldDecl *FD, QualType FieldTy) final {
     Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) << FieldTy;
     IsInvalid = true;

clang/test/SemaSYCL/pointer-to-vla.cpp

@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -verify %s
+//
+// This test checks if compiler reports compilation error on an attempt to pass
+// a pointer to VLA as kernel argument
+
+#include "Inputs/sycl.hpp"
+
+void foo(unsigned X) {
+  using VLATy = float(*)[X];
+  VLATy PtrToVLA;
+  cl::sycl::kernel_single_task<class Kernel>([=]() {
+    // expected-error@+1 {{variable length arrays are not supported for the current target}}
+    (void)PtrToVLA;
+  });
+}