[SYCLomatic] Enable cooperativate_groups::reduce function (#863)

Signed-off-by: Chen, Sheng S <[email protected]>

Author: ShengchenJ

clang/lib/DPCT/APINamesCooperativeGroups.inc

@@ -69,6 +69,21 @@ CONDITIONAL_FACTORY_ENTRY(
                 UNSUPPORT_FACTORY_ENTRY("sync",
                                         Diagnostics::API_NOT_MIGRATED,
                                         ARG("sync"))))))
+
+/*
+   --- cg::reduce(group, val, op) ---
+
+   cg::reduce(cg::thread_block_tile<32>, val, op)
+   => sycl::reduce_over_group(item.get_sub_group(), val, op)
+*/
+CONDITIONAL_FACTORY_ENTRY(
+    argHasThreadBlockTileType(0, 32),
+        CALL_FACTORY_ENTRY("reduce", CALL(MapNames::getClNamespace() + "reduce_over_group",
+                                    SUBGROUP, ARG(1), ARG(2))),
+        UNSUPPORT_FACTORY_ENTRY("reduce",
+                                    Diagnostics::API_NOT_MIGRATED,
+                                    ARG("reduce")))
+
 /*
   --- cg::thread_rank(X) ---
   cg::thread_rank(cg::thread_block)

clang/lib/DPCT/APINamesTemplateType.inc

@@ -173,3 +173,27 @@ TYPE_REWRITE_ENTRY("thrust::identity",
                    TYPE_FACTORY(STR("oneapi::dpl::identity")))
 
 TYPE_REWRITE_ENTRY("thrust::null_type", TYPE_FACTORY(STR("dpct::null_type")))
+
+TYPE_REWRITE_ENTRY("cooperative_groups::__v1::plus",
+                   TYPE_FACTORY(STR(MapNames::getClNamespace() + "plus"),
+                                TEMPLATE_ARG(0)))
+
+TYPE_REWRITE_ENTRY("cooperative_groups::__v1::less",
+                   TYPE_FACTORY(STR(MapNames::getClNamespace() + "minimum"),
+                                TEMPLATE_ARG(0)))
+
+TYPE_REWRITE_ENTRY("cooperative_groups::__v1::greater",
+                   TYPE_FACTORY(STR(MapNames::getClNamespace() + "maximum"),
+                                TEMPLATE_ARG(0)))
+
+TYPE_REWRITE_ENTRY("cooperative_groups::__v1::bit_and",
+                   TYPE_FACTORY(STR(MapNames::getClNamespace() + "bit_and"),
+                                TEMPLATE_ARG(0)))
+
+TYPE_REWRITE_ENTRY("cooperative_groups::__v1::bit_xor",
+                   TYPE_FACTORY(STR(MapNames::getClNamespace() + "bit_xor"),
+                                TEMPLATE_ARG(0)))
+
+TYPE_REWRITE_ENTRY("cooperative_groups::__v1::bit_or",
+                   TYPE_FACTORY(STR(MapNames::getClNamespace() + "bit_or"),
+                                TEMPLATE_ARG(0)))

clang/lib/DPCT/ASTTraversal.cpp

@@ -11933,6 +11933,11 @@ void CooperativeGroupsFunctionRule::runRule(
         }
       }
     }
+  } else if (FuncName == "reduce") {
+    RUW.NeedReport = false;
+    ExprAnalysis EA(CE);
+    emplaceTransformation(EA.getReplacement());
+    EA.applyAllSubExprRepl();
   }
 }
 

clang/lib/DPCT/ExprAnalysis.cpp

@@ -573,7 +573,8 @@ void ExprAnalysis::analyzeExpr(const CXXTemporaryObjectExpr *Temp) {
       Temp->getType().getCanonicalType());
   if ((StringRef(TypeName).startswith("cub::") &&
        CubTypeRule::CanMappingToSyclType(TypeName)) ||
-      StringRef(TypeName).startswith("thrust::")) {
+       StringRef(TypeName).startswith("thrust::") ||
+       StringRef(TypeName).startswith("cooperative_groups::")) {
     analyzeType(Temp->getTypeSourceInfo()->getTypeLoc());
   }
   analyzeExpr(static_cast<const CXXConstructExpr *>(Temp));
@@ -1034,8 +1035,11 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE) {
     }
   case TypeLoc::TemplateSpecialization: {
     llvm::raw_string_ostream OS(TyName);
+    TyName.clear();
     auto &TSTL = TYPELOC_CAST(TemplateSpecializationTypeLoc);
-    TSTL.getTypePtr()->getTemplateName().print(OS, Context.getPrintingPolicy());
+    auto PP = Context.getPrintingPolicy();
+    PP.PrintCanonicalTypes = 1;
+    TSTL.getTypePtr()->getTemplateName().print(OS, PP, TemplateName::Qualified::Fully);
     if (!TypeLocRewriterFactoryBase::TypeLocRewriterMap)
       return;
     auto Itr = TypeLocRewriterFactoryBase::TypeLocRewriterMap->find(OS.str());

clang/lib/DPCT/MapNames.cpp

@@ -4428,4 +4428,4 @@ const std::vector<std::string> MemoryDataTypeRule::RemoveMember{
 
 const std::unordered_set<std::string> MapNames::CooperativeGroupsAPISet{
     "this_thread_block", "sync", "tiled_partition",
-    "thread_rank",       "size", "shfl_down"};
+    "thread_rank",       "size", "shfl_down", "reduce"};

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

@@ -559,4 +559,15 @@ extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
 #pragma pop_macro("__clang_major__")
 #endif // SYCLomatic_CUSTOMIZATION
 
+#if defined(SYCLomatic_CUSTOMIZATION)
+// Fixed the parsing error when source include <cooperative_groups/reduce.h>.
+#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
+# if CUDA_VERSION >= 11800
+#include <crt/sm_90_rt.h>
+#endif
+#if CUDA_VERSION >= 11000
+#include <crt/sm_80_rt.h>
+#endif
+#endif // SYCLomatic_CUSTOMIZATION
+
 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__

clang/test/dpct/cooperative_groups_reduce.cu

@@ -0,0 +1,30 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0 
+// RUN: dpct --format-range=none -out-root %T/cooperative_groups_reduce %s --cuda-include-path="%cuda-path/include" --extra-arg="-std=c++14"
+// RUN: FileCheck %s --match-full-lines --input-file %T/cooperative_groups_reduce/cooperative_groups_reduce.dp.cpp
+
+
+#include <cooperative_groups.h>
+#include <cooperative_groups/reduce.h>
+
+namespace cg = cooperative_groups;
+
+__device__ void testReduce(double *sdata, const cg::thread_block &cta) {
+  const unsigned int tid = cta.thread_rank();
+  cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);
+  int *idata;
+  // CHECK: sycl::reduce_over_group(item_ct1.get_sub_group(), sdata[tid], sycl::plus<double>());
+  cg::reduce(tile32, sdata[tid], cg::plus<double>());  
+  // CHECK: sycl::reduce_over_group(item_ct1.get_sub_group(), sdata[tid], sycl::minimum<double>());
+  cg::reduce(tile32, sdata[tid], cg::less<double>());  
+  // CHECK: sycl::reduce_over_group(item_ct1.get_sub_group(), sdata[tid], sycl::maximum<double>());
+  cg::reduce(tile32, sdata[tid], cg::greater<double>());  
+  // CHECK: sycl::reduce_over_group(item_ct1.get_sub_group(), idata[tid], sycl::bit_and<int>());
+  cg::reduce(tile32, idata[tid], cg::bit_and<int>());  
+  // CHECK: sycl::reduce_over_group(item_ct1.get_sub_group(), idata[tid], sycl::bit_xor<int>());
+  cg::reduce(tile32, idata[tid], cg::bit_xor<int>());
+  // CHECK: sycl::reduce_over_group(item_ct1.get_sub_group(), idata[tid], sycl::bit_or<int>());
+  cg::reduce(tile32, idata[tid], cg::bit_or<int>());
+  cg::sync(cta);
+
+}