[SYCLomatic] Add query api mapping for 22 cub blocklevel API and 4 cub devicelevel API (#2900)

intwanghao · web-flow · commit 5d261f049b29 · 2025-06-20T09:20:46.000+08:00
Signed-off-by: intwanghao &lt;hao3.wang@intel.com&gt;
diff --git a/clang/examples/DPCT/CUB/cub$$BlockExchange$$BlockedToStriped.cu b/clang/examples/DPCT/CUB/cub$$BlockExchange$$BlockedToStriped.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockExchange<int, 128, 4>(temp_storage).BlockedToStriped(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockExchange$$BlockedToWarpStriped.cu b/clang/examples/DPCT/CUB/cub$$BlockExchange$$BlockedToWarpStriped.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockExchange<int, 128, 4>(temp_storage).BlockedToWarpStriped(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockExchange$$ScatterToBlocked.cu b/clang/examples/DPCT/CUB/cub$$BlockExchange$$ScatterToBlocked.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4], int (&thread_rank)[4]) {
+  // Start
+  __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockExchange<int, 128, 4>(temp_storage).ScatterToBlocked(thread_data/*int(&)[4]*/, thread_rank/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockExchange$$ScatterToStriped.cu b/clang/examples/DPCT/CUB/cub$$BlockExchange$$ScatterToStriped.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4], int (&thread_rank)[4]) {
+  // Start
+  __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockExchange<int, 128, 4>(temp_storage).ScatterToStriped(thread_data/*int(&)[4]*/, thread_rank/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockExchange$$StripedToBlocked.cu b/clang/examples/DPCT/CUB/cub$$BlockExchange$$StripedToBlocked.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockExchange<int, 128, 4>(temp_storage).StripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockExchange$$WarpStripedToBlocked.cu b/clang/examples/DPCT/CUB/cub$$BlockExchange$$WarpStripedToBlocked.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockExchange<int, 128, 4>(temp_storage).WarpStripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockLoad$$Load.cu b/clang/examples/DPCT/CUB/cub$$BlockLoad$$Load.cu
@@ -0,0 +1,12 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int *src, int (&thread_data)[4], int end, int default_value) {
+  // Start
+  __shared__ typename cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>::TempStorage temp_storage;
+  cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>(temp_storage).Load(src/*int **/, thread_data/*int(&)[4]*/);
+  cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>(temp_storage).Load(src/*int **/, thread_data/*int(&)[4]*/, end/*int*/);
+  cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>(temp_storage).Load(src/*int **/, thread_data/*int(&)[4]*/, end/*int*/, default_value/*int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$Sort.cu b/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$Sort.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockRadixSort<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockRadixSort<int, 128, 4>(temp_storage).Sort(thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$SortBlockedToStriped.cu b/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$SortBlockedToStriped.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockRadixSort<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockRadixSort<int, 128, 4>(temp_storage).SortBlockedToStriped(thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$SortDescending.cu b/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$SortDescending.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockRadixSort<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockRadixSort<int, 128, 4>(temp_storage).SortDescending(thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$SortDescendingBlockedToStriped.cu b/clang/examples/DPCT/CUB/cub$$BlockRadixSort$$SortDescendingBlockedToStriped.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int (&thread_data)[4]) {
+  // Start
+  __shared__ typename cub::BlockRadixSort<int, 128, 4>::TempStorage temp_storage;
+  cub::BlockRadixSort<int, 128, 4>(temp_storage).SortDescendingBlockedToStriped(thread_data/*int(&)[4]*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockReduce$$Reduce.cu b/clang/examples/DPCT/CUB/cub$$BlockReduce$$Reduce.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int data) {
+  // Start
+  __shared__ typename cub::BlockReduce<int, 4>::TempStorage temp_storage;
+  cub::BlockReduce<int, 4>(temp_storage).Reduce(data/*int*/, cub::Sum()/*ReduceOp*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockReduce$$Sum.cu b/clang/examples/DPCT/CUB/cub$$BlockReduce$$Sum.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int data) {
+  // Start
+  __shared__ typename cub::BlockReduce<int, 4>::TempStorage temp_storage;
+  cub::BlockReduce<int, 4>(temp_storage).Sum(data/*int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockScan$$ExclusiveScan.cu b/clang/examples/DPCT/CUB/cub$$BlockScan$$ExclusiveScan.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int input, int output, int init) {
+  // Start
+  __shared__ typename cub::BlockScan<int, 4>::TempStorage temp_storage;
+  cub::BlockScan<int, 4>(temp_storage).ExclusiveScan(input/*int*/, output/*int &*/, init/*int*/, cub::Sum()/*ScanOp*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockScan$$ExclusiveSum.cu b/clang/examples/DPCT/CUB/cub$$BlockScan$$ExclusiveSum.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int input, int output) {
+  // Start
+  __shared__ typename cub::BlockScan<int, 4>::TempStorage temp_storage;
+  cub::BlockScan<int, 4>(temp_storage).ExclusiveSum(input/*int*/, output/*int &*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockScan$$InclusiveScan.cu b/clang/examples/DPCT/CUB/cub$$BlockScan$$InclusiveScan.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int input, int output) {
+  // Start
+  __shared__ typename cub::BlockScan<int, 4>::TempStorage temp_storage;
+  cub::BlockScan<int, 4>(temp_storage).InclusiveScan(input/*int*/, output/*int &*/, cub::Sum()/*ScanOp*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockScan$$InclusiveSum.cu b/clang/examples/DPCT/CUB/cub$$BlockScan$$InclusiveSum.cu
@@ -0,0 +1,10 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int input, int output) {
+  // Start
+  __shared__ typename cub::BlockScan<int, 4>::TempStorage temp_storage;
+  cub::BlockScan<int, 4>(temp_storage).InclusiveSum(input/*int*/, output/*int &*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Down.cu b/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Down.cu
@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+#include <cub/block/block_shuffle.cuh>
+
+__device__ void test(int (&input)[4], int (&output)[4]) {
+  // Start
+  __shared__ typename cub::BlockShuffle<int, 128>::TempStorage temp_storage;
+  cub::BlockShuffle<int, 128>(temp_storage).Down(input/*int*/, output/*int &*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Offset.cu b/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Offset.cu
@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+#include <cub/block/block_shuffle.cuh>
+
+__device__ void test(int input, int output, int distance) {
+  // Start
+  __shared__ typename cub::BlockShuffle<int, 128>::TempStorage temp_storage;
+  cub::BlockShuffle<int, 128>(temp_storage).Offset(input/*int*/, output/*int &*/, distance/*int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Rotate.cu b/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Rotate.cu
@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+#include <cub/block/block_shuffle.cuh>
+
+__device__ void test(int input, int output, unsigned int distance) {
+  // Start
+  __shared__ typename cub::BlockShuffle<int, 128>::TempStorage temp_storage;
+  cub::BlockShuffle<int, 128>(temp_storage).Rotate(input/*int*/, output/*int &*/, distance/*unsigned int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Up.cu b/clang/examples/DPCT/CUB/cub$$BlockShuffle$$Up.cu
@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+#include <cub/block/block_shuffle.cuh>
+
+__device__ void test(int (&input)[4], int (&output)[4]) {
+  // Start
+  __shared__ typename cub::BlockShuffle<int, 128>::TempStorage temp_storage;
+  cub::BlockShuffle<int, 128>(temp_storage).Up(input/*int*/, output/*int &*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$BlockStore$$Store.cu b/clang/examples/DPCT/CUB/cub$$BlockStore$$Store.cu
@@ -0,0 +1,11 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+__device__ void test(int *dst, int (&thread_data)[4], int end) {
+  // Start
+  __shared__ typename cub::BlockStore<int, 128, 4, cub::BLOCK_STORE_DIRECT>::TempStorage temp_storage;
+  cub::BlockStore<int, 128, 4, cub::BLOCK_STORE_DIRECT>(temp_storage).Store(dst/*int **/, thread_data/*int(&)[4]*/);
+  cub::BlockStore<int, 128, 4, cub::BLOCK_STORE_DIRECT>(temp_storage).Store(dst/*int **/, thread_data/*int(&)[4]*/, end/*int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$HistogramEven.cu b/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$HistogramEven.cu
@@ -0,0 +1,13 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+void test(int num_samples, float *d_samples, int *d_histogram, int num_levels, float lower_level, float upper_level) {
+  // Start
+  void *d_temp_storage = nullptr;
+  size_t temp_storage_bytes = 0;
+  cub::DeviceHistogram::HistogramEven(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*float **/, d_histogram/*int **/, num_levels/*int*/, lower_level/*float*/, upper_level/*float*/, num_samples/*int*/);
+  cudaMalloc(&d_temp_storage, temp_storage_bytes);
+  cub::DeviceHistogram::HistogramEven(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*float **/, d_histogram/*int **/, num_levels/*int*/, lower_level/*float*/, upper_level/*float*/, num_samples/*int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$HistogramRange.cu b/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$HistogramRange.cu
@@ -0,0 +1,13 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+void test(int num_samples, float *d_samples, int *d_histogram, int num_levels, float *d_levels) {
+  // Start
+  void *d_temp_storage = nullptr;
+  size_t temp_storage_bytes = 0;
+  cub::DeviceHistogram::HistogramRange(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*float **/, d_histogram/*int **/, num_levels/*int*/, d_levels/*float **/, num_samples/*int*/);
+  cudaMalloc(&d_temp_storage, temp_storage_bytes);
+  cub::DeviceHistogram::HistogramRange(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*float **/, d_histogram/*int **/, num_levels/*int*/, d_levels/*float **/, num_samples/*int*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$MultiHistogramEven.cu b/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$MultiHistogramEven.cu
@@ -0,0 +1,13 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+void test(int num_pixels, unsigned char *d_samples, int *(&d_histogram)[3], int (&num_levels)[3], unsigned int (&lower_level)[3], unsigned int (&upper_level)[3], cudaStream_t S) {
+  // Start
+  void *d_temp_storage = nullptr;
+  size_t temp_storage_bytes = 0;
+  cub::DeviceHistogram::MultiHistogramEven<4, 3>(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*unsigned char **/, d_histogram/*int *(&)[3]*/, num_levels/*int(&)[3]*/, lower_level/*unsigned int(&)[3]*/, upper_level/*unsigned int(&)[3]*/, num_pixels/*int*/, S/*cudaStream_t*/);
+  cudaMalloc(&d_temp_storage, temp_storage_bytes);
+  cub::DeviceHistogram::MultiHistogramEven<4, 3>(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*unsigned char **/, d_histogram/*int *(&)[3]*/, num_levels/*int(&)[3]*/, lower_level/*unsigned int(&)[3]*/, upper_level/*unsigned int(&)[3]*/, num_pixels/*int*/, S/*cudaStream_t*/);
+  // End
+}
diff --git a/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$MultiHistogramRange.cu b/clang/examples/DPCT/CUB/cub$$DeviceHistogram$$MultiHistogramRange.cu
@@ -0,0 +1,13 @@
+// clang-format off
+#include <cstddef>
+#include <cub/cub.cuh>
+
+void test(int num_pixels, unsigned char *d_samples, int *(&d_histogram)[3], int (&num_levels)[3], unsigned int *(&d_levels)[3]) {
+  // Start
+  void *d_temp_storage = nullptr;
+  size_t temp_storage_bytes = 0;
+  cub::DeviceHistogram::MultiHistogramRange<4, 3>(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*unsigned char **/, d_histogram/*int *(&)[3]*/, num_levels/*int(&)[3]*/, d_levels/*unsigned int *(&)[3]*/, num_pixels/*int*/);
+  cudaMalloc(&d_temp_storage, temp_storage_bytes);
+  cub::DeviceHistogram::MultiHistogramRange<4, 3>(d_temp_storage/*void **/, temp_storage_bytes/*size_t*/, d_samples/*unsigned char **/, d_histogram/*int *(&)[3]*/, num_levels/*int(&)[3]*/, d_levels/*unsigned int *(&)[3]*/, num_pixels/*int*/);
+  // End
+}
diff --git a/clang/test/dpct/query_api_mapping/CUB/cub_block.cu b/clang/test/dpct/query_api_mapping/CUB/cub_block.cu
diff --git a/clang/test/dpct/query_api_mapping/CUB/cub_device.cu b/clang/test/dpct/query_api_mapping/CUB/cub_device.cu
diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu
