amd: fix implicit capture of this warnings

oneapi-src · Feb 6, 2025 · 6f9f97e · 6f9f97e
1 parent 7193163
commit 6f9f97e
Show file tree

Hide file tree

Showing 11 changed files with 373 additions and 312 deletions.
diff --git a/src/gpu/amd/miopen_binary.cpp b/src/gpu/amd/miopen_binary.cpp
@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2020-2024 Intel Corporation
+* Copyright 2020-2025 Intel Corporation
 * Copyright 2020 Codeplay Software Limited
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
@@ -15,6 +15,8 @@
 * limitations under the License.
 *******************************************************************************/
 
+#include "common/compiler_workarounds.hpp"
+
 #include "gpu/amd/miopen_binary.hpp"
 #include "gpu/amd/stream.hpp"
 #include "gpu/amd/sycl_hip_scoped_context.hpp"
@@ -40,20 +42,21 @@ status_t miopen_binary_t::execute(const exec_ctx_t &ctx) const {
         auto arg_scale1
                 = CTX_IN_SYCL_MEMORY(DNNL_ARG_ATTR_SCALES | DNNL_ARG_SRC_1);
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
 
-            void *a = arg_src_0.get_native_pointer(ih);
-            void *b = arg_src_1.get_native_pointer(ih);
-            void *c = arg_dst.get_native_pointer(ih);
-            void *s0 = arg_scale0.get_native_pointer(ih);
-            void *s1 = arg_scale1.get_native_pointer(ih);
+                    void *a = arg_src_0.get_native_pointer(ih);
+                    void *b = arg_src_1.get_native_pointer(ih);
+                    void *c = arg_dst.get_native_pointer(ih);
+                    void *s0 = arg_scale0.get_native_pointer(ih);
+                    void *s1 = arg_scale1.get_native_pointer(ih);
 
-            pd()->binary_impl_->execute(handle, a, b, c, s0, s1);
-        });
+                    pd()->binary_impl_->execute(handle, a, b, c, s0, s1);
+                });
     });
 }
 

diff --git a/src/gpu/amd/miopen_convolution.cpp b/src/gpu/amd/miopen_convolution.cpp
@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2020-2024 Intel Corporation
+* Copyright 2020-2025 Intel Corporation
 * Copyright 2020 Codeplay Software Limited
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
@@ -15,6 +15,8 @@
 * limitations under the License.
 *******************************************************************************/
 
+#include "common/compiler_workarounds.hpp"
+
 #include "gpu/amd/miopen_convolution.hpp"
 #include "gpu/amd/stream.hpp"
 #include "gpu/amd/sycl_hip_scoped_context.hpp"
@@ -55,25 +57,26 @@ status_t miopen_convolution_fwd_t::execute_convolution(
                     ::sycl::access::mode::read_write>(temp_reorder_mem, cgh);
         }
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
-
-            std::vector<void *> args;
-            args.push_back(arg_src.get_native_pointer(ih));
-            args.push_back(arg_weights.get_native_pointer(ih));
-            args.push_back(arg_dst.get_native_pointer(ih));
-            args.push_back(arg_bias.get_native_pointer(ih));
-            args.push_back(arg_scratch.get_native_pointer(ih));
-            args.push_back(arg_filter_scratch.get_native_pointer(ih));
-            args.push_back(temp_dst.get_native_pointer(ih));
-            args.push_back(temp_reorder.get_native_pointer(ih));
-            args.push_back(arg_oscale.get_native_pointer(ih));
-
-            pd()->impl_->execute(handle, args);
-        });
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
+
+                    std::vector<void *> args;
+                    args.push_back(arg_src.get_native_pointer(ih));
+                    args.push_back(arg_weights.get_native_pointer(ih));
+                    args.push_back(arg_dst.get_native_pointer(ih));
+                    args.push_back(arg_bias.get_native_pointer(ih));
+                    args.push_back(arg_scratch.get_native_pointer(ih));
+                    args.push_back(arg_filter_scratch.get_native_pointer(ih));
+                    args.push_back(temp_dst.get_native_pointer(ih));
+                    args.push_back(temp_reorder.get_native_pointer(ih));
+                    args.push_back(arg_oscale.get_native_pointer(ih));
+
+                    pd()->impl_->execute(handle, args);
+                });
     });
 }
 
