Skip to content

Commit d422526

Browse files
bagrorgkurapov-peter
bagrorg
authored and
kurapov-peter
committed
Fix formatting
1 parent 28d2754 commit d422526

7 files changed

+52
-46
lines changed

‎common/dpcpp/dpl_wrapper/dpl_wrapper.hpp

+15-14
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#pragma once
22

3-
#include <oneapi/dpl/execution>
43
#include <oneapi/dpl/algorithm>
4+
#include <oneapi/dpl/execution>
55
#include <oneapi/dpl/iterator>
66
#include <oneapi/dpl/numeric>
77

@@ -16,25 +16,26 @@
1616
namespace DPLWrapper {
1717

1818
template <class Kernel, typename T>
19-
void exclusive_scan(sycl::queue &q, sycl::buffer<T> &src, sycl::buffer<T> &dst, T init_value) {
20-
auto policy = oneapi::dpl::execution::make_device_policy<Kernel>(q);
21-
oneapi::dpl::exclusive_scan(policy, oneapi::dpl::begin(src),
22-
oneapi::dpl::end(src),
23-
oneapi::dpl::begin(dst), init_value);
19+
void exclusive_scan(sycl::queue &q, sycl::buffer<T> &src, sycl::buffer<T> &dst,
20+
T init_value) {
21+
auto policy = oneapi::dpl::execution::make_device_policy<Kernel>(q);
22+
oneapi::dpl::exclusive_scan(policy, oneapi::dpl::begin(src),
23+
oneapi::dpl::end(src), oneapi::dpl::begin(dst),
24+
init_value);
2425
}
2526

2627
template <class Kernel, typename T, typename F>
27-
void copy_if(cl::sycl::device_selector &sel, sycl::buffer<T> &src, sycl::buffer<T> &dst, F func) {
28-
auto policy = oneapi::dpl::execution::device_policy<Kernel>{sel};
29-
std::copy_if(
30-
policy, oneapi::dpl::begin(src), oneapi::dpl::end(src),
31-
oneapi::dpl::begin(dst), func);
28+
void copy_if(cl::sycl::device_selector &sel, sycl::buffer<T> &src,
29+
sycl::buffer<T> &dst, F func) {
30+
auto policy = oneapi::dpl::execution::device_policy<Kernel>{sel};
31+
std::copy_if(policy, oneapi::dpl::begin(src), oneapi::dpl::end(src),
32+
oneapi::dpl::begin(dst), func);
3233
}
3334

3435
template <class Kernel, typename T>
3536
void sort(cl::sycl::device_selector &sel, sycl::buffer<T> &src) {
36-
auto policy = oneapi::dpl::execution::device_policy<Kernel>{sel};
37-
std::sort(policy, oneapi::dpl::begin(src), oneapi::dpl::end(src));
37+
auto policy = oneapi::dpl::execution::device_policy<Kernel>{sel};
38+
std::sort(policy, oneapi::dpl::begin(src), oneapi::dpl::end(src));
3839
}
3940

40-
} //namespace DPLWrapper
41+
} // namespace DPLWrapper

‎common/dpcpp/omnisci_hashtable.hpp

+17-18
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,8 @@ std::vector<size_t> get_from_device(sycl::global_ptr<size_t> p, size_t size,
3838
}
3939

4040
namespace OmniSci {
41-
template <typename K, typename V, typename H, class ConstructorKernel> class HashTable {
41+
template <typename K, typename V, typename H, class ConstructorKernel>
42+
class HashTable {
4243
public:
4344
HashTable(const std::vector<K> &keys_vec, K empty_key, H hasher,
4445
size_t ht_size, size_t distinct_size, sycl::queue &q)
@@ -61,22 +62,22 @@ template <typename K, typename V, typename H, class ConstructorKernel> class Has
6162
auto pos = position_buffer.get_access(cgh);
6263
auto id = id_buffer.get_access(cgh);
6364

64-
cgh.parallel_for<ConstructorKernel>(std::max(ht_size, keys_vec.size()), [=](auto i) {
65-
if (i < loc_f.ht_size) {
66-
ht[i] = loc_f.empty_key;
67-
cnt[i] = 0;
68-
pos[i] = 0;
69-
}
70-
71-
if (i < loc_f.keys_size) {
72-
id[i] = 0;
73-
}
74-
});
65+
cgh.parallel_for<ConstructorKernel>(std::max(ht_size, keys_vec.size()),
66+
[=](auto i) {
67+
if (i < loc_f.ht_size) {
68+
ht[i] = loc_f.empty_key;
69+
cnt[i] = 0;
70+
pos[i] = 0;
71+
}
72+
73+
if (i < loc_f.keys_size) {
74+
id[i] = 0;
75+
}
76+
});
7577
}).wait();
7678
}
7779

78-
template <class Kernel>
79-
void build_table() {
80+
template <class Kernel> void build_table() {
8081
q.submit([&](sycl::handler &cgh) {
8182
hash_table_pimpl loc_f = *f;
8283
auto ht = hash_table.get_access(cgh);
@@ -219,8 +220,7 @@ template <typename K, typename V, typename H, class ConstructorKernel> class Has
219220
size_t get_buffer_size() { return f->ht_size * 3 + f->keys_size; }
220221

221222
private:
222-
template <class Kernel>
223-
void build_count_buffer() {
223+
template <class Kernel> void build_count_buffer() {
224224
q.submit([&](sycl::handler &cgh) {
225225
hash_table_pimpl loc_f = *f;
226226
auto ht = hash_table.get_access(cgh);
@@ -247,8 +247,7 @@ template <typename K, typename V, typename H, class ConstructorKernel> class Has
247247
}).wait();
248248
}
249249

250-
template <class Kernel>
251-
void build_pos_buffer() {
250+
template <class Kernel> void build_pos_buffer() {
252251
auto policy = oneapi::dpl::execution::make_device_policy<Kernel>(q);
253252
oneapi::dpl::exclusive_scan(policy, oneapi::dpl::begin(count_buffer),
254253
oneapi::dpl::end(count_buffer),

‎join/join_omnisci.cpp

+6-4
Original file line numberDiff line numberDiff line change
@@ -68,13 +68,15 @@ void JoinOmnisci::_run(const size_t buf_size, Meter &meter) {
6868

6969
for (unsigned it = 0; it < opts.iterations; ++it) {
7070
std::unique_ptr<HashJoinResult> result = std::make_unique<HashJoinResult>();
71-
OmniSci::HashTable<uint32_t, uint32_t, SimpleHasher<uint32_t>, class JoinOmnisciTable> ht(
72-
table_a_keys, std::numeric_limits<uint32_t>::max(), hasher, ht_size,
73-
unique_keys, q);
71+
OmniSci::HashTable<uint32_t, uint32_t, SimpleHasher<uint32_t>,
72+
class JoinOmnisciTable>
73+
ht(table_a_keys, std::numeric_limits<uint32_t>::max(), hasher, ht_size,
74+
unique_keys, q);
7475
auto host_start = std::chrono::steady_clock::now();
7576

7677
ht.build_table<class JoinOmnisciBuildTable>();
77-
ht.build_id_buffer<class JoinOmnisciBuildID, class JoinOmnisciBuildCnt, class JoinOmnisciBuildPos>();
78+
ht.build_id_buffer<class JoinOmnisciBuildID, class JoinOmnisciBuildCnt,
79+
class JoinOmnisciBuildPos>();
7880

7981
auto build_end = std::chrono::steady_clock::now();
8082

‎join/join_omnisci.hpp

-1
Original file line numberDiff line numberDiff line change
@@ -20,4 +20,3 @@ class JoinOmnisciCuda : public Dwarf {
2020
private:
2121
void _run(const size_t buffer_size, Meter &meter);
2222
};
23-

‎join/join_omnisci_cuda.cpp

+10-7
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#include "common/dpcpp/omnisci_hashtable.hpp"
22

3-
#include "join_omnisci.hpp"
43
#include "join_helpers/join_helpers.hpp"
4+
#include "join_omnisci.hpp"
55
#include <unordered_set>
66

77
using JoinOneToManySet = JoinOneToMany<std::unordered_set<size_t>>;
@@ -46,7 +46,7 @@ bool are_equal(const std::vector<JoinOneToManySet> &expected,
4646
JoinOmnisciCuda::JoinOmnisciCuda() : Dwarf("JoinOmnisciCuda") {}
4747
using namespace join_helpers;
4848
void JoinOmnisciCuda::_run(const size_t buf_size, Meter &meter) {
49-
std::cout << "CUDA" << std::endl;
49+
std::cout << "CUDA" << std::endl;
5050
auto opts = meter.opts();
5151

5252
constexpr uint32_t empty_element = std::numeric_limits<uint32_t>::max();
@@ -58,7 +58,7 @@ void JoinOmnisciCuda::_run(const size_t buf_size, Meter &meter) {
5858
helpers::make_random<uint32_t>(buf_size);
5959

6060
auto sel = get_device_selector(opts);
61-
sycl::queue q{ *sel };
61+
sycl::queue q{*sel};
6262
std::cout << "Selected device: "
6363
<< q.get_device().get_info<sycl::info::device::name>() << "\n";
6464
auto expected = build_join_id_buffer(table_a_keys, table_b_keys);
@@ -68,14 +68,17 @@ void JoinOmnisciCuda::_run(const size_t buf_size, Meter &meter) {
6868

6969
for (unsigned it = 0; it < opts.iterations; ++it) {
7070
std::unique_ptr<HashJoinResult> result = std::make_unique<HashJoinResult>();
71-
OmniSci::HashTable<uint32_t, uint32_t, SimpleHasher<uint32_t>, class JoinOmnisciCudaTable> ht(
72-
table_a_keys, std::numeric_limits<uint32_t>::max(), hasher, ht_size,
73-
unique_keys, q);
71+
OmniSci::HashTable<uint32_t, uint32_t, SimpleHasher<uint32_t>,
72+
class JoinOmnisciCudaTable>
73+
ht(table_a_keys, std::numeric_limits<uint32_t>::max(), hasher, ht_size,
74+
unique_keys, q);
7475
auto host_start = std::chrono::steady_clock::now();
7576

7677
ht.build_table<class JoinOmnisciCudaBuildTable>();
7778
std::cout << "built" << std::endl;
78-
ht.build_id_buffer<class JoinOmnisciCudaBuildID, class JoinOmnisciCudaBuildCnt, class JoinOmnisciCudaBuildPos>();
79+
ht.build_id_buffer<class JoinOmnisciCudaBuildID,
80+
class JoinOmnisciCudaBuildCnt,
81+
class JoinOmnisciCudaBuildPos>();
7982
std::cout << "id" << std::endl;
8083
auto build_end = std::chrono::steady_clock::now();
8184

‎scan/dplscan.cpp

+2-1
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,8 @@ void DPLScan::run_scan(const size_t buf_size, Meter &meter) {
3838

3939
auto host_start = std::chrono::steady_clock::now();
4040

41-
DPLWrapper::copy_if<class DPLScanKernel>(*sel, src_buf, out_buf, [](auto &x) { return x < 5; });
41+
DPLWrapper::copy_if<class DPLScanKernel>(*sel, src_buf, out_buf,
42+
[](auto &x) { return x < 5; });
4243

4344
auto host_end = std::chrono::steady_clock::now();
4445
auto host_exe_time = std::chrono::duration_cast<std::chrono::microseconds>(

‎scan/dplscan_cuda.cpp

+2-1
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,8 @@ void DPLScanCuda::run_scan(const size_t buf_size, Meter &meter) {
3838

3939
auto host_start = std::chrono::steady_clock::now();
4040

41-
DPLWrapper::copy_if<class DPLScanCudaKernel>(*sel, src_buf, out_buf, [](auto &x) { return x < 5; });
41+
DPLWrapper::copy_if<class DPLScanCudaKernel>(*sel, src_buf, out_buf,
42+
[](auto &x) { return x < 5; });
4243

4344
auto host_end = std::chrono::steady_clock::now();
4445
auto host_exe_time = std::chrono::duration_cast<std::chrono::microseconds>(

0 commit comments

Comments
 (0)