[SYCLomatic] Add query api mapping for 2 cub warpreduce API (#2899)

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

Author: intwanghao

clang/examples/DPCT/CUB/cub$$WarpReduce$$Reduce.cu

@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int thread_data, int valid_items) {
+  // Start
+  __shared__ typename cub::WarpReduce<int>::TempStorage temp_storage;
+  int result1 = cub::WarpReduce<int>(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/);
+  int result2 = cub::WarpReduce<int>(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/, valid_items/*int*/);
+  // End
+}

clang/examples/DPCT/CUB/cub$$WarpReduce$$Sum.cu

@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int thread_data, int valid_items) {
+  // Start
+  __shared__ typename cub::WarpReduce<int>::TempStorage temp_storage;
+  int result1 = cub::WarpReduce<int>(temp_storage).Sum(thread_data/*int*/);
+  int result2 = cub::WarpReduce<int>(temp_storage).Sum(thread_data/*int*/, valid_items/*int*/);
+  // End
+}

clang/test/dpct/query_api_mapping/CUB/cub_warp.cu

@@ -37,3 +37,20 @@
 // CHECK_WARPSCAN_BROADCAST:  Is migrated to:
 // CHECK_WARPSCAN_BROADCAST:    sycl::group_broadcast(sycl::ext::oneapi::this_work_item::get_sub_group(), thread_data, 0);
 
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::WarpReduce::Sum | FileCheck %s -check-prefix=CHECK_WARPREDUCE_SUM
+// CHECK_WARPREDUCE_SUM:  CUDA API:
+// CHECK_WARPREDUCE_SUM:    __shared__ typename cub::WarpReduce<int>::TempStorage temp_storage;
+// CHECK_WARPREDUCE_SUM:    int result1 = cub::WarpReduce<int>(temp_storage).Sum(thread_data/*int*/);
+// CHECK_WARPREDUCE_SUM:    int result2 = cub::WarpReduce<int>(temp_storage).Sum(thread_data/*int*/, valid_items/*int*/);
+// CHECK_WARPREDUCE_SUM:  Is migrated to:
+// CHECK_WARPREDUCE_SUM:    int result1 = sycl::reduce_over_group(sycl::ext::oneapi::this_work_item::get_sub_group(), thread_data, sycl::plus<>());
+// CHECK_WARPREDUCE_SUM:    int result2 = dpct::group::reduce_over_partial_group(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, valid_items, sycl::plus<>());
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::WarpReduce::Reduce | FileCheck %s -check-prefix=CHECK_WARPREDUCE_REDUCE
+// CHECK_WARPREDUCE_REDUCE:  CUDA API:
+// CHECK_WARPREDUCE_REDUCE:    __shared__ typename cub::WarpReduce<int>::TempStorage temp_storage;
+// CHECK_WARPREDUCE_REDUCE:    int result1 = cub::WarpReduce<int>(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/);
+// CHECK_WARPREDUCE_REDUCE:    int result2 = cub::WarpReduce<int>(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/, valid_items/*int*/);
+// CHECK_WARPREDUCE_REDUCE:  Is migrated to:
+// CHECK_WARPREDUCE_REDUCE:    int result1 = sycl::reduce_over_group(sycl::ext::oneapi::this_work_item::get_sub_group(), thread_data, sycl::minimum<>());
+// CHECK_WARPREDUCE_REDUCE:    int result2 = dpct::group::reduce_over_partial_group(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, valid_items, sycl::minimum<>());

clang/test/dpct/query_api_mapping/test_all.cu

@@ -833,6 +833,8 @@
 // CHECK-NEXT: cub::DeviceSelect::Flagged
 // CHECK-NEXT: cub::DeviceSelect::If
 // CHECK-NEXT: cub::DeviceSelect::Unique
+// CHECK-NEXT: cub::WarpReduce::Reduce
+// CHECK-NEXT: cub::WarpReduce::Sum
 // CHECK-NEXT: cub::WarpScan::Broadcast
 // CHECK-NEXT: cub::WarpScan::ExclusiveScan
 // CHECK-NEXT: cub::WarpScan::ExclusiveSum