@@ -91,22 +94,23 @@ status_t miopen_convolution_bwd_data_t::execute_convolution(
         auto arg_filter_scratch = CTX_SCRATCH_SYCL_MEMORY(
                 memory_tracking::names::key_conv_miopen_filter);
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
-
-            std::vector<void *> args;
-            args.push_back(arg_diff_src.get_native_pointer(ih));
-            args.push_back(arg_weights.get_native_pointer(ih));
-            args.push_back(arg_diff_dst.get_native_pointer(ih));
-            args.push_back(arg_bias.get_native_pointer(ih));
-            args.push_back(arg_scratch.get_native_pointer(ih));
-            args.push_back(arg_filter_scratch.get_native_pointer(ih));
-
-            pd()->impl_->execute(handle, args);
-        });
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
+
+                    std::vector<void *> args;
+                    args.push_back(arg_diff_src.get_native_pointer(ih));
+                    args.push_back(arg_weights.get_native_pointer(ih));
+                    args.push_back(arg_diff_dst.get_native_pointer(ih));
+                    args.push_back(arg_bias.get_native_pointer(ih));
+                    args.push_back(arg_scratch.get_native_pointer(ih));
+                    args.push_back(arg_filter_scratch.get_native_pointer(ih));
+
+                    pd()->impl_->execute(handle, args);
+                });
     });
 }
 
@@ -118,17 +122,19 @@ status_t miopen_convolution_bwd_weights_t::execute_zero_dims(
         auto arg_diff_weights = CTX_OUT_SYCL_MEMORY(DNNL_ARG_DIFF_WEIGHTS);
         auto arg_diff_bias = CTX_OUT_SYCL_MEMORY(DNNL_ARG_DIFF_BIAS);
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
 
-            void *weights = arg_diff_weights.get_native_pointer(ih);
-            void *bias = arg_diff_bias.get_native_pointer(ih);
+                    void *weights = arg_diff_weights.get_native_pointer(ih);
+                    void *bias = arg_diff_bias.get_native_pointer(ih);
 
-            pd()->impl_->execute_set_weights_bias(handle, weights, bias, 0.f);
-        });
+                    pd()->impl_->execute_set_weights_bias(
+                            handle, weights, bias, 0.f);
+                });
     });
 }
 
@@ -152,22 +158,23 @@ status_t miopen_convolution_bwd_weights_t::execute_convolution(
             arg_diff_bias = CTX_OUT_SYCL_MEMORY(DNNL_ARG_DIFF_BIAS);
         }
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
-
-            std::vector<void *> args;
-            args.push_back(arg_src.get_native_pointer(ih));
-            args.push_back(arg_diff_weights.get_native_pointer(ih));
-            args.push_back(arg_diff_dst.get_native_pointer(ih));
-            args.push_back(arg_diff_bias.get_native_pointer(ih));
-            args.push_back(arg_scratch.get_native_pointer(ih));
-            args.push_back(arg_filter_scratch.get_native_pointer(ih));
-
-            pd()->impl_->execute(handle, args);
-        });
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
+
+                    std::vector<void *> args;
+                    args.push_back(arg_src.get_native_pointer(ih));
+                    args.push_back(arg_diff_weights.get_native_pointer(ih));
+                    args.push_back(arg_diff_dst.get_native_pointer(ih));
+                    args.push_back(arg_diff_bias.get_native_pointer(ih));
+                    args.push_back(arg_scratch.get_native_pointer(ih));
+                    args.push_back(arg_filter_scratch.get_native_pointer(ih));
+
+                    pd()->impl_->execute(handle, args);
+                });
     });
 }
 

diff --git a/src/gpu/amd/miopen_deconvolution.cpp b/src/gpu/amd/miopen_deconvolution.cpp
@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2020-2024 Intel Corporation
+* Copyright 2020-2025 Intel Corporation
 * Copyright 2020 Codeplay Software Limited
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
@@ -15,6 +15,8 @@
 * limitations under the License.
 *******************************************************************************/
 
+#include "common/compiler_workarounds.hpp"
+
 #include "gpu/amd/miopen_deconvolution.hpp"
 #include "gpu/amd/stream.hpp"
 #include "gpu/amd/sycl_hip_scoped_context.hpp"
