Skip to content

Commit

Permalink
Merge pull request #1797 from IntelPython/sycl-free-noexcept
Browse files Browse the repository at this point in the history
  • Loading branch information
@oleksandr-pavlyk
oleksandr-pavlyk authored Aug 14, 2024
2 parents 52edb6d + 3dd4469 commit 76b2eb0
Show file tree
Hide file tree
Showing 24 changed files with 316 additions and 139 deletions.
15 changes: 11 additions & 4 deletions dpctl/tensor/libtensor/include/kernels/accumulators.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@

#include "dpctl_tensor_types.hpp"
#include "utils/offset_utils.hpp"
#include "utils/sycl_alloc_utils.hpp"
#include "utils/sycl_utils.hpp"
#include "utils/type_dispatch_building.hpp"
#include "utils/type_utils.hpp"
Expand Down Expand Up @@ -436,7 +437,8 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
sycl::event free_ev = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(dependent_event);
const auto &ctx = exec_q.get_context();
cgh.host_task([ctx, temp]() { sycl::free(temp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, temp]() { sycl_free_noexcept(temp, ctx); });
});
host_tasks.push_back(free_ev);
}
Expand Down Expand Up @@ -765,7 +767,8 @@ sycl::event inclusive_scan_iter(sycl::queue &exec_q,
sycl::event free_ev = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(dependent_event);
const auto &ctx = exec_q.get_context();
cgh.host_task([ctx, temp]() { sycl::free(temp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, temp]() { sycl_free_noexcept(temp, ctx); });
});
host_tasks.push_back(free_ev);
}
Expand Down Expand Up @@ -917,7 +920,9 @@ size_t cumsum_val_contig_impl(sycl::queue &q,
});
copy_e.wait();
size_t return_val = static_cast<size_t>(*last_elem_host_usm);
sycl::free(last_elem_host_usm, q);

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
sycl_free_noexcept(last_elem_host_usm, q);

return return_val;
}
Expand Down Expand Up @@ -1026,7 +1031,9 @@ size_t cumsum_val_strided_impl(sycl::queue &q,
});
copy_e.wait();
size_t return_val = static_cast<size_t>(*last_elem_host_usm);
sycl::free(last_elem_host_usm, q);

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
sycl_free_noexcept(last_elem_host_usm, q);

return return_val;
}
Expand Down
9 changes: 7 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/common.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@
#include "kernels/alignment.hpp"
#include "kernels/dpctl_tensor_types.hpp"
#include "utils/offset_utils.hpp"
#include "utils/sycl_alloc_utils.hpp"

namespace dpctl
{
Expand Down Expand Up @@ -947,7 +948,9 @@ sycl::event binary_contig_matrix_contig_row_broadcast_impl(
sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(comp_ev);
const sycl::context &ctx = exec_q.get_context();
cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task(
[ctx, padded_vec]() { sycl_free_noexcept(padded_vec, ctx); });
});
host_tasks.push_back(tmp_cleanup_ev);

Expand Down Expand Up @@ -1026,7 +1029,9 @@ sycl::event binary_contig_row_contig_matrix_broadcast_impl(
sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(comp_ev);
const sycl::context &ctx = exec_q.get_context();
cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task(
[ctx, padded_vec]() { sycl_free_noexcept(padded_vec, ctx); });
});
host_tasks.push_back(tmp_cleanup_ev);

Expand Down
6 changes: 5 additions & 1 deletion dpctl/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,8 @@

#include "kernels/alignment.hpp"
#include "kernels/dpctl_tensor_types.hpp"
#include "utils/offset_utils.hpp"
#include "utils/sycl_alloc_utils.hpp"

