Skip to content

Commit 429bdbc

Browse files
oleksandr-pavlykoleksandr-pavlyk
committed
Replace sycl::free with sycl_free_noexcept
Add implementation of dpctl::tensor::offset_utils::sycl_free_noexcept that calls sycl::free from within try/catch. The exception is logged to std::cerr, but otherwise ignored.
1 parent 43f5aea commit 429bdbc

23 files changed

+217
-110
lines changed

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

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -436,7 +436,8 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
436436
sycl::event free_ev = exec_q.submit([&](sycl::handler &cgh) {
437437
cgh.depends_on(dependent_event);
438438
const auto &ctx = exec_q.get_context();
439-
cgh.host_task([ctx, temp]() { sycl::free(temp, ctx); });
439+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
440+
cgh.host_task([ctx, temp]() { sycl_free_noexcept(temp, ctx); });
440441
});
441442
host_tasks.push_back(free_ev);
442443
}
@@ -765,7 +766,8 @@ sycl::event inclusive_scan_iter(sycl::queue &exec_q,
765766
sycl::event free_ev = exec_q.submit([&](sycl::handler &cgh) {
766767
cgh.depends_on(dependent_event);
767768
const auto &ctx = exec_q.get_context();
768-
cgh.host_task([ctx, temp]() { sycl::free(temp, ctx); });
769+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
770+
cgh.host_task([ctx, temp]() { sycl_free_noexcept(temp, ctx); });
769771
});
770772
host_tasks.push_back(free_ev);
771773
}
@@ -917,7 +919,9 @@ size_t cumsum_val_contig_impl(sycl::queue &q,
917919
});
918920
copy_e.wait();
919921
size_t return_val = static_cast<size_t>(*last_elem_host_usm);
920-
sycl::free(last_elem_host_usm, q);
922+
923+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
924+
sycl_free_noexcept(last_elem_host_usm, q);
921925

922926
return return_val;
923927
}
@@ -1026,7 +1030,9 @@ size_t cumsum_val_strided_impl(sycl::queue &q,
10261030
});
10271031
copy_e.wait();
10281032
size_t return_val = static_cast<size_t>(*last_elem_host_usm);
1029-
sycl::free(last_elem_host_usm, q);
1033+
1034+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
1035+
sycl_free_noexcept(last_elem_host_usm, q);
10301036

10311037
return return_val;
10321038
}

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -947,7 +947,9 @@ sycl::event binary_contig_matrix_contig_row_broadcast_impl(
947947
sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
948948
cgh.depends_on(comp_ev);
949949
const sycl::context &ctx = exec_q.get_context();
950-
cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
950+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
951+
cgh.host_task(
952+
[ctx, padded_vec]() { sycl_free_noexcept(padded_vec, ctx); });
951953
});
952954
host_tasks.push_back(tmp_cleanup_ev);
953955

@@ -1026,7 +1028,9 @@ sycl::event binary_contig_row_contig_matrix_broadcast_impl(
10261028
sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
10271029
cgh.depends_on(comp_ev);
10281030
const sycl::context &ctx = exec_q.get_context();
1029-
cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
1031+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
1032+
cgh.host_task(
1033+
[ctx, padded_vec]() { sycl_free_noexcept(padded_vec, ctx); });
10301034
});
10311035
host_tasks.push_back(tmp_cleanup_ev);
10321036

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,7 @@
3030

3131
#include "kernels/alignment.hpp"
3232
#include "kernels/dpctl_tensor_types.hpp"
33+
#include "utils/offset_utils.hpp"
3334

3435
namespace dpctl
3536
{
@@ -458,7 +459,9 @@ sycl::event binary_inplace_row_matrix_broadcast_impl(
458459
sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
459460
cgh.depends_on(comp_ev);
460461
const sycl::context &ctx = exec_q.get_context();
461-
cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
462+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
463+
cgh.host_task(
464+
[ctx, padded_vec]() { sycl_free_noexcept(padded_vec, ctx); });
462465
});
463466
host_tasks.push_back(tmp_cleanup_ev);
464467