@@ -37,17 +39,18 @@ status_t miopen_deconvolution_bwd_weights_t::execute_bias(
         auto arg_diff_bias = CTX_OUT_SYCL_MEMORY(DNNL_ARG_DIFF_BIAS);
         auto arg_diff_dst = CTX_IN_SYCL_MEMORY(DNNL_ARG_DIFF_DST);
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
 
-            void *bias = arg_diff_bias.get_native_pointer(ih);
-            void *y = arg_diff_dst.get_native_pointer(ih);
+                    void *bias = arg_diff_bias.get_native_pointer(ih);
+                    void *y = arg_diff_dst.get_native_pointer(ih);
 
-            impl_->execute_bias(handle, y, bias);
-        });
+                    impl_->execute_bias(handle, y, bias);
+                });
     });
 }
 

diff --git a/src/gpu/amd/miopen_eltwise.cpp b/src/gpu/amd/miopen_eltwise.cpp
@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2020-2024 Intel Corporation
+* Copyright 2020-2025 Intel Corporation
 * Copyright 2020-2022 Codeplay Software Limited
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
@@ -15,6 +15,8 @@
 * limitations under the License.
 *******************************************************************************/
 
+#include "common/compiler_workarounds.hpp"
+
 #include "gpu/amd/miopen_eltwise.hpp"
 #include "gpu/amd/stream.hpp"
 #include "gpu/amd/sycl_hip_scoped_context.hpp"
@@ -36,18 +38,20 @@ status_t miopen_eltwise_fwd_t::execute(const exec_ctx_t &ctx) const {
         auto arg_src = CTX_IN_SYCL_MEMORY(DNNL_ARG_SRC);
         auto arg_dst = CTX_OUT_SYCL_MEMORY(DNNL_ARG_DST);
 
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            std::vector<void *> args;
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    std::vector<void *> args;
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
 
-            args.push_back(arg_src.get_native_pointer(ih));
-            args.push_back(arg_dst.get_native_pointer(ih));
+                    args.push_back(arg_src.get_native_pointer(ih));
+                    args.push_back(arg_dst.get_native_pointer(ih));
 
-            pd()->eltwise_fwd_impl_->execute(handle, args.data(), args.size());
-        });
+                    pd()->eltwise_fwd_impl_->execute(
+                            handle, args.data(), args.size());
+                });
     });
 }
 
@@ -61,19 +65,21 @@ status_t miopen_eltwise_bwd_t::execute(const exec_ctx_t &ctx) const {
         auto arg_src = CTX_IN_SYCL_MEMORY(DNNL_ARG_SRC);
         auto arg_diff_dst = CTX_IN_SYCL_MEMORY(DNNL_ARG_DIFF_DST);
         auto arg_diff_src = CTX_OUT_SYCL_MEMORY(DNNL_ARG_DIFF_SRC);
-        compat::host_task(cgh, [=](const compat::interop_handle &ih) {
-            std::vector<void *> args;
-            auto &sycl_engine
-                    = *utils::downcast<amd::engine_t *>(hip_stream->engine());
-            auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
-            auto handle = hip_stream->get_miopen_handle();
+        compat::host_task(cgh,
+                [= WA_THIS_COPY_CAPTURE](const compat::interop_handle &ih) {
+                    std::vector<void *> args;
+                    auto &sycl_engine = *utils::downcast<amd::engine_t *>(
+                            hip_stream->engine());
+                    auto sc = hip_sycl_scoped_context_handler_t(sycl_engine);
+                    auto handle = hip_stream->get_miopen_handle();
 
-            args.push_back(arg_src.get_native_pointer(ih));
-            args.push_back(arg_diff_dst.get_native_pointer(ih));
-            args.push_back(arg_diff_src.get_native_pointer(ih));
+                    args.push_back(arg_src.get_native_pointer(ih));
+                    args.push_back(arg_diff_dst.get_native_pointer(ih));
+                    args.push_back(arg_diff_src.get_native_pointer(ih));
 
-            pd()->eltwise_bwd_impl_->execute(handle, args.data(), args.size());
-        });
+                    pd()->eltwise_bwd_impl_->execute(
+                            handle, args.data(), args.size());
+                });
     });
 }