diff --git a/clang/examples/DPCT/CUB/cub$$WarpReduce$$Reduce.cu b/clang/examples/DPCT/CUB/cub$$WarpReduce$$Reduce.cu new file mode 100644 index 000000000000..ba9ea0e1183a --- /dev/null +++ b/clang/examples/DPCT/CUB/cub$$WarpReduce$$Reduce.cu @@ -0,0 +1,11 @@ +// clang-format off +#include +#include + +__device__ void test(int thread_data, int valid_items) { + // Start + __shared__ typename cub::WarpReduce::TempStorage temp_storage; + int result1 = cub::WarpReduce(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/); + int result2 = cub::WarpReduce(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/, valid_items/*int*/); + // End +} diff --git a/clang/examples/DPCT/CUB/cub$$WarpReduce$$Sum.cu b/clang/examples/DPCT/CUB/cub$$WarpReduce$$Sum.cu new file mode 100644 index 000000000000..686fdc88a700 --- /dev/null +++ b/clang/examples/DPCT/CUB/cub$$WarpReduce$$Sum.cu @@ -0,0 +1,11 @@ +// clang-format off +#include +#include + +__device__ void test(int thread_data, int valid_items) { + // Start + __shared__ typename cub::WarpReduce::TempStorage temp_storage; + int result1 = cub::WarpReduce(temp_storage).Sum(thread_data/*int*/); + int result2 = cub::WarpReduce(temp_storage).Sum(thread_data/*int*/, valid_items/*int*/); + // End +} diff --git a/clang/test/dpct/query_api_mapping/CUB/cub_warp.cu b/clang/test/dpct/query_api_mapping/CUB/cub_warp.cu index b67684483279..e1668a568b53 100644 --- a/clang/test/dpct/query_api_mapping/CUB/cub_warp.cu +++ b/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::TempStorage temp_storage; +// CHECK_WARPREDUCE_SUM: int result1 = cub::WarpReduce(temp_storage).Sum(thread_data/*int*/); +// CHECK_WARPREDUCE_SUM: int result2 = cub::WarpReduce(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::TempStorage temp_storage; +// CHECK_WARPREDUCE_REDUCE: int result1 = cub::WarpReduce(temp_storage).Reduce(thread_data/*int*/, cub::Min()/*ReductionOp*/); +// CHECK_WARPREDUCE_REDUCE: int result2 = cub::WarpReduce(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<>()); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index a53aec96afee..2c2659e74af1 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -759,6 +759,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