[SYCLomatic] enhance cub scan API migration in case thrust::plus/maximum/minimum() used as BinaryOp functor. (#2852)

Signed-off-by: intwanghao <hao3.wang@intel.com>

Author: intwanghao

clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp

@@ -846,11 +846,17 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
   auto processOperatorExpr = [&](const Expr *Obj) {
     std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
         Obj->getType().getCanonicalType());
-    if (OpType == "cub::Sum" || OpType == "cuda::std::plus<void>") {
+    if (OpType.find("cub::Sum") != std::string::npos ||
+        OpType.find("cuda::std::plus") != std::string::npos ||
+        OpType.find("thrust::plus") != std::string::npos) {
       OpRepl = MapNames::getClNamespace() + "plus<>()";
-    } else if (OpType == "cub::Max" || OpType == "cuda::maximum<void>") {
+    } else if (OpType.find("cub::Max") != std::string::npos ||
+               OpType.find("cuda::maximum") != std::string::npos ||
+               OpType.find("thrust::maximum") != std::string::npos) {
       OpRepl = MapNames::getClNamespace() + "maximum<>()";
-    } else if (OpType == "cub::Min" || OpType == "cuda::minimum<void>") {
+    } else if (OpType.find("cub::Min") != std::string::npos ||
+               OpType.find("cuda::minimum") != std::string::npos ||
+               OpType.find("thrust::minimum") != std::string::npos) {
       OpRepl = MapNames::getClNamespace() + "minimum<>()";
     }
   };
@@ -861,17 +867,21 @@ std::string CubRule::getOpRepl(const Expr *Operator) {
     } else {
       auto CtorArg = Op->getArg(0)->IgnoreImplicitAsWritten();
       if (auto DRE = dyn_cast<DeclRefExpr>(CtorArg)) {
-        auto D = DRE->getDecl();
-        if (!D)
-          return OpRepl;
-        std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
-            D->getType().getCanonicalType());
-        if (OpType == "cub::Sum" || OpType == "cub::Max" ||
-            OpType == "cub::Min" || OpType == "cuda::std::plus<void>" ||
-            OpType == "cuda::maximum<void>" ||
-            OpType == "cuda::minimum<void>") {
-          ExprAnalysis EA(Operator);
-          OpRepl = EA.getReplacedString();
+        if (auto D = DRE->getDecl()) {
+          std::string OpType = DpctGlobalInfo::getUnqualifiedTypeName(
+              D->getType().getCanonicalType());
+          if (OpType.find("cub::Sum") != std::string::npos ||
+              OpType.find("cub::Max") != std::string::npos ||
+              OpType.find("cub::Min") != std::string::npos ||
+              OpType.find("cuda::std::plus") != std::string::npos ||
+              OpType.find("cuda::maximum") != std::string::npos ||
+              OpType.find("cuda::minimum") != std::string::npos ||
+              OpType.find("thrust::plus") != std::string::npos ||
+              OpType.find("thrust::maximum") != std::string::npos ||
+              OpType.find("thrust::minimum") != std::string::npos) {
+            ExprAnalysis EA(Operator);
+            OpRepl = EA.getReplacedString();
+          }
         }
       } else if (auto CXXTempObj = dyn_cast<CXXTemporaryObjectExpr>(CtorArg)) {
         processOperatorExpr(CXXTempObj);

clang/test/dpct/cub/blocklevel/blockscan.cu

@@ -135,6 +135,26 @@ __global__ void InclusiveScanKernel(int* data) {
   data[threadid] = output;
 }
 
+// CHECK: void InclusiveScanKernelThrustFunctor(int* data) {
+// CHECK:   int threadid = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_local_id(2);
+// CHECK:   int input = data[threadid];
+// CHECK:   int output = 0;
+// CHECK:   output = sycl::inclusive_scan_over_group(sycl::ext::oneapi::this_work_item::get_work_group<3>(), input, sycl::plus<>());
+// CHECK:   data[threadid] = output;
+// CHECK: }
+__global__ void InclusiveScanKernelThrustFunctor(int* data) {
+  typedef cub::BlockScan<int, 4> BlockScan;
+
+  __shared__ typename BlockScan::TempStorage temp1;
+
+  int threadid = threadIdx.x;
+
+  int input = data[threadid];
+  int output = 0;
+  BlockScan(temp1).InclusiveScan(input, output, thrust::plus<>());
+  data[threadid] = output;
+}
+
 //CHECK: void InclusiveScanKernel_Max(int* data) {
 //CHECK-EMPTY:
 //CHECK-NEXT:  int threadid = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_local_id(2);