[SYCLomatic] Match the typeAliasDecl for Cub Type, match template scan and reduce for block and warp. (#2887)

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: ShengchenJ

clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp

@@ -946,10 +946,14 @@ void CubRule::processCubDeclStmt(const DeclStmt *DS) {
       if (isTypeInAnalysisScope(ObjCanonicalType.getTypePtr())) {
         continue;
       } else if (ObjTypeStr.find("class cub::WarpScan") == 0 ||
-                 ObjTypeStr.find("class cub::WarpReduce") == 0) {
+                 ObjTypeStr.find("class cub::WarpReduce") == 0 ||
+                 ObjTypeStr.find("WarpScan<") == 0 ||
+                 ObjTypeStr.find("WarpReduce<") == 0) {
         Repl = DpctGlobalInfo::getSubGroup(DRE);
       } else if (ObjTypeStr.find("class cub::BlockScan") == 0 ||
-                 ObjTypeStr.find("class cub::BlockReduce") == 0) {
+                 ObjTypeStr.find("class cub::BlockReduce") == 0 ||
+                 ObjTypeStr.find("BlockScan<") == 0 ||
+                 ObjTypeStr.find("BlockReduce<") == 0) {
         Repl = DpctGlobalInfo::getGroup(DRE);
       } else {
         continue;
@@ -975,9 +979,12 @@ void CubRule::processCubTypeDefOrUsing(const TypedefNameDecl *TD) {
   std::string TypeName = TD->getNameAsString();
   auto &Context = dpct::DpctGlobalInfo::getContext();
   auto &SM = dpct::DpctGlobalInfo::getSourceManager();
-  auto MyMatcher = compoundStmt(forEachDescendant(
-      typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasName(TypeName))))))
-          .bind("typeLoc")));
+
+  auto MyMatcher = compoundStmt(
+      forEachDescendant(typeLoc(loc(qualType(hasDeclaration(
+                                    anyOf(typedefDecl(hasName(TypeName)),
+                                          typeAliasDecl(hasName(TypeName)))))))
+                            .bind("typeLoc")));
   auto MatcherScope = DpctGlobalInfo::findAncestor<CompoundStmt>(TD);
   if (!MatcherScope)
     return;
@@ -1033,6 +1040,7 @@ void CubRule::processCubTypeDefOrUsing(const TypedefNameDecl *TD) {
       }
     }
   }
+
   if (DeleteFlag) {
     emplaceTransformation(new ReplaceDecl(TD, ""));
   } else {

clang/test/dpct/cub/type/fix_alias_bug.cu

@@ -0,0 +1,40 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
+// RUN: dpct --format-range=none -in-root %S -out-root %T/type/fix_alias_bug %S/fix_alias_bug.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/type/fix_alias_bug/fix_alias_bug.dp.cpp %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/type/fix_alias_bug/fix_alias_bug.dp.cpp -o %T/type/fix_alias_bug/fix_alias_bug.dp.o %}
+
+#include <cub/cub.cuh>
+#include <cuda.h>
+template <bool> struct warp_reduce;
+
+template <> struct warp_reduce<true> {
+  template <typename T, typename reducer_t>
+  inline T operator()(const T &value_, bool all, const reducer_t &r) {
+    // CHECK: using warp_reduce_t = sycl::sub_group;
+
+    using warp_reduce_t = cub::WarpReduce<T, 10>;
+
+    typename warp_reduce_t::TempStorage dummy_storage;
+    // CHECK: warp_reduce_t warp_reduce(sycl::ext::oneapi::this_work_item::get_sub_group());
+
+    warp_reduce_t warp_reduce(dummy_storage);
+  }
+};
+
+template <> struct warp_reduce<false> {
+  template <typename T, typename reducer_t>
+  inline T operator()(const T &value_, bool all, const reducer_t &r) {
+    // CHECK: using warp_reduce_t = sycl::sub_group;
+
+    using warp_reduce_t = cub::WarpReduce<T, 10>;
+
+    typename warp_reduce_t::TempStorage dummy_storage;
+    // CHECK: warp_reduce_t warp_reduce(sycl::ext::oneapi::this_work_item::get_sub_group());
+    warp_reduce_t warp_reduce(dummy_storage);
+  }
+};
+int main() {
+  int a, b;
+  warp_reduce<true>()(a, true, b);
+}