Skip to content

Commit 21ec5c2

Browse files
vlad-perevezentsevvlad-perevezentsev
committed
Merge master into update_cuda_build
2 parents 2ee05d2 + eed5274 commit 21ec5c2

File tree

128 files changed

+767
-704
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

128 files changed

+767
-704
lines changed

‎dpctl/_host_task_util.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -77,19 +77,19 @@ DPCTLSyclEventRef async_dec_ref(DPCTLSyclQueueRef QRef,
7777
});
7878
});
7979

80-
constexpr int result_ok = 0;
80+
static constexpr int result_ok = 0;
8181

8282
*status = result_ok;
8383
auto e_ptr = new sycl::event(ht_ev);
8484
return wrap<sycl::event>(e_ptr);
8585
} catch (const std::exception &e) {
86-
constexpr int result_std_exception = 1;
86+
static constexpr int result_std_exception = 1;
8787

8888
*status = result_std_exception;
8989
return nullptr;
9090
}
9191

92-
constexpr int result_other_abnormal = 2;
92+
static constexpr int result_other_abnormal = 2;
9393

9494
*status = result_other_abnormal;
9595
return nullptr;

‎dpctl/tensor/libtensor/include/kernels/accumulators.hpp‎

Lines changed: 47 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -60,9 +60,9 @@ template <typename inputT, typename outputT> struct NonZeroIndicator
6060

