[SYCLomatic] Support migration for cub::{StoreDirectBlocked, StoreDirectStriped} API

clang/lib/DPCT/APINames_CUB.inc

@@ -107,7 +107,6 @@ ENTRY_MEMBER_FUNCTION(cub::BlockExchange, cub::BlockExchange, ScatterToStripedFl
 ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, InitHistogram, InitHistogram, false, NO_FLAG, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Histogram, Histogram, false, NO_FLAG, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Composite, Composite, false, NO_FLAG, P4, "Comment")
-ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, false, NO_FLAG, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::BlockStore, cub::BlockStore, Store, Store, false, NO_FLAG, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, Sort, Sort, false, NO_FLAG, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, SortDescending, SortDescending, false, NO_FLAG, P4, "Comment")
@@ -124,7 +123,6 @@ ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Rotate, Rotate, fals
 ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Up, Up, false, NO_FLAG, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Down, Down, false, NO_FLAG, P4, "Comment")
 
-
 // Device Level
 ENTRY_MEMBER_FUNCTION(cub::DeviceAdjacentDifference, cub::DeviceAdjacentDifference, SubtractLeftCopy, SubtractLeftCopy, false, NO_FLAGE, P4, "Comment")
 ENTRY_MEMBER_FUNCTION(cub::DeviceAdjacentDifference, cub::DeviceAdjacentDifference, SubtractLeft, SubtractLeft, false, NO_FLAGE, P4, "Comment")
@@ -192,10 +190,10 @@ ENTRY(cub::ThreadLoad, cub::ThreadLoad, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::ThreadStore, cub::ThreadStore, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::LoadDirectBlocked, cub::LoadDirectBlocked, false, NO_FLAG, P4, "Comment")
 ENTRY(cub::LoadDirectBlockedVectorized, cub::LoadDirectBlockedVectorized, false, NO_FLAG, P4, "Comment")
-ENTRY(cub::StoreDirectBlocked, cub::StoreDirectBlocked, false, NO_FLAG, P4, "Comment")
+ENTRY(cub::StoreDirectBlocked, cub::StoreDirectBlocked, true, NO_FLAG, P4, "Comment")
 ENTRY(cub::StoreDirectBlockedVectorized, cub::StoreDirectBlockedVectorized, false, NO_FLAG, P4, "Comment")
 ENTRY(cub::LoadDirectStriped, cub::LoadDirectStriped, false, NO_FLAG, P4, "Comment")
-ENTRY(cub::StoreDirectStriped, cub::StoreDirectStriped, false, NO_FLAG, P4, "Comment")
+ENTRY(cub::StoreDirectStriped, cub::StoreDirectStriped, true, NO_FLAG, P4, "Comment")
 ENTRY(cub::LoadDirectWarpStriped, cub::LoadDirectWarpStriped, false, NO_FLAG, P4, "Comment")
 ENTRY(cub::StoreDirectWarpStriped, cub::StoreDirectWarpStriped, false, NO_FLAG, P4, "Comment")
 

clang/lib/DPCT/CUBAPIMigration.cpp

@@ -192,7 +192,8 @@ void CubIntrinsicRule::registerMatcher(ast_matchers::MatchFinder &MF) {
                          "WarpId", "SyncStream", "CurrentDevice", "DeviceCount",
                          "DeviceCountUncached", "DeviceCountCachedValue",
                          "PtxVersion", "PtxVersionUncached", "SmVersion",
-                         "SmVersionUncached", "RowMajorTid"),
+                         "SmVersionUncached", "RowMajorTid", "StoreDirectBlocked",
+                         "StoreDirectStriped"),
               hasAncestor(namespaceDecl(hasName("cub")))))))
           .bind("IntrinsicCall"),
       this);
@@ -1508,7 +1509,7 @@ void CubRule::processTypeLoc(const TypeLoc *TL) {
                                       MapNames::getClNamespace() + "sub_group",
                                       SM));
   } else if (TypeName.find("class cub::BlockScan") == 0 ||
-             TypeName.find("class cub::BlockReduce") == 0) {
+             TypeName.find("class cub::BlockReduce") == 0 ) {
     auto DeviceFuncDecl = DpctGlobalInfo::findAncestor<FunctionDecl>(TL);
     if (DeviceFuncDecl && (DeviceFuncDecl->hasAttr<CUDADeviceAttr>() ||
                            DeviceFuncDecl->hasAttr<CUDAGlobalAttr>())) {

clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp

@@ -8,6 +8,7 @@
 
 #include "CallExprRewriterCUB.h"
 #include "CallExprRewriterCommon.h"
+#include "InclusionHeaders.h"
 
 using namespace clang::dpct;
 
@@ -112,6 +113,20 @@ RewriterMap dpct::createUtilityFunctionsRewriterMap() {
                 CALL(MapNames::getDpctNamespace() + "get_minor_version",
                      makeDeviceStr()),
                 LITERAL("10"))))
+      // cub::StoreDirectBlocked
+      HEADER_INSERT_FACTORY(
+          HeaderType::HT_DPCT_GROUP_Utils,
+          CALL_FACTORY_ENTRY(
+              "cub::StoreDirectBlocked",
+              CALL(MapNames::getDpctNamespace() + "group::store_blocked", NDITEM,
+                   ARG(1), ARG(2))))
+      // cub::StoreDirectStriped
+      HEADER_INSERT_FACTORY(
+          HeaderType::HT_DPCT_GROUP_Utils,
+          CALL_FACTORY_ENTRY(
+              "cub::StoreDirectStriped",
+              CALL(MapNames::getDpctNamespace() + "group::store_striped", NDITEM,
+                   ARG(1), ARG(2))))
       // cub::RowMajorTid
       MEMBER_CALL_FACTORY_ENTRY("cub::RowMajorTid", NDITEM, /*IsArrow=*/false,
                                 "get_local_linear_id")};

clang/test/dpct/cub/intrinsic/store.cu

@@ -0,0 +1,23 @@
+// 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 -in-root %S -out-root %T/intrinsic/store %S/store.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/intrinsic/store/store.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/intrinsic/store/store.dp.cpp -o %T/intrinsic/store/store.dp.o %}
+
+// CHECK:#include <sycl/sycl.hpp>
+// CHECK:#include <dpct/dpct.hpp>
+// CHECK:#include <dpct/group_utils.hpp>
+#include <cub/cub.cuh>
+
+__global__ void TestStoreStriped(int *d_data) {
+  int thread_data[4];
+  // CHECK: dpct::group::load_striped(item_ct1, d_data, thread_data);
+  cub::StoreDirectStriped<128>(threadIdx.x, d_data, thread_data);
+}
+
+
+__global__ void BlockedToStripedKernel(int *d_data) {
+  int thread_data[4];
+  // CHECK: dpct::group::load_blocked(item_ct1, d_data, thread_data);
+  cub::StoreDirectBlocked(threadIdx.x, d_data, thread_data);
+}