namespace dpctl
{
Expand Down Expand Up @@ -458,7 +460,9 @@ sycl::event binary_inplace_row_matrix_broadcast_impl(
sycl::event tmp_cleanup_ev = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(comp_ev);
const sycl::context &ctx = exec_q.get_context();
cgh.host_task([ctx, padded_vec]() { sycl::free(padded_vec, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task(
[ctx, padded_vec]() { sycl_free_noexcept(padded_vec, ctx); });
});
host_tasks.push_back(tmp_cleanup_ev);

Expand Down
7 changes: 5 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/linalg_functions/dot_product.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,7 @@
#include "kernels/dpctl_tensor_types.hpp"
#include "kernels/reductions.hpp"
#include "utils/offset_utils.hpp"
#include "utils/sycl_alloc_utils.hpp"
#include "utils/sycl_utils.hpp"
#include "utils/type_utils.hpp"

Expand Down Expand Up @@ -1153,8 +1154,9 @@ sycl::event dot_product_tree_impl(sycl::queue &exec_q,
cgh.depends_on(final_reduction_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -1403,8 +1405,9 @@ dot_product_contig_tree_impl(sycl::queue &exec_q,
cgh.depends_on(final_reduction_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down
49 changes: 33 additions & 16 deletions dpctl/tensor/libtensor/include/kernels/linalg_functions/gemm.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,7 @@
#include "kernels/dpctl_tensor_types.hpp"
#include "kernels/reductions.hpp"
#include "utils/offset_utils.hpp"
#include "utils/sycl_alloc_utils.hpp"
#include "utils/sycl_utils.hpp"
#include "utils/type_utils.hpp"

Expand Down Expand Up @@ -2364,7 +2365,8 @@ gemm_batch_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -2427,8 +2429,9 @@ gemm_batch_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -2661,7 +2664,8 @@ gemm_batch_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -2728,8 +2732,9 @@ gemm_batch_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -3038,7 +3043,8 @@ gemm_batch_contig_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -3097,8 +3103,9 @@ gemm_batch_contig_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -3238,7 +3245,8 @@ gemm_batch_contig_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -3299,8 +3307,9 @@ gemm_batch_contig_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -3603,7 +3612,8 @@ sycl::event gemm_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -3646,8 +3656,9 @@ sycl::event gemm_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -3769,7 +3780,8 @@ sycl::event gemm_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -3812,8 +3824,9 @@ sycl::event gemm_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -4016,7 +4029,8 @@ sycl::event gemm_contig_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -4058,8 +4072,9 @@ sycl::event gemm_contig_tree_k_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -4170,7 +4185,8 @@ sycl::event gemm_contig_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

cgh.host_task([ctx, tmp] { sycl::free(tmp, ctx); });
using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, tmp] { sycl_free_noexcept(tmp, ctx); });
});
return cleanup_host_task_event;
}
Expand Down Expand Up @@ -4211,8 +4227,9 @@ sycl::event gemm_contig_tree_nm_impl(sycl::queue &exec_q,
cgh.depends_on(red_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down
25 changes: 16 additions & 9 deletions dpctl/tensor/libtensor/include/kernels/reductions.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@
#include "dpctl_tensor_types.hpp"
#include "utils/math_utils.hpp"
#include "utils/offset_utils.hpp"
#include "utils/sycl_alloc_utils.hpp"
#include "utils/sycl_utils.hpp"
#include "utils/type_dispatch_building.hpp"
#include "utils/type_utils.hpp"
Expand Down Expand Up @@ -1374,8 +1375,9 @@ sycl::event reduction_over_group_temps_strided_impl(
cgh.depends_on(final_reduction_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -1617,8 +1619,9 @@ sycl::event reduction_axis1_over_group_temps_contig_impl(
cgh.depends_on(final_reduction_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -1861,8 +1864,9 @@ sycl::event reduction_axis0_over_group_temps_contig_impl(
cgh.depends_on(final_reduction_ev);
const sycl::context &ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task([ctx, partially_reduced_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
});
});

Expand Down Expand Up @@ -2796,10 +2800,11 @@ sycl::event search_over_group_temps_strided_impl(
cgh.depends_on(final_reduction_ev);
sycl::context ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task(
[ctx, partially_reduced_tmp, partially_reduced_vals_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl::free(partially_reduced_vals_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_vals_tmp, ctx);
});
});

Expand Down Expand Up @@ -3087,10 +3092,11 @@ sycl::event search_axis1_over_group_temps_contig_impl(
cgh.depends_on(final_reduction_ev);
sycl::context ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task(
[ctx, partially_reduced_tmp, partially_reduced_vals_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl::free(partially_reduced_vals_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_vals_tmp, ctx);
});
});

Expand Down Expand Up @@ -3374,10 +3380,11 @@ sycl::event search_axis0_over_group_temps_contig_impl(
cgh.depends_on(final_reduction_ev);
sycl::context ctx = exec_q.get_context();

using dpctl::tensor::alloc_utils::sycl_free_noexcept;
cgh.host_task(
[ctx, partially_reduced_tmp, partially_reduced_vals_tmp] {
sycl::free(partially_reduced_tmp, ctx);
sycl::free(partially_reduced_vals_tmp, ctx);
sycl_free_noexcept(partially_reduced_tmp, ctx);
sycl_free_noexcept(partially_reduced_vals_tmp, ctx);
});
});

Expand Down
Loading

0 comments on commit 76b2eb0

Please sign in to comment.