6161
outputT operator()(const inputT &val) const
6262
{
63-
constexpr outputT out_one(1);
64-
constexpr outputT out_zero(0);
65-
constexpr inputT val_zero(0);
63+
static constexpr outputT out_one(1);
64+
static constexpr outputT out_zero(0);
65+
static constexpr inputT val_zero(0);
6666

6767
return (val == val_zero) ? out_zero : out_one;
6868
}
@@ -583,7 +583,7 @@ sycl::event update_local_chunks_1d(sycl::queue &exec_q,
583583
cgh.depends_on(dependent_event);
584584
cgh.use_kernel_bundle(kb);
585585

586-
constexpr nwiT updates_per_wi = n_wi;
586+
static constexpr nwiT updates_per_wi = n_wi;
587587
const std::size_t n_items =
588588
ceiling_quotient<std::size_t>(src_size, sg_size * n_wi) * sg_size;
589589

@@ -594,8 +594,8 @@ sycl::event update_local_chunks_1d(sycl::queue &exec_q,
594594
cgh.parallel_for<UpdateKernelName>(
595595
ndRange,
596596
[chunk_size, src, src_size, local_scans](sycl::nd_item<1> ndit) {
597-
constexpr ScanOpT scan_op{};
598-
constexpr outputT identity =
597+
static constexpr ScanOpT scan_op{};
598+
static constexpr outputT identity =
599599
su_ns::Identity<ScanOpT, outputT>::value;
600600

601601
const std::uint32_t lws = ndit.get_local_range(0);
@@ -640,16 +640,17 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
640640
std::vector<sycl::event> &host_tasks,
641641
const std::vector<sycl::event> &depends = {})
642642
{
643-
constexpr ScanOpT scan_op{};
644-
constexpr outputT identity = su_ns::Identity<ScanOpT, outputT>::value;
643+
static constexpr ScanOpT scan_op{};
644+
static constexpr outputT identity =
645+
su_ns::Identity<ScanOpT, outputT>::value;
645646

646-
constexpr std::size_t _iter_nelems = 1;
647+
static constexpr std::size_t _iter_nelems = 1;
647648

648649
using IterIndexerT = dpctl::tensor::offset_utils::TwoZeroOffsets_Indexer;
649-
constexpr IterIndexerT _no_op_iter_indexer{};
650+
static constexpr IterIndexerT _no_op_iter_indexer{};
650651

651652
using NoOpIndexerT = dpctl::tensor::offset_utils::NoOpIndexer;
652-
constexpr NoOpIndexerT _no_op_indexer{};
653+
static constexpr NoOpIndexerT _no_op_indexer{};
653654

654655
std::size_t n_groups;
655656
sycl::event inc_scan_phase1_ev =
@@ -687,7 +688,7 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
687688
outputT *local_scans = temp;
688689

689690
using NoOpTransformerT = NoOpTransformer<outputT>;
690-
constexpr NoOpTransformerT _no_op_transformer{};
691+
static constexpr NoOpTransformerT _no_op_transformer{};
691692
std::size_t size_to_update = n_elems;
692693
while (n_groups_ > 1) {
693694

@@ -761,16 +762,16 @@ accumulate_1d_contig_impl(sycl::queue &q,
761762
dstT *dst_data_ptr = reinterpret_cast<dstT *>(dst);
762763

763764
using NoOpIndexerT = dpctl::tensor::offset_utils::NoOpIndexer;
764-
constexpr NoOpIndexerT flat_indexer{};
765-
constexpr transformerT transformer{};
765+
static constexpr NoOpIndexerT flat_indexer{};
766+
static constexpr transformerT transformer{};
766767

767-
constexpr std::size_t s0 = 0;
768-
constexpr std::size_t s1 = 1;
768+
static constexpr std::size_t s0 = 0;
769+
static constexpr std::size_t s1 = 1;
769770

770771
sycl::event comp_ev;
771772
const sycl::device &dev = q.get_device();
772773
if (dev.has(sycl::aspect::cpu)) {
773-
constexpr nwiT n_wi_for_cpu = 8;
774+
static constexpr nwiT n_wi_for_cpu = 8;
774775
const std::uint32_t wg_size = 256;
775776
comp_ev = inclusive_scan_iter_1d<srcT, dstT, n_wi_for_cpu, NoOpIndexerT,
776777
transformerT, AccumulateOpT,
@@ -779,7 +780,7 @@ accumulate_1d_contig_impl(sycl::queue &q,
779780
flat_indexer, transformer, host_tasks, depends);
780781
}
781782
else {
782-
constexpr nwiT n_wi_for_gpu = 4;
783+
static constexpr nwiT n_wi_for_gpu = 4;
783784
// base_scan_striped algorithm does not execute correctly
784785
// on HIP device with wg_size > 64
785786
const std::uint32_t wg_size =
@@ -829,7 +830,7 @@ sycl::event final_update_local_chunks(sycl::queue &exec_q,
829830
const std::uint32_t sg_size = krn.template get_info<
830831
sycl::info::kernel_device_specific::max_sub_group_size>(dev);
831832

832-
constexpr nwiT updates_per_wi = n_wi;
833+
static constexpr nwiT updates_per_wi = n_wi;
833834
const std::size_t updates_per_sg = sg_size * updates_per_wi;
834835
const std::size_t update_nelems =
835836
ceiling_quotient(src_size, updates_per_sg) * sg_size;
@@ -845,8 +846,8 @@ sycl::event final_update_local_chunks(sycl::queue &exec_q,
845846
cgh.parallel_for<UpdateKernelName>(
846847
ndRange, [chunk_size, src_size, local_stride, src, local_scans,
847848
out_iter_indexer, out_indexer](sycl::nd_item<2> ndit) {
848-
constexpr ScanOpT scan_op{};
849-
constexpr outputT identity =
849+
static constexpr ScanOpT scan_op{};
850+
static constexpr outputT identity =
850851
su_ns::Identity<ScanOpT, outputT>::value;
851852

852853
const std::uint32_t lws = ndit.get_local_range(1);
@@ -898,8 +899,8 @@ sycl::event update_local_chunks(sycl::queue &exec_q,
898899
std::size_t local_stride,
899900
sycl::event dependent_event)
900901
{
901-
constexpr NoOpIndexer out_indexer{};
902-
constexpr NoOpIndexer iter_out_indexer{};
902+
static constexpr NoOpIndexer out_indexer{};
903+
static constexpr NoOpIndexer iter_out_indexer{};
903904

904905
return final_update_local_chunks<UpdateKernelName, outputT, n_wi,
905906
NoOpIndexer, NoOpIndexer, ScanOpT>(
@@ -933,8 +934,9 @@ sycl::event inclusive_scan_iter(sycl::queue &exec_q,
933934
std::vector<sycl::event> &host_tasks,
934935
const std::vector<sycl::event> &depends = {})
935936
{
936-
constexpr ScanOpT scan_op{};
937-
constexpr outputT identity = su_ns::Identity<ScanOpT, outputT>::value;
937+
static constexpr ScanOpT scan_op{};
938+
static constexpr outputT identity =
939+
su_ns::Identity<ScanOpT, outputT>::value;
938940

939941
using IterIndexerT =
940942
dpctl::tensor::offset_utils::TwoOffsets_CombinedIndexer<
@@ -977,9 +979,9 @@ sycl::event inclusive_scan_iter(sycl::queue &exec_q,
977979
outputT *local_scans = temp;
978980

979981
using NoOpIndexerT = dpctl::tensor::offset_utils::NoOpIndexer;
980-
constexpr NoOpIndexerT _no_op_indexer{};
982+
static constexpr NoOpIndexerT _no_op_indexer{};
981983
using NoOpTransformerT = NoOpTransformer<outputT>;
982-
constexpr NoOpTransformerT _no_op_transformer{};
984+
static constexpr NoOpTransformerT _no_op_transformer{};
983985
std::size_t size_to_update = acc_nelems;
984986

985987
{
@@ -1142,15 +1144,15 @@ accumulate_strided_impl(sycl::queue &q,
11421144
iter_shape_strides,
11431145
iter_shape_strides + 2 * iter_nd};
11441146

1145-
constexpr transformerT transformer{};
1147+
static constexpr transformerT transformer{};
11461148

1147-
constexpr std::size_t s0 = 0;
1148-
constexpr std::size_t s1 = 1;
1149+
static constexpr std::size_t s0 = 0;
1150+
static constexpr std::size_t s1 = 1;
11491151

11501152
const sycl::device &dev = q.get_device();
11511153
sycl::event comp_ev;
11521154
if (dev.has(sycl::aspect::cpu)) {
1153-
constexpr nwiT n_wi_for_cpu = 8;
1155+
static constexpr nwiT n_wi_for_cpu = 8;
11541156
const std::uint32_t wg_size = 256;
11551157
comp_ev =
11561158
inclusive_scan_iter<srcT, dstT, n_wi_for_cpu, InpIndexerT,
@@ -1161,7 +1163,7 @@ accumulate_strided_impl(sycl::queue &q,
11611163
out_axis_indexer, transformer, host_tasks, depends);
11621164
}
11631165
else {
1164-
constexpr nwiT n_wi_for_gpu = 4;
1166+
static constexpr nwiT n_wi_for_gpu = 4;
11651167
// base_scan_striped algorithm does not execute correctly
11661168
// on HIP device with wg_size > 64
11671169
const std::uint32_t wg_size =
@@ -1198,18 +1200,18 @@ std::size_t cumsum_val_contig_impl(sycl::queue &q,
11981200
cumsumT *cumsum_data_ptr = reinterpret_cast<cumsumT *>(cumsum);
11991201

12001202
using NoOpIndexerT = dpctl::tensor::offset_utils::NoOpIndexer;
1201-
constexpr NoOpIndexerT flat_indexer{};
1202-
constexpr transformerT transformer{};
1203+
static constexpr NoOpIndexerT flat_indexer{};
1204+
static constexpr transformerT transformer{};
12031205

1204-
constexpr std::size_t s0 = 0;
1205-
constexpr std::size_t s1 = 1;
1206-
constexpr bool include_initial = false;
1206+
static constexpr std::size_t s0 = 0;
1207+
static constexpr std::size_t s1 = 1;
1208+
static constexpr bool include_initial = false;
12071209
using AccumulateOpT = sycl::plus<cumsumT>;
12081210

12091211
sycl::event comp_ev;
12101212
const sycl::device &dev = q.get_device();
12111213
if (dev.has(sycl::aspect::cpu)) {
1212-
constexpr nwiT n_wi_for_cpu = 8;
1214+
static constexpr nwiT n_wi_for_cpu = 8;
12131215
const std::uint32_t wg_size = 256;
12141216
comp_ev = inclusive_scan_iter_1d<maskT, cumsumT, n_wi_for_cpu,
12151217
NoOpIndexerT, transformerT,
@@ -1218,7 +1220,7 @@ std::size_t cumsum_val_contig_impl(sycl::queue &q,
12181220
flat_indexer, transformer, host_tasks, depends);
12191221
}
12201222
else {
1221-
constexpr nwiT n_wi_for_gpu = 4;
1223+
static constexpr nwiT n_wi_for_gpu = 4;
12221224
// base_scan_striped algorithm does not execute correctly
12231225
// on HIP device with wg_size > 64
12241226
const std::uint32_t wg_size =
@@ -1313,17 +1315,17 @@ cumsum_val_strided_impl(sycl::queue &q,
13131315

13141316
using StridedIndexerT = dpctl::tensor::offset_utils::StridedIndexer;
13151317
const StridedIndexerT strided_indexer{nd, 0, shape_strides};
1316-
constexpr transformerT transformer{};
1318+
static constexpr transformerT transformer{};
13171319

1318-
constexpr std::size_t s0 = 0;
1319-
constexpr std::size_t s1 = 1;
1320-
constexpr bool include_initial = false;
1320+
static constexpr std::size_t s0 = 0;
1321+
static constexpr std::size_t s1 = 1;
1322+
static constexpr bool include_initial = false;
13211323
using AccumulateOpT = sycl::plus<cumsumT>;
13221324

13231325
const sycl::device &dev = q.get_device();
13241326
sycl::event comp_ev;
13251327
if (dev.has(sycl::aspect::cpu)) {
1326-
constexpr nwiT n_wi_for_cpu = 8;
1328+
static constexpr nwiT n_wi_for_cpu = 8;
13271329
const std::uint32_t wg_size = 256;
13281330
comp_ev = inclusive_scan_iter_1d<maskT, cumsumT, n_wi_for_cpu,
13291331
StridedIndexerT, transformerT,
@@ -1332,7 +1334,7 @@ cumsum_val_strided_impl(sycl::queue &q,
13321334
strided_indexer, transformer, host_tasks, depends);
13331335
}
13341336
else {
1335-
constexpr nwiT n_wi_for_gpu = 4;
1337+
static constexpr nwiT n_wi_for_gpu = 4;
13361338
// base_scan_striped algorithm does not execute correctly
13371339
// on HIP device with wg_size > 64
13381340
const std::uint32_t wg_size =

‎dpctl/tensor/libtensor/include/kernels/alignment.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ namespace kernels
3030
namespace alignment_utils
3131
{
3232

33-
static constexpr std::size_t required_alignment = 64UL;
33+
inline constexpr std::size_t required_alignment = 64UL;
3434

3535
template <std::uintptr_t alignment, typename Ptr> bool is_aligned(Ptr p)
3636
{

‎dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp‎

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -227,11 +227,11 @@ std::size_t _get_lws_impl(std::size_t n)
227227
}
228228
}
229229

230-
std::size_t get_lws(std::size_t n)
230+
inline std::size_t get_lws(std::size_t n)
231231
{
232-
constexpr std::size_t lws0 = 256u;
233-
constexpr std::size_t lws1 = 128u;
234-
constexpr std::size_t lws2 = 64u;
232+
static constexpr std::size_t lws0 = 256u;
233+
static constexpr std::size_t lws1 = 128u;
234+
static constexpr std::size_t lws2 = 64u;
235235
return _get_lws_impl<lws0, lws1, lws2>(n);
236236
}
237237

@@ -261,9 +261,9 @@ sycl::event masked_extract_all_slices_contig_impl(
261261
ssize_t dst_stride,
262262
const std::vector<sycl::event> &depends = {})
263263
{
264-
constexpr TwoZeroOffsets_Indexer orthog_src_dst_indexer{};
264+
static constexpr TwoZeroOffsets_Indexer orthog_src_dst_indexer{};
265265

266-
constexpr NoOpIndexer masked_src_indexer{};
266+
static constexpr NoOpIndexer masked_src_indexer{};
267267
const Strided1DIndexer masked_dst_indexer(/* size */ dst_size,
268268
/* step */ dst_stride);
269269

@@ -339,7 +339,7 @@ sycl::event masked_extract_all_slices_strided_impl(
339339
ssize_t dst_stride,
340340
const std::vector<sycl::event> &depends = {})
341341
{
342-
constexpr TwoZeroOffsets_Indexer orthog_src_dst_indexer{};
342+
static constexpr TwoZeroOffsets_Indexer orthog_src_dst_indexer{};
343343

344344
/* StridedIndexer(int _nd, ssize_t _offset, ssize_t const
345345
* *_packed_shape_strides) */
@@ -578,7 +578,7 @@ sycl::event masked_place_all_slices_strided_impl(
578578
ssize_t rhs_stride,
579579
const std::vector<sycl::event> &depends = {})
580580
{
581-
constexpr TwoZeroOffsets_Indexer orthog_dst_rhs_indexer{};
581+
static constexpr TwoZeroOffsets_Indexer orthog_dst_rhs_indexer{};
582582

583583
/* StridedIndexer(int _nd, ssize_t _offset, ssize_t const
584584
* *_packed_shape_strides) */
@@ -589,7 +589,7 @@ sycl::event masked_place_all_slices_strided_impl(
589589
TwoZeroOffsets_Indexer, StridedIndexer, Strided1DCyclicIndexer, dataT,
590590
indT>;
591591

592-
constexpr std::size_t nominal_lws = 256;
592+
static constexpr std::size_t nominal_lws = 256;
593593
const std::size_t masked_extent = iteration_size;
594594
const std::size_t lws = std::min(masked_extent, nominal_lws);
595595

@@ -685,7 +685,7 @@ sycl::event masked_place_some_slices_strided_impl(
685685
TwoOffsets_StridedIndexer, StridedIndexer, Strided1DCyclicIndexer,
686686
dataT, indT>;
687687

688-
constexpr std::size_t nominal_lws = 256;
688+
static constexpr std::size_t nominal_lws = 256;
689689
const std::size_t orthog_extent = orthog_nelems;
690690
const std::size_t masked_extent = masked_nelems;
691691
const std::size_t lws = std::min(masked_extent, nominal_lws);
@@ -788,7 +788,7 @@ sycl::event non_zero_indexes_impl(sycl::queue &exec_q,
788788
const indT1 *cumsum_data = reinterpret_cast<const indT1 *>(cumsum_cp);
789789
indT2 *indexes_data = reinterpret_cast<indT2 *>(indexes_cp);
790790

791-
constexpr std::size_t nominal_lws = 256u;
791+
static constexpr std::size_t nominal_lws = 256u;
792792
const std::size_t masked_extent = iter_size;
793793
const std::size_t lws = std::min(masked_extent, nominal_lws);
794794

‎dpctl/tensor/libtensor/include/kernels/clip.hpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -106,7 +106,7 @@ class ClipContigFunctor
106106

107107
void operator()(sycl::nd_item<1> ndit) const
108108
{
109-
constexpr std::uint8_t nelems_per_wi = n_vecs * vec_sz;
109+
static constexpr std::uint8_t nelems_per_wi = n_vecs * vec_sz;
110110

111111
using dpctl::tensor::type_utils::is_complex;
112112
if constexpr (is_complex<T>::value || !enable_sg_loadstore) {
@@ -202,8 +202,8 @@ sycl::event clip_contig_impl(sycl::queue &q,
202202
cgh.depends_on(depends);
203203

204204
std::size_t lws = 64;
205-
constexpr std::uint8_t vec_sz = 4;
206-
constexpr std::uint8_t n_vecs = 2;
205+
static constexpr std::uint8_t vec_sz = 4;
206+
static constexpr std::uint8_t n_vecs = 2;
207207
const std::size_t n_groups =
208208
((nelems + lws * n_vecs * vec_sz - 1) / (lws * n_vecs * vec_sz));
209209
const auto gws_range = sycl::range<1>(n_groups * lws);
@@ -214,7 +214,7 @@ sycl::event clip_contig_impl(sycl::queue &q,
214214
is_aligned<required_alignment>(max_cp) &&
215215
is_aligned<required_alignment>(dst_cp))
216216
{
217-
constexpr bool enable_sg_loadstore = true;
217+
static constexpr bool enable_sg_loadstore = true;
218218
using KernelName = clip_contig_kernel<T, vec_sz, n_vecs>;
219219
using Impl =
220220
ClipContigFunctor<T, vec_sz, n_vecs, enable_sg_loadstore>;
@@ -224,7 +224,7 @@ sycl::event clip_contig_impl(sycl::queue &q,
224224
Impl(nelems, x_tp, min_tp, max_tp, dst_tp));
225225
}
226226
else {
227-
constexpr bool disable_sg_loadstore = false;
227+
static constexpr bool disable_sg_loadstore = false;
228228
using InnerKernelName = clip_contig_kernel<T, vec_sz, n_vecs>;
229229
using KernelName =
230230
disabled_sg_loadstore_wrapper_krn<InnerKernelName>;

‎dpctl/tensor/libtensor/include/kernels/constructors.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -468,7 +468,7 @@ sycl::event tri_impl(sycl::queue &exec_q,
468468
const std::vector<sycl::event> &depends,
469469
const std::vector<sycl::event> &additional_depends)
470470
{
471-
constexpr int d2 = 2;
471+
static constexpr int d2 = 2;
472472
ssize_t src_s = nd;
473473
ssize_t dst_s = 2 * nd;
474474
ssize_t nd_1 = nd - 1;

0 commit comments

Comments
 (0)