Skip to content

Commit 37eddd9

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

4 files changed

Lines changed: 47 additions & 0 deletions

File tree

‎clang/examples/DPCT/CUB/cub$$BlockLoad$$Load.cu‎

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

‎clang/examples/DPCT/CUB/cub$$BlockStore$$Store.cu‎

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

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,3 +44,25 @@
4444
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: cub::BlockExchange<int, 128, 4>(temp_storage).WarpStripedToBlocked(thread_data/*int(&)[4]*/, thread_data/*int(&)[4]*/);
4545
// CHECK_BLOCKEXCHANGE_WARPSTRIPEDTOBLOCKED: Is migrated to:
4646
// 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);
47+
48+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockLoad::Load | FileCheck %s -check-prefix=CHECK_BLOCKLOAD_LOAD
49+
// CHECK_BLOCKLOAD_LOAD: CUDA API:
50+
// CHECK_BLOCKLOAD_LOAD: __shared__ typename cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>::TempStorage temp_storage;
51+
// CHECK_BLOCKLOAD_LOAD: cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>(temp_storage).Load(src/*int **/, thread_data/*int(&)[4]*/);
52+
// CHECK_BLOCKLOAD_LOAD: cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>(temp_storage).Load(src/*int **/, thread_data/*int(&)[4]*/, end/*int*/);
53+
// CHECK_BLOCKLOAD_LOAD: cub::BlockLoad<int, 128, 4, cub::BLOCK_LOAD_DIRECT>(temp_storage).Load(src/*int **/, thread_data/*int(&)[4]*/, end/*int*/, default_value/*int*/);
54+
// CHECK_BLOCKLOAD_LOAD: Is migrated to:
55+
// CHECK_BLOCKLOAD_LOAD: auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
56+
// CHECK_BLOCKLOAD_LOAD: dpct::group::group_load<int, 4, dpct::group::group_load_algorithm::blocked>(temp_storage).load(item_ct1, src, thread_data);
57+
// CHECK_BLOCKLOAD_LOAD: dpct::group::group_load<int, 4, dpct::group::group_load_algorithm::blocked>(temp_storage).load(item_ct1, src, thread_data, end);
58+
// CHECK_BLOCKLOAD_LOAD: dpct::group::group_load<int, 4, dpct::group::group_load_algorithm::blocked>(temp_storage).load(item_ct1, src, thread_data, end, default_value);
59+
60+
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cub::BlockStore::Store | FileCheck %s -check-prefix=CHECK_BLOCKSTORE_STORE
61+
// CHECK_BLOCKSTORE_STORE: CUDA API:
62+
// CHECK_BLOCKSTORE_STORE: __shared__ typename cub::BlockStore<int, 128, 4, cub::BLOCK_STORE_DIRECT>::TempStorage temp_storage;
63+
// CHECK_BLOCKSTORE_STORE: cub::BlockStore<int, 128, 4, cub::BLOCK_STORE_DIRECT>(temp_storage).Store(dst/*int **/, thread_data/*int(&)[4]*/);
64+
// CHECK_BLOCKSTORE_STORE: cub::BlockStore<int, 128, 4, cub::BLOCK_STORE_DIRECT>(temp_storage).Store(dst/*int **/, thread_data/*int(&)[4]*/, end/*int*/);
65+
// CHECK_BLOCKSTORE_STORE: Is migrated to:
66+
// CHECK_BLOCKSTORE_STORE: auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
67+
// CHECK_BLOCKSTORE_STORE: dpct::group::group_store<int, 4, dpct::group::group_store_algorithm::blocked>(temp_storage).store(item_ct1, dst, thread_data);
68+
// CHECK_BLOCKSTORE_STORE: dpct::group::group_store<int, 4, dpct::group::group_store_algorithm::blocked>(temp_storage).store(item_ct1, dst, thread_data, end);

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -745,6 +745,8 @@
745745
// CHECK-NEXT: cub::BlockExchange::ScatterToStriped
746746
// CHECK-NEXT: cub::BlockExchange::StripedToBlocked
747747
// CHECK-NEXT: cub::BlockExchange::WarpStripedToBlocked
748+
// CHECK-NEXT: cub::BlockLoad::Load
749+
// CHECK-NEXT: cub::BlockStore::Store
748750
// CHECK-NEXT: cub::DeviceReduce::ArgMax
749751
// CHECK-NEXT: cub::DeviceReduce::ArgMin
750752
// CHECK-NEXT: cub::DeviceReduce::Max

0 commit comments

Comments
 (0)