‎dpctl/tensor/libtensor/include/kernels/linalg_functions/dot_product.hpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1153,8 +1153,9 @@ sycl::event dot_product_tree_impl(sycl::queue &exec_q,
11531153
cgh.depends_on(final_reduction_ev);
11541154
const sycl::context &ctx = exec_q.get_context();
11551155

1156+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
11561157
cgh.host_task([ctx, partially_reduced_tmp] {
1157-
sycl::free(partially_reduced_tmp, ctx);
1158+
sycl_free_noexcept(partially_reduced_tmp, ctx);
11581159
});
11591160
});
11601161

@@ -1403,8 +1404,9 @@ dot_product_contig_tree_impl(sycl::queue &exec_q,
14031404
cgh.depends_on(final_reduction_ev);
14041405
const sycl::context &ctx = exec_q.get_context();
14051406

1407+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
14061408
cgh.host_task([ctx, partially_reduced_tmp] {
1407-
sycl::free(partially_reduced_tmp, ctx);
1409+
sycl_free_noexcept(partially_reduced_tmp, ctx);
14081410
});
14091411
});
14101412

‎dpctl/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp

Lines changed: 32 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -2364,7 +2364,8 @@ gemm_batch_tree_k_impl(sycl::queue &exec_q,
23642364
cgh.depends_on(red_ev);
23652365
const sycl::context &ctx = exec_q.get_context();
23662366

2367-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
2367+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
2368+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
23682369
});
23692370
return cleanup_host_task_event;
23702371
}
@@ -2427,8 +2428,9 @@ gemm_batch_tree_k_impl(sycl::queue &exec_q,
24272428
cgh.depends_on(red_ev);
24282429
const sycl::context &ctx = exec_q.get_context();
24292430

2431+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
24302432
cgh.host_task([ctx, partially_reduced_tmp] {
2431-
sycl::free(partially_reduced_tmp, ctx);
2433+
sycl_free_noexcept(partially_reduced_tmp, ctx);
24322434
});
24332435
});
24342436

@@ -2661,7 +2663,8 @@ gemm_batch_tree_nm_impl(sycl::queue &exec_q,
26612663
cgh.depends_on(red_ev);
26622664
const sycl::context &ctx = exec_q.get_context();
26632665

2664-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
2666+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
2667+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
26652668
});
26662669
return cleanup_host_task_event;
26672670
}
@@ -2728,8 +2731,9 @@ gemm_batch_tree_nm_impl(sycl::queue &exec_q,
27282731
cgh.depends_on(red_ev);
27292732
const sycl::context &ctx = exec_q.get_context();
27302733

2734+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
27312735
cgh.host_task([ctx, partially_reduced_tmp] {
2732-
sycl::free(partially_reduced_tmp, ctx);
2736+
sycl_free_noexcept(partially_reduced_tmp, ctx);
27332737
});
27342738
});
27352739

@@ -3038,7 +3042,8 @@ gemm_batch_contig_tree_k_impl(sycl::queue &exec_q,
30383042
cgh.depends_on(red_ev);
30393043
const sycl::context &ctx = exec_q.get_context();
30403044

3041-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
3045+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
3046+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
30423047
});
30433048
return cleanup_host_task_event;
30443049
}
@@ -3097,8 +3102,9 @@ gemm_batch_contig_tree_k_impl(sycl::queue &exec_q,
30973102
cgh.depends_on(red_ev);
30983103
const sycl::context &ctx = exec_q.get_context();
30993104

3105+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
31003106
cgh.host_task([ctx, partially_reduced_tmp] {
3101-
sycl::free(partially_reduced_tmp, ctx);
3107+
sycl_free_noexcept(partially_reduced_tmp, ctx);
31023108
});
31033109
});
31043110

@@ -3238,7 +3244,8 @@ gemm_batch_contig_tree_nm_impl(sycl::queue &exec_q,
32383244
cgh.depends_on(red_ev);
32393245
const sycl::context &ctx = exec_q.get_context();
32403246

3241-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
3247+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
3248+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
32423249
});
32433250
return cleanup_host_task_event;
32443251
}
@@ -3299,8 +3306,9 @@ gemm_batch_contig_tree_nm_impl(sycl::queue &exec_q,
32993306
cgh.depends_on(red_ev);
33003307
const sycl::context &ctx = exec_q.get_context();
33013308

3309+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
33023310
cgh.host_task([ctx, partially_reduced_tmp] {
3303-
sycl::free(partially_reduced_tmp, ctx);
3311+
sycl_free_noexcept(partially_reduced_tmp, ctx);
33043312
});
33053313
});
33063314

