|
1 |
| -#include <CL/sycl.hpp> |
2 | 1 | #include "use_sycl_buffer.h"
|
3 |
| -#include <oneapi/mkl.hpp> |
4 | 2 | #include "dpctl_sycl_types.h"
|
| 3 | +#include <CL/sycl.hpp> |
| 4 | +#include <oneapi/mkl.hpp> |
5 | 5 |
|
6 |
| -int |
7 |
| -c_columnwise_total(DPCTLSyclQueueRef q_ref, size_t n, size_t m, double *mat, double *ct) { |
| 6 | +int c_columnwise_total(DPCTLSyclQueueRef q_ref, |
| 7 | + size_t n, |
| 8 | + size_t m, |
| 9 | + double *mat, |
| 10 | + double *ct) |
| 11 | +{ |
8 | 12 |
|
9 | 13 | sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));
|
10 | 14 |
|
11 |
| - sycl::buffer<double, 1> mat_buffer = sycl::buffer(mat, sycl::range<1>(n * m)); |
| 15 | + sycl::buffer<double, 1> mat_buffer = |
| 16 | + sycl::buffer(mat, sycl::range<1>(n * m)); |
12 | 17 | sycl::buffer<double, 1> ct_buffer = sycl::buffer(ct, sycl::range<1>(m));
|
13 | 18 |
|
14 | 19 | double *ones = reinterpret_cast<double *>(malloc(n * sizeof(double)));
|
15 | 20 | {
|
16 |
| - sycl::buffer<double, 1> ones_buffer = sycl::buffer(ones, sycl::range<1>(n)); |
17 |
| - |
18 |
| - try { |
19 |
| - auto ev = q.submit([&](sycl::handler &cgh) { |
20 |
| - auto ones_acc = ones_buffer.get_access<sycl::access::mode::read_write>(cgh); |
21 |
| - cgh.fill(ones_acc, double(1.0)); |
22 |
| - }); |
23 |
| - |
24 |
| - ev.wait_and_throw(); |
25 |
| - } |
26 |
| - catch (sycl::exception const& e) { |
27 |
| - std::cout << "\t\tCaught synchronous SYCL exception during fill:\n" |
28 |
| - << e.what() << std::endl << "OpenCL status: " << e.get_cl_code() << std::endl; |
29 |
| - goto cleanup; |
30 |
| - } |
31 |
| - |
32 |
| - try { |
33 |
| - oneapi::mkl::blas::row_major::gemv( |
34 |
| - q, |
35 |
| - oneapi::mkl::transpose::trans, |
36 |
| - n, m, double(1.0), mat_buffer, m, |
37 |
| - ones_buffer, 1, |
38 |
| - double(0.0), ct_buffer, 1); |
39 |
| - q.wait(); |
40 |
| - } |
41 |
| - catch (sycl::exception const &e) { |
42 |
| - std::cout << "\t\tCaught synchronous SYCL exception during GEMV:\n" |
43 |
| - << e.what() << std::endl << "OpenCL status: " << e.get_cl_code() << std::endl; |
44 |
| - goto cleanup; |
45 |
| - } |
| 21 | + sycl::buffer<double, 1> ones_buffer = |
| 22 | + sycl::buffer(ones, sycl::range<1>(n)); |
| 23 | + |
| 24 | + try { |
| 25 | + auto ev = q.submit([&](sycl::handler &cgh) { |
| 26 | + auto ones_acc = |
| 27 | + ones_buffer.get_access<sycl::access::mode::read_write>(cgh); |
| 28 | + cgh.fill(ones_acc, double(1.0)); |
| 29 | + }); |
| 30 | + |
| 31 | + ev.wait_and_throw(); |
| 32 | + } catch (sycl::exception const &e) { |
| 33 | + std::cout << "\t\tCaught synchronous SYCL exception during fill:\n" |
| 34 | + << e.what() << std::endl |
| 35 | + << "OpenCL status: " << e.get_cl_code() << std::endl; |
| 36 | + goto cleanup; |
| 37 | + } |
| 38 | + |
| 39 | + try { |
| 40 | + oneapi::mkl::blas::row_major::gemv( |
| 41 | + q, oneapi::mkl::transpose::trans, n, m, double(1.0), mat_buffer, |
| 42 | + m, ones_buffer, 1, double(0.0), ct_buffer, 1); |
| 43 | + q.wait(); |
| 44 | + } catch (sycl::exception const &e) { |
| 45 | + std::cout << "\t\tCaught synchronous SYCL exception during GEMV:\n" |
| 46 | + << e.what() << std::endl |
| 47 | + << "OpenCL status: " << e.get_cl_code() << std::endl; |
| 48 | + goto cleanup; |
| 49 | + } |
46 | 50 | }
|
47 |
| - |
| 51 | + |
48 | 52 | free(ones);
|
49 | 53 | return 0;
|
50 | 54 |
|
51 |
| - cleanup: |
| 55 | +cleanup: |
52 | 56 | free(ones);
|
53 | 57 | return -1;
|
54 | 58 | }
|
55 | 59 |
|
56 |
| -inline size_t upper_multiple(size_t n, size_t wg) { return wg * ((n + wg - 1)/wg); } |
| 60 | +inline size_t upper_multiple(size_t n, size_t wg) |
| 61 | +{ |
| 62 | + return wg * ((n + wg - 1) / wg); |
| 63 | +} |
57 | 64 |
|
58 |
| -int |
59 |
| -c_columnwise_total_no_mkl(DPCTLSyclQueueRef q_ref, size_t n, size_t m, double *mat, double *ct) { |
| 65 | +int c_columnwise_total_no_mkl(DPCTLSyclQueueRef q_ref, |
| 66 | + size_t n, |
| 67 | + size_t m, |
| 68 | + double *mat, |
| 69 | + double *ct) |
| 70 | +{ |
60 | 71 |
|
61 | 72 | sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));
|
62 | 73 |
|
63 |
| - sycl::buffer<double, 2> mat_buffer = sycl::buffer(mat, sycl::range<2>(n, m)); |
| 74 | + sycl::buffer<double, 2> mat_buffer = |
| 75 | + sycl::buffer(mat, sycl::range<2>(n, m)); |
64 | 76 | sycl::buffer<double, 1> ct_buffer = sycl::buffer(ct, sycl::range<1>(m));
|
65 | 77 |
|
66 |
| - auto e = q.submit( |
67 |
| - [&](sycl::handler &h) { |
68 |
| - sycl::accessor ct_acc {ct_buffer, h, sycl::write_only}; |
69 |
| - h.parallel_for( |
70 |
| - sycl::range<1>(m), |
71 |
| - [=](sycl::id<1> i){ |
72 |
| - ct_acc[i] = 0.0; |
73 |
| - }); |
74 |
| - }); |
| 78 | + auto e = q.submit([&](sycl::handler &h) { |
| 79 | + sycl::accessor ct_acc{ct_buffer, h, sycl::write_only}; |
| 80 | + h.parallel_for(sycl::range<1>(m), |
| 81 | + [=](sycl::id<1> i) { ct_acc[i] = 0.0; }); |
| 82 | + }); |
75 | 83 |
|
76 | 84 | constexpr size_t wg = 256;
|
77 |
| - auto e2 = q.submit( |
78 |
| - [&](sycl::handler &h) { |
79 |
| - |
80 |
| - sycl::accessor mat_acc {mat_buffer, h, sycl::read_only}; |
81 |
| - sycl::accessor ct_acc {ct_buffer, h}; |
82 |
| - h.depends_on(e); |
83 |
| - |
84 |
| - sycl::range<2> global {upper_multiple(n, wg), m}; |
85 |
| - sycl::range<2> local {wg, 1}; |
86 |
| - |
87 |
| - h.parallel_for( |
88 |
| - sycl::nd_range<2>(global, local), |
89 |
| - [=](sycl::nd_item<2> it) { |
90 |
| - size_t i = it.get_global_id(0); |
91 |
| - size_t j = it.get_global_id(1); |
92 |
| - double group_sum = sycl::ONEAPI::reduce( |
93 |
| - it.get_group(), |
94 |
| - (i < n) ? mat_acc[it.get_global_id()] : 0.0, |
95 |
| - std::plus<double>() |
96 |
| - ); |
97 |
| - if (it.get_local_id(0) == 0) { |
98 |
| - sycl::ONEAPI::atomic_ref< |
99 |
| - double, |
100 |
| - sycl::ONEAPI::memory_order::relaxed, |
101 |
| - sycl::ONEAPI::memory_scope::system, |
102 |
| - sycl::access::address_space::global_space>(ct_acc[j]) += group_sum; |
103 |
| - } |
104 |
| - }); |
105 |
| - }); |
| 85 | + auto e2 = q.submit([&](sycl::handler &h) { |
| 86 | + sycl::accessor mat_acc{mat_buffer, h, sycl::read_only}; |
| 87 | + sycl::accessor ct_acc{ct_buffer, h}; |
| 88 | + h.depends_on(e); |
| 89 | + |
| 90 | + sycl::range<2> global{upper_multiple(n, wg), m}; |
| 91 | + sycl::range<2> local{wg, 1}; |
| 92 | + |
| 93 | + h.parallel_for( |
| 94 | + sycl::nd_range<2>(global, local), [=](sycl::nd_item<2> it) { |
| 95 | + size_t i = it.get_global_id(0); |
| 96 | + size_t j = it.get_global_id(1); |
| 97 | + double group_sum = sycl::ONEAPI::reduce( |
| 98 | + it.get_group(), (i < n) ? mat_acc[it.get_global_id()] : 0.0, |
| 99 | + std::plus<double>()); |
| 100 | + if (it.get_local_id(0) == 0) { |
| 101 | + sycl::ONEAPI::atomic_ref< |
| 102 | + double, sycl::ONEAPI::memory_order::relaxed, |
| 103 | + sycl::ONEAPI::memory_scope::system, |
| 104 | + sycl::access::address_space::global_space>(ct_acc[j]) += |
| 105 | + group_sum; |
| 106 | + } |
| 107 | + }); |
| 108 | + }); |
106 | 109 |
|
107 | 110 | e2.wait_and_throw();
|
108 | 111 | return 0;
|
|
0 commit comments