[SYCLomatic] Do not migrate __half_raw if it is a user-defined type (#1922)

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

Author: zhiweij1

clang/lib/DPCT/ASTTraversal.cpp

@@ -2684,6 +2684,11 @@ void VectorTypeNamespaceRule::runRule(const MatchFinder::MatchResult &Result) {
   }
   // Runrule for __half_raw implicitly convert to half.
   if (auto DRE = getNodeAsType<DeclRefExpr>(Result, "halfRawExpr")) {
+    if (const auto *RT =
+            DRE->getType().getCanonicalType()->getAs<RecordType>()) {
+      if (isUserDefinedDecl(RT->getDecl()))
+        return;
+    }
     ExprAnalysis EA;
     std::string Replacement;
     llvm::raw_string_ostream OS(Replacement);

clang/test/dpct/user_defined_half_raw.cu

@@ -0,0 +1,20 @@
+// RUN: dpct --format-range=none --out-root %T/user_defined_half_raw %s --cuda-include-path="%cuda-path/include"
+// RUN: FileCheck %s --match-full-lines --input-file %T/user_defined_half_raw/user_defined_half_raw.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/user_defined_half_raw/user_defined_half_raw.dp.cpp -o %T/user_defined_half_raw/user_defined_half_raw.dp.o %}
+
+struct __half_raw {
+  __device__ __half_raw() : _raw(0) {}
+  explicit __device__ __half_raw(unsigned short raw) : _raw(raw) {}
+  unsigned short _raw;
+};
+
+//      CHECK: __half_raw foo(unsigned short x) {
+// CHECK-NEXT:   __half_raw h;
+// CHECK-NEXT:   h._raw = x;
+// CHECK-NEXT:   return h;
+// CHECK-NEXT: }
+__device__ __half_raw foo(unsigned short x) {
+  __half_raw h;
+  h._raw = x;
+  return h;
+}