@@ -3603,7 +3611,8 @@ sycl::event gemm_tree_k_impl(sycl::queue &exec_q,
36033611
cgh.depends_on(red_ev);
36043612
const sycl::context &ctx = exec_q.get_context();
36053613

3606-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
3614+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
3615+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
36073616
});
36083617
return cleanup_host_task_event;
36093618
}
@@ -3646,8 +3655,9 @@ sycl::event gemm_tree_k_impl(sycl::queue &exec_q,
36463655
cgh.depends_on(red_ev);
36473656
const sycl::context &ctx = exec_q.get_context();
36483657

3658+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
36493659
cgh.host_task([ctx, partially_reduced_tmp] {
3650-
sycl::free(partially_reduced_tmp, ctx);
3660+
sycl_free_noexcept(partially_reduced_tmp, ctx);
36513661
});
36523662
});
36533663

@@ -3769,7 +3779,8 @@ sycl::event gemm_tree_nm_impl(sycl::queue &exec_q,
37693779
cgh.depends_on(red_ev);
37703780
const sycl::context &ctx = exec_q.get_context();
37713781

3772-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
3782+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
3783+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
37733784
});
37743785
return cleanup_host_task_event;
37753786
}
@@ -3812,8 +3823,9 @@ sycl::event gemm_tree_nm_impl(sycl::queue &exec_q,
38123823
cgh.depends_on(red_ev);
38133824
const sycl::context &ctx = exec_q.get_context();
38143825

3826+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
38153827
cgh.host_task([ctx, partially_reduced_tmp] {
3816-
sycl::free(partially_reduced_tmp, ctx);
3828+
sycl_free_noexcept(partially_reduced_tmp, ctx);
38173829
});
38183830
});
38193831

@@ -4016,7 +4028,8 @@ sycl::event gemm_contig_tree_k_impl(sycl::queue &exec_q,
40164028
cgh.depends_on(red_ev);
40174029
const sycl::context &ctx = exec_q.get_context();
40184030

4019-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
4031+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
4032+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
40204033
});
40214034
return cleanup_host_task_event;
40224035
}
@@ -4058,8 +4071,9 @@ sycl::event gemm_contig_tree_k_impl(sycl::queue &exec_q,
40584071
cgh.depends_on(red_ev);
40594072
const sycl::context &ctx = exec_q.get_context();
40604073

4074+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
40614075
cgh.host_task([ctx, partially_reduced_tmp] {
4062-
sycl::free(partially_reduced_tmp, ctx);
4076+
sycl_free_noexcept(partially_reduced_tmp, ctx);
40634077
});
40644078
});
40654079

@@ -4170,7 +4184,8 @@ sycl::event gemm_contig_tree_nm_impl(sycl::queue &exec_q,
41704184
cgh.depends_on(red_ev);
41714185
const sycl::context &ctx = exec_q.get_context();
41724186

4173-
cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
4187+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
4188+
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
41744189
});
41754190
return cleanup_host_task_event;
41764191
}
@@ -4211,8 +4226,9 @@ sycl::event gemm_contig_tree_nm_impl(sycl::queue &exec_q,
42114226
cgh.depends_on(red_ev);
42124227
const sycl::context &ctx = exec_q.get_context();
42134228

4229+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
42144230
cgh.host_task([ctx, partially_reduced_tmp] {
4215-
sycl::free(partially_reduced_tmp, ctx);
4231+
sycl_free_noexcept(partially_reduced_tmp, ctx);
42164232
});
42174233
});
42184234

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

Lines changed: 15 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1374,8 +1374,9 @@ sycl::event reduction_over_group_temps_strided_impl(
13741374
cgh.depends_on(final_reduction_ev);
13751375
const sycl::context &ctx = exec_q.get_context();
13761376

1377+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
13771378
cgh.host_task([ctx, partially_reduced_tmp] {
1378-
sycl::free(partially_reduced_tmp, ctx);
1379+
sycl_free_noexcept(partially_reduced_tmp, ctx);
13791380
});
13801381
});
13811382

