Skip to content

Commit 64c6a87

Browse files
intwanghaointwanghao
committed
fix
Signed-off-by: intwanghao <hao3.wang@intel.com>
1 parent 6c89f6c commit 64c6a87

8 files changed

Lines changed: 112 additions & 0 deletions

‎clang/examples/DPCT/CUB/cub$$BlockExchange$$BlockedToStriped.cu‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// clang-format off
2+
#include <cstddef>
3+
#include <cub/cub.cuh>
4+
5+
__device__ void test(int (&thread_data)[4]) {
6+
// Start
7+
__shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
8+
cub::BlockExchange<int, 128, 4>(temp_storage).BlockedToStriped(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
9+
// End
10+
}

‎clang/examples/DPCT/CUB/cub$$BlockExchange$$BlockedToWarpStriped.cu‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// clang-format off
2+
#include <cstddef>
3+
#include <cub/cub.cuh>
4+
5+
__device__ void test(int (&thread_data)[4]) {
6+
// Start
7+
__shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
8+
cub::BlockExchange<int, 128, 4>(temp_storage).BlockedToWarpStriped(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
9+
// End
10+
}

‎clang/examples/DPCT/CUB/cub$$BlockExchange$$ScatterToBlocked.cu‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// clang-format off
2+
#include <cstddef>
3+
#include <cub/cub.cuh>
4+
5+
__device__ void test(int (&thread_data)[4], int (&thread_rank)[4]) {
6+
// Start
7+
__shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
8+
cub::BlockExchange<int, 128, 4>(temp_storage).ScatterToBlocked(thread_data/*int(&)[4]*/, thread_rank/*int(&)[4]*/);
9+
// End
10+
}

‎clang/examples/DPCT/CUB/cub$$BlockExchange$$ScatterToStriped.cu‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// clang-format off
2+
#include <cstddef>
3+
#include <cub/cub.cuh>
4+
5+
__device__ void test(int (&thread_data)[4], int (&thread_rank)[4]) {
6+
// Start
7+
__shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
8+
cub::BlockExchange<int, 128, 4>(temp_storage).ScatterToStriped(thread_data/*int(&)[4]*/, thread_rank/*int(&)[4]*/);
9+
// End
10+
}

‎clang/examples/DPCT/CUB/cub$$BlockExchange$$StripedToBlocked.cu‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// clang-format off
2+
#include <cstddef>
3+
#include <cub/cub.cuh>
4+
5+
__device__ void test(int (&thread_data)[4]) {
6+
// Start
7+
__shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
8+
cub::BlockExchange<int, 128, 4>(temp_storage).StripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
9+
// End
10+
}

‎clang/examples/DPCT/CUB/cub$$BlockExchange$$WarpStripedToBlocked.cu‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// clang-format off
2+
#include <cstddef>
3+
#include <cub/cub.cuh>
4+
5+
__device__ void test(int (&thread_data)[4]) {
6+
// Start
7+
__shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
8+
cub::BlockExchange<int, 128, 4>(temp_storage).WarpStripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
9+
// End
10+
}

‎clang/test/dpct/query_api_mapping/CUB/cub_block.cu‎

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.2, cuda-11.4
2+
// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.2, v11.4
3+
// UNSUPPORTED: system-windows
4+
// clang-format off
5+
6+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockExchange::BlockedToStriped | FileCheck %s -check-prefix=CHECK_BLOCKEXCHANGE_BLOCKEDTOSTRIPED
7+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOSTRIPED: CUDA API:
8+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOSTRIPED: __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
9+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOSTRIPED: cub::BlockExchange<int, 128, 4>(temp_storage).BlockedToStriped(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
10+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOSTRIPED: Is migrated to:
11+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOSTRIPED: dpct::group::exchange<int, 4>(temp_storage).blocked_to_striped(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, thread_data);
12+
13+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockExchange::BlockedToWarpStriped | FileCheck %s -check-prefix=CHECK_BLOCKEXCHANGE_BLOCKEDTOWARPSTRIPED
14+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOWARPSTRIPED: CUDA API:
15+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOWARPSTRIPED: __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
16+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOWARPSTRIPED: cub::BlockExchange<int, 128, 4>(temp_storage).BlockedToWarpStriped(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
17+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOWARPSTRIPED: Is migrated to:
18+
// CHECK_BLOCKEXCHANGE_BLOCKEDTOWARPSTRIPED: dpct::group::exchange<int, 4>(temp_storage).blocked_to_sub_group_striped(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, thread_data);
19+
20+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockExchange::ScatterToBlocked | FileCheck %s -check-prefix=CHECK_BLOCKEXCHANGE_SCATTERTOBLOCKED
21+
// CHECK_BLOCKEXCHANGE_SCATTERTOBLOCKED: CUDA API:
22+
// CHECK_BLOCKEXCHANGE_SCATTERTOBLOCKED: __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
23+
// CHECK_BLOCKEXCHANGE_SCATTERTOBLOCKED: cub::BlockExchange<int, 128, 4>(temp_storage).ScatterToBlocked(thread_data/*int(&)[4]*/, thread_rank/*int(&)[4]*/);
24+
// CHECK_BLOCKEXCHANGE_SCATTERTOBLOCKED: Is migrated to:
25+
// CHECK_BLOCKEXCHANGE_SCATTERTOBLOCKED: dpct::group::exchange<int, 4>(temp_storage).scatter_to_blocked(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, thread_rank);
26+
27+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockExchange::ScatterToStriped | FileCheck %s -check-prefix=CHECK_BLOCKEXCHANGE_SCATTERTOSTRIPED
28+
// CHECK_BLOCKEXCHANGE_SCATTERTOSTRIPED: CUDA API:
29+
// CHECK_BLOCKEXCHANGE_SCATTERTOSTRIPED: __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
30+
// CHECK_BLOCKEXCHANGE_SCATTERTOSTRIPED: cub::BlockExchange<int, 128, 4>(temp_storage).ScatterToStriped(thread_data/*int(&)[4]*/, thread_rank/*int(&)[4]*/);
31+
// CHECK_BLOCKEXCHANGE_SCATTERTOSTRIPED: Is migrated to:
32+
// CHECK_BLOCKEXCHANGE_SCATTERTOSTRIPED: dpct::group::exchange<int, 4>(temp_storage).scatter_to_striped(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, thread_rank);
33+
34+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockExchange::StripedToBlocked | FileCheck %s -check-prefix=CHECK_BLOCKEXCHANGE_STRIPEDTOBLOCKED
35+
// CHECK_BLOCKEXCHANGE_STRIPEDTOBLOCKED: CUDA API:
36+
// CHECK_BLOCKEXCHANGE_STRIPEDTOBLOCKED: __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
37+
// CHECK_BLOCKEXCHANGE_STRIPEDTOBLOCKED: cub::BlockExchange<int, 128, 4>(temp_storage).StripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
38+
// CHECK_BLOCKEXCHANGE_STRIPEDTOBLOCKED: Is migrated to:
39+
// CHECK_BLOCKEXCHANGE_STRIPEDTOBLOCKED: dpct::group::exchange<int, 4>(temp_storage).striped_to_blocked(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, thread_data);
40+
41+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockExchange::WarpStripedToBlocked | FileCheck %s -check-prefix=CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED
42+
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: CUDA API:
43+
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: __shared__ typename cub::BlockExchange<int, 128, 4>::TempStorage temp_storage;
44+
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: cub::BlockExchange<int, 128, 4>(temp_storage).WarpStripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
45+
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: Is migrated to:
46+
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: dpct::group::exchange<int, 4>(temp_storage).sub_group_striped_to_blocked(sycl::ext::oneapi::this_work_item::get_nd_item<3>(), thread_data, thread_data);

‎clang/test/dpct/query_api_mapping/test_all.cu‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -739,6 +739,12 @@
739739
// CHECK-NEXT: cuTexRefSetFilterMode
740740
// CHECK-NEXT: cuTexRefSetFlags
741741
// CHECK-NEXT: cuTexRefSetFormat
742+
// CHECK-NEXT: cub::BlockExchange::BlockedToStriped
743+
// CHECK-NEXT: cub::BlockExchange::BlockedToWarpStriped
744+
// CHECK-NEXT: cub::BlockExchange::ScatterToBlocked
745+
// CHECK-NEXT: cub::BlockExchange::ScatterToStriped
746+
// CHECK-NEXT: cub::BlockExchange::StripedToBlocked
747+
// CHECK-NEXT: cub::BlockExchange::WarpStripedToBlocked
742748
// CHECK-NEXT: cub::DeviceReduce::ArgMax
743749
// CHECK-NEXT: cub::DeviceReduce::ArgMin
744750
// CHECK-NEXT: cub::DeviceReduce::Max

0 commit comments

Comments
 (0)