[SYCLomatic] In the migration of BlockLoad/Store, fix the issue that local memory size is not correct calculated without load/store algorithm template type (#2773)

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

Author: intwanghao

clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp

@@ -258,8 +258,9 @@ void CubMemberCallRule::runRule(
         Name == "BlockedToWarpStriped";
     bool isBlockShuffle =
         Name == "Offset" || Name == "Rotate" || Name == "Up" || Name == "Down";
+    bool isBlockLoadStore = Name == "Load" || Name == "Store";
     if (isBlockRadixSort || isBlockExchange || isBlockShuffle ||
-        Name == "Load" || Name == "Store") {
+        isBlockLoadStore) {
       std::string HelpFuncName;
       if (isBlockRadixSort)
         HelpFuncName = "group_radix_sort";
@@ -300,6 +301,32 @@ void CubMemberCallRule::runRule(
         const auto &ItemsPreThreadArg = ClassSpecDecl->getTemplateArgs()[2];
         OS << ", " << ItemsPreThreadArg.getAsIntegral();
       }
+      if (isBlockLoadStore &&
+          !ClassSpecDecl->getTemplateArgs()[3].getIsDefaulted()) {
+        int AlgoType =
+            ClassSpecDecl->getTemplateArgs()[3].getAsIntegral().getExtValue();
+        if (Name == "Load") {
+          if (AlgoType == 3) {
+            OS << ", "
+               << MapNames::getDpctNamespace() +
+                      "group::group_load_algorithm::transpose";
+          } else if (AlgoType == 4) {
+            OS << ", "
+               << MapNames::getDpctNamespace() +
+                      "group::group_load_algorithm::sub_group_transpose";
+          }
+        } else {
+          if (AlgoType == 3) {
+            OS << ", "
+               << MapNames::getDpctNamespace() +
+                      "group::group_store_algorithm::transpose";
+          } else if (AlgoType == 4) {
+            OS << ", "
+               << MapNames::getDpctNamespace() +
+                      "group::group_store_algorithm::sub_group_transpose";
+          }
+        }
+      }
       OS << ">::get_local_memory_size";
       if (auto FuncInfo = DeviceFunctionDecl::LinkRedecls(FD)) {
         auto LocInfo = DpctGlobalInfo::getLocInfo(TempStorage);

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

@@ -125,5 +125,32 @@ int main() {
   // CHECK-NEXT:   });
   StripedKernel<<<1, 128>>>(d_data, 128);
   cudaStreamSynchronize(0);
+
+// CHECK:  q_ct1.submit(
+// CHECK:    [&](sycl::handler &cgh) {
+// CHECK:      sycl::stream stream_ct1(64 * 1024, 80, cgh);
+// CHECK:      sycl::local_accessor<uint8_t, 1> temp_storage_acc(dpct::group::group_load<int, 4, dpct::group::group_load_algorithm::transpose>::get_local_memory_size(sycl::range<3>(1, 1, 128).size()), cgh);
+// CHECK:      cgh.parallel_for(
+// CHECK:        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+// CHECK:        [=](sycl::nd_item<3> item_ct1) {
+// CHECK:          TransposeKernel(d_data, 128, item_ct1, stream_ct1, &temp_storage_acc[0]);
+// CHECK:        });
+// CHECK:    });
+  TransposeKernel<<<1, 128>>>(d_data, 128);
+  cudaStreamSynchronize(0);
+
+// CHECK:  q_ct1.submit(
+// CHECK:    [&](sycl::handler &cgh) {
+// CHECK:      sycl::stream stream_ct1(64 * 1024, 80, cgh);
+// CHECK:      sycl::local_accessor<uint8_t, 1> temp_storage_acc(dpct::group::group_load<int, 4, dpct::group::group_load_algorithm::sub_group_transpose>::get_local_memory_size(sycl::range<3>(1, 1, 128).size()), cgh);
+// CHECK:      cgh.parallel_for(
+// CHECK:        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+// CHECK:        [=](sycl::nd_item<3> item_ct1) {
+// CHECK:          SubGroupTransposeKernel(d_data, 128, item_ct1, stream_ct1, &temp_storage_acc[0]);
+// CHECK:        });
+// CHECK:    });
+  SubGroupTransposeKernel<<<1, 128>>>(d_data, 128);
+  cudaStreamSynchronize(0);
+
   return 0;
 }

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

@@ -135,6 +135,31 @@ int main() {
   // CHECK-NEXT:   });
   StripedKernel<<<1, 128>>>(d_data, 5);
   cudaStreamSynchronize(0);
+
+// CHECK:   q_ct1.submit(
+// CHECK:     [&](sycl::handler &cgh) {
+// CHECK:       sycl::local_accessor<uint8_t, 1> temp_storage_acc(dpct::group::group_store<int, 4, dpct::group::group_store_algorithm::transpose>::get_local_memory_size(sycl::range<3>(1, 1, 128).size()), cgh);
+// CHECK:       cgh.parallel_for(
+// CHECK:         sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+// CHECK:         [=](sycl::nd_item<3> item_ct1) {
+// CHECK:           TransposeKernel(d_data, 5, item_ct1, &temp_storage_acc[0]);
+// CHECK:         });
+// CHECK:     });
+  TransposeKernel<<<1, 128>>>(d_data, 5);
+  cudaStreamSynchronize(0);
+
+// CHECK:  q_ct1.submit(
+// CHECK:    [&](sycl::handler &cgh) {
+// CHECK:      sycl::local_accessor<uint8_t, 1> temp_storage_acc(dpct::group::group_store<int, 4, dpct::group::group_store_algorithm::sub_group_transpose>::get_local_memory_size(sycl::range<3>(1, 1, 128).size()), cgh);
+// CHECK:      cgh.parallel_for(
+// CHECK:        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
+// CHECK:        [=](sycl::nd_item<3> item_ct1) {
+// CHECK:          SubGroupTransposeKernel(d_data, 5, item_ct1, &temp_storage_acc[0]);
+// CHECK:        });
+// CHECK:    });
+  SubGroupTransposeKernel<<<1, 128>>>(d_data, 5);
+  cudaStreamSynchronize(0);
+
   for (int i = 0; i < 512; ++i)
     printf("%d%c", d_data[i], (i == 511 ? '\n' : ' '));
   return 0;