@@ -1617,8 +1618,9 @@ sycl::event reduction_axis1_over_group_temps_contig_impl(
16171618
cgh.depends_on(final_reduction_ev);
16181619
const sycl::context &ctx = exec_q.get_context();
16191620

1621+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
16201622
cgh.host_task([ctx, partially_reduced_tmp] {
1621-
sycl::free(partially_reduced_tmp, ctx);
1623+
sycl_free_noexcept(partially_reduced_tmp, ctx);
16221624
});
16231625
});
16241626

@@ -1861,8 +1863,9 @@ sycl::event reduction_axis0_over_group_temps_contig_impl(
18611863
cgh.depends_on(final_reduction_ev);
18621864
const sycl::context &ctx = exec_q.get_context();
18631865

1866+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
18641867
cgh.host_task([ctx, partially_reduced_tmp] {
1865-
sycl::free(partially_reduced_tmp, ctx);
1868+
sycl_free_noexcept(partially_reduced_tmp, ctx);
18661869
});
18671870
});
18681871

@@ -2796,10 +2799,11 @@ sycl::event search_over_group_temps_strided_impl(
27962799
cgh.depends_on(final_reduction_ev);
27972800
sycl::context ctx = exec_q.get_context();
27982801

2802+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
27992803
cgh.host_task(
28002804
[ctx, partially_reduced_tmp, partially_reduced_vals_tmp] {
2801-
sycl::free(partially_reduced_tmp, ctx);
2802-
sycl::free(partially_reduced_vals_tmp, ctx);
2805+
sycl_free_noexcept(partially_reduced_tmp, ctx);
2806+
sycl_free_noexcept(partially_reduced_vals_tmp, ctx);
28032807
});
28042808
});
28052809

@@ -3087,10 +3091,11 @@ sycl::event search_axis1_over_group_temps_contig_impl(
30873091
cgh.depends_on(final_reduction_ev);
30883092
sycl::context ctx = exec_q.get_context();
30893093

3094+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
30903095
cgh.host_task(
30913096
[ctx, partially_reduced_tmp, partially_reduced_vals_tmp] {
3092-
sycl::free(partially_reduced_tmp, ctx);
3093-
sycl::free(partially_reduced_vals_tmp, ctx);
3097+
sycl_free_noexcept(partially_reduced_tmp, ctx);
3098+
sycl_free_noexcept(partially_reduced_vals_tmp, ctx);
30943099
});
30953100
});
30963101

@@ -3374,10 +3379,11 @@ sycl::event search_axis0_over_group_temps_contig_impl(
33743379
cgh.depends_on(final_reduction_ev);
33753380
sycl::context ctx = exec_q.get_context();
33763381

3382+
using dpctl::tensor::offset_utils::sycl_free_noexcept;
33773383
cgh.host_task(
33783384
[ctx, partially_reduced_tmp, partially_reduced_vals_tmp] {
3379-
sycl::free(partially_reduced_tmp, ctx);
3380-
sycl::free(partially_reduced_vals_tmp, ctx);
3385+
sycl_free_noexcept(partially_reduced_tmp, ctx);
3386+
sycl_free_noexcept(partially_reduced_vals_tmp, ctx);
33813387
});
33823388
});
33833389

‎dpctl/tensor/libtensor/include/utils/offset_utils.hpp

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -106,6 +106,22 @@ class usm_host_allocator : public sycl::usm_allocator<T, sycl::usm::alloc::host>
106106
}
107107
};
108108

109+
template <typename T>
110+
void sycl_free_noexcept(T *ptr, const sycl::context &ctx) noexcept
111+
{
112+
try {
113+
sycl::free(ptr, ctx);
114+
} catch (const std::exception &e) {
115+
std::cerr << "Call to sycl::free caught exception: " << e.what()
116+
<< std::endl;
117+
}
118+
}
119+
120+
template <typename T> void sycl_free_noexcept(T *ptr, sycl::queue &q) noexcept
121+
{
122+
sycl_free_noexcept(ptr, q.get_context());
123+
}
124+
109125
template <typename indT, typename... Vs>
110126
std::tuple<indT *, size_t, sycl::event>
111127
device_allocate_and_pack(sycl::queue &q,

0 commit comments

Comments
 (0)