diff --git a/mindspore/ccsrc/device/gpu/gpu_kernel_runtime.cc b/mindspore/ccsrc/device/gpu/gpu_kernel_runtime.cc index 387977d345..dc4487ccee 100644 --- a/mindspore/ccsrc/device/gpu/gpu_kernel_runtime.cc +++ b/mindspore/ccsrc/device/gpu/gpu_kernel_runtime.cc @@ -199,7 +199,7 @@ bool GPUKernelRuntime::LaunchKernelDynamic(const session::KernelGraph *graph) { AddressPtrList kernel_workspaces; AddressPtrList kernel_outputs; AllocKernelDynamicRes(*kernel_mod, kernel, &kernel_inputs, &kernel_workspaces, &kernel_outputs); - if (!kernel_mod->Launch(kernel_inputs, kernel_workspaces, kernel_outputs, reinterpret_cast(stream_))) { + if (!kernel_mod->Launch(kernel_inputs, kernel_workspaces, kernel_outputs, stream_)) { MS_LOG(ERROR) << "Launch kernel failed."; return false; } diff --git a/mindspore/ccsrc/device/kernel_runtime.cc b/mindspore/ccsrc/device/kernel_runtime.cc index 42f56af8d5..0f5f282fd1 100644 --- a/mindspore/ccsrc/device/kernel_runtime.cc +++ b/mindspore/ccsrc/device/kernel_runtime.cc @@ -664,8 +664,7 @@ bool KernelRuntime::LaunchKernelMod(const session::KernelGraph &graph) { struct timeval start_time, end_time; (void)gettimeofday(&start_time, nullptr); #endif - auto ret = - kernel_mod->Launch(kernel_inputs, kernel_workspaces, kernel_outputs, reinterpret_cast(stream_)); + auto ret = kernel_mod->Launch(kernel_inputs, kernel_workspaces, kernel_outputs, stream_); if (!ret) { MS_LOG(ERROR) << "Launch kernel failed."; return false; diff --git a/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.cc b/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.cc index 5e07cd0ad7..7875baaf0e 100644 --- a/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.cc +++ b/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.cc @@ -103,14 +103,13 @@ void AicpuOpKernelMod::CreateCpuKernelInfo(const std::vector &inputs } bool AicpuOpKernelMod::Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) { - if (stream_ptr == 0) { + const std::vector &outputs, void *stream_ptr) { + if (stream_ptr == nullptr) { MS_LOG(ERROR) << "stream_ptr should not be nullptr."; return false; } CreateCpuKernelInfo(inputs, outputs); - auto *stream = reinterpret_cast(stream_ptr); if (node_name_ == "TopK") { node_name_ = "TopKV2"; } @@ -119,7 +118,7 @@ bool AicpuOpKernelMod::Launch(const std::vector &inputs, const std:: if (rtCpuKernelLaunch(reinterpret_cast(node_so_.c_str()), reinterpret_cast(node_name_.c_str()), 1, reinterpret_cast(args_.data()), static_cast(args_.length()), nullptr, - stream) != RT_ERROR_NONE) { + stream_ptr) != RT_ERROR_NONE) { MS_LOG(ERROR) << "Aicpu op launch failed!"; return false; diff --git a/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.h b/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.h index f7469c5e72..3ee9bd2a15 100644 --- a/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.h +++ b/mindspore/ccsrc/kernel/aicpu/aicpu_kernel_mod.h @@ -27,7 +27,7 @@ class AicpuOpKernelMod : public AscendKernelMod { AicpuOpKernelMod(); ~AicpuOpKernelMod() override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.cc b/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.cc index 7ce434b2d3..64590cd9b8 100644 --- a/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.cc +++ b/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.cc @@ -82,7 +82,7 @@ const std::vector &GpuKernelMod::GetOutputSizeList() const { return outp const std::vector &GpuKernelMod::GetWorkspaceSizeList() const { return workspace_size_list_; } bool GpuKernelMod::Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { if (stream_ptr == 0) { MS_LOG(ERROR) << "stream_ptr should not be nullptr."; return false; diff --git a/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.h b/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.h index 4ef903f1f3..df9cb069f7 100644 --- a/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.h +++ b/mindspore/ccsrc/kernel/akg/gpu/akg_gpu_kernel_mod.h @@ -64,7 +64,7 @@ class GpuKernelMod : public KernelMod { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; static GpuKernelManagerPtr kernelmanager_; diff --git a/mindspore/ccsrc/kernel/cpu/cpu_kernel.h b/mindspore/ccsrc/kernel/cpu/cpu_kernel.h index 378f36ac5b..edd133bf13 100644 --- a/mindspore/ccsrc/kernel/cpu/cpu_kernel.h +++ b/mindspore/ccsrc/kernel/cpu/cpu_kernel.h @@ -56,7 +56,7 @@ class CPUKernel : public kernel::KernelMod { void Init(const CNodePtr &kernel_node); virtual void InitKernel(const CNodePtr &kernel_node) = 0; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t /*stream_ptr*/) override { + const std::vector &outputs, void * /*stream_ptr*/) override { return Launch(inputs, workspace, outputs); }; virtual bool Launch(const std::vector &inputs, const std::vector &workspace, diff --git a/mindspore/ccsrc/kernel/gpu/arrays/argmax_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/argmax_gpu_kernel.h index ccd691ae91..3df70d0960 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/argmax_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/argmax_gpu_kernel.h @@ -35,7 +35,7 @@ class ArgmaxGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input = GetDeviceAddress(inputs, 0); int *output = GetDeviceAddress(outputs, 0); CalArgmax(input, SizeToInt(batch_size_), SizeToInt(channel_size_), axis_, output, diff --git a/mindspore/ccsrc/kernel/gpu/arrays/array_reduce_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/array_reduce_gpu_kernel.h index c8410c419d..224a3da8ad 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/array_reduce_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/array_reduce_gpu_kernel.h @@ -55,7 +55,7 @@ class ArrayReduceGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/arrays/concatv2_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/concatv2_gpu_kernel.h index 5dabb3045c..a91c50ce69 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/concatv2_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/concatv2_gpu_kernel.h @@ -34,7 +34,7 @@ class ConcatV2GpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (inputs.size() == 2) { T *input_0 = GetDeviceAddress(inputs, 0); T *input_1 = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/arrays/gather_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/gather_gpu_kernel.h index c4424df59c..72a05b0915 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/gather_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/gather_gpu_kernel.h @@ -35,7 +35,7 @@ class GatherGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); T *input_addr = GetDeviceAddress(inputs, 0); S *indices_addr = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/arrays/one_hot_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/one_hot_gpu_kernel.h index d8059869f2..c8b64e7243 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/one_hot_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/one_hot_gpu_kernel.h @@ -34,7 +34,7 @@ class OneHotGpuFwdKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); const S *indices = GetDeviceAddress(inputs, 0); const T *on_value = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/arrays/select_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/select_gpu_kernel.h index ba0bea4dee..f1b6c5853a 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/select_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/select_gpu_kernel.h @@ -34,7 +34,7 @@ class SelectGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { bool *input_cond = GetDeviceAddress(inputs, 0); T *input_x = GetDeviceAddress(inputs, 1); T *input_y = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/arrays/slice_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/slice_gpu_kernel.h index f71ec23d2e..eb829f73c6 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/slice_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/slice_gpu_kernel.h @@ -34,7 +34,7 @@ class SliceGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input = GetDeviceAddress(inputs, 0); T *output = GetDeviceAddress(outputs, 0); if (is_strided_slice_) { diff --git a/mindspore/ccsrc/kernel/gpu/arrays/slice_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/slice_grad_gpu_kernel.h index 80eef23112..bf24272d93 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/slice_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/slice_grad_gpu_kernel.h @@ -34,7 +34,7 @@ class SliceGradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *dy = GetDeviceAddress(inputs, 0); T *dx = GetDeviceAddress(outputs, 0); FillDeviceArray(outputs[0]->size / sizeof(T), dx, 0.f, reinterpret_cast(stream_ptr)); diff --git a/mindspore/ccsrc/kernel/gpu/arrays/transpose_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/transpose_gpu_kernel.h index 1c9cf925ea..61be9b68fe 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/transpose_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/transpose_gpu_kernel.h @@ -34,7 +34,7 @@ class TransposeGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input = GetDeviceAddress(inputs, 0); T *output = GetDeviceAddress(outputs, 0); int *input_shape = GetDeviceAddress(workspace, 0); diff --git a/mindspore/ccsrc/kernel/gpu/arrays/unsorted_segment_sum_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/arrays/unsorted_segment_sum_gpu_kernel.h index 24c1f09097..a20375ee29 100644 --- a/mindspore/ccsrc/kernel/gpu/arrays/unsorted_segment_sum_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/arrays/unsorted_segment_sum_gpu_kernel.h @@ -35,7 +35,7 @@ class UnsortedSegmentSumGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input_addr = GetDeviceAddress(inputs, 0); S *indices_addr = GetDeviceAddress(inputs, 1); T *output_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/control/recv_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/control/recv_gpu_kernel.h index 206eac5bd9..12b4eed132 100644 --- a/mindspore/ccsrc/kernel/gpu/control/recv_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/control/recv_gpu_kernel.h @@ -33,7 +33,7 @@ class RecvGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &, const std::vector &, const std::vector &, - uintptr_t) override { + void *) override { CHECK_CUDA_RET_WITH_EXCEPT(cudaStreamWaitEvent(wait_stream_, wait_event_, 0), "Waiting cuda event failed."); return true; } diff --git a/mindspore/ccsrc/kernel/gpu/control/send_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/control/send_gpu_kernel.h index 156ec4160d..a26e41aa1e 100644 --- a/mindspore/ccsrc/kernel/gpu/control/send_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/control/send_gpu_kernel.h @@ -33,7 +33,7 @@ class SendGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &, const std::vector &, const std::vector &, - uintptr_t) override { + void *) override { CHECK_CUDA_RET_WITH_EXCEPT(cudaEventRecord(record_event_, record_stream_), "Recording cuda event failed."); return true; } diff --git a/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.cc b/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.cc index d87a6cc726..777310cebc 100644 --- a/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.cc @@ -53,7 +53,7 @@ bool DatasetInitKernel::Init(const CNodePtr &kernel_node) { void DatasetInitKernel::InitSizeLists() { return; } bool DatasetInitKernel::Launch(const std::vector &, const std::vector &, - const std::vector &, uintptr_t) { + const std::vector &, void *) { void *addr = nullptr; size_t len = total_bytes_ * buffer_q_capacity_; diff --git a/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.h b/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.h index f7ffb41995..318049f4ad 100644 --- a/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/data/dataset_init_kernel.h @@ -33,7 +33,7 @@ class DatasetInitKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel_node) override; protected: diff --git a/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.cc b/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.cc index f8ee134018..69e35585cf 100644 --- a/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.cc @@ -64,7 +64,7 @@ bool DatasetIteratorKernel::Init(const CNodePtr &kernel_node) { void DatasetIteratorKernel::InitSizeLists() { return; } bool DatasetIteratorKernel::Launch(const std::vector &, const std::vector &, - const std::vector &outputs, uintptr_t) { + const std::vector &outputs, void *) { void *addr = nullptr; size_t len = 0; diff --git a/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.h b/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.h index d3231cab3c..cdd7a47e7b 100644 --- a/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/data/dataset_iterator_kernel.h @@ -33,7 +33,7 @@ class DatasetIteratorKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel_node) override; protected: diff --git a/mindspore/ccsrc/kernel/gpu/math/addn_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/addn_gpu_kernel.h index 485d3b4f72..0b27602761 100644 --- a/mindspore/ccsrc/kernel/gpu/math/addn_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/addn_gpu_kernel.h @@ -43,7 +43,7 @@ class AddNGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t) override { + const std::vector &outputs, void *) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/math/assign_add_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/assign_add_gpu_kernel.h index 191bd0b1c8..db69fd7be6 100644 --- a/mindspore/ccsrc/kernel/gpu/math/assign_add_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/assign_add_gpu_kernel.h @@ -35,7 +35,7 @@ class AssignAddGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/math/bias_add_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/bias_add_gpu_kernel.h index 03192a36a3..d7ed8e1072 100644 --- a/mindspore/ccsrc/kernel/gpu/math/bias_add_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/bias_add_gpu_kernel.h @@ -42,7 +42,7 @@ class BiasAddGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); VARIABLE_NOT_USED(stream_ptr); T *x_addr = GetDeviceAddress(inputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/math/broadcast_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/broadcast_gpu_kernel.h index 1b597cea70..c652d9aae4 100644 --- a/mindspore/ccsrc/kernel/gpu/math/broadcast_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/broadcast_gpu_kernel.h @@ -39,7 +39,7 @@ class BroadcastOpGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *lhs = GetDeviceAddress(inputs, 0); T *rhs = GetDeviceAddress(inputs, 1); S *output = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/math/broadcast_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/broadcast_grad_gpu_kernel.h index 2c15e68829..3e1f91b5b7 100644 --- a/mindspore/ccsrc/kernel/gpu/math/broadcast_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/broadcast_grad_gpu_kernel.h @@ -39,7 +39,7 @@ class BroadcastOpGradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *x1 = GetDeviceAddress(inputs, 0); T *x2 = GetDeviceAddress(inputs, 1); T *dy = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/math/equalcount_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/equalcount_gpu_kernel.h index 9e98d6be0f..7d3f74970f 100644 --- a/mindspore/ccsrc/kernel/gpu/math/equalcount_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/equalcount_gpu_kernel.h @@ -35,7 +35,7 @@ class EqualCountGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); T *input1 = GetDeviceAddress(inputs, 0); T *input2 = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/math/float_status_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/float_status_gpu_kernel.h index bdd93d5d54..1aa9b18684 100644 --- a/mindspore/ccsrc/kernel/gpu/math/float_status_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/float_status_gpu_kernel.h @@ -40,7 +40,7 @@ class FloatStatusGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input = GetDeviceAddress(inputs, 0); switch (kernel_name_) { diff --git a/mindspore/ccsrc/kernel/gpu/math/matmul_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/matmul_gpu_kernel.h index 2dc164b457..765fb3d7d4 100644 --- a/mindspore/ccsrc/kernel/gpu/math/matmul_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/matmul_gpu_kernel.h @@ -48,7 +48,7 @@ class MatMulGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); VARIABLE_NOT_USED(stream_ptr); auto input1_addr = GetDeviceAddress(inputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/math/tensoradd_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/tensoradd_gpu_kernel.h index 52480b8c70..67c6a34f3f 100644 --- a/mindspore/ccsrc/kernel/gpu/math/tensoradd_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/tensoradd_gpu_kernel.h @@ -43,7 +43,7 @@ class TensorAddGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t) { + const std::vector &outputs, void *) { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/math/unary_op_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/math/unary_op_gpu_kernel.h index d4c8a9b29c..4503b805f6 100644 --- a/mindspore/ccsrc/kernel/gpu/math/unary_op_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/math/unary_op_gpu_kernel.h @@ -62,7 +62,7 @@ class UnaryOpGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); T *input_addr = GetDeviceAddress(inputs, 0); T *output_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nccl/nccl_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nccl/nccl_gpu_kernel.h index 4ea332784d..b5ab46a67d 100644 --- a/mindspore/ccsrc/kernel/gpu/nccl/nccl_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nccl/nccl_gpu_kernel.h @@ -60,7 +60,7 @@ class NcclGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input_addr = GetDeviceAddress(inputs, 0); T *output_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nn/bias_add_grad_gpu_kenel.h b/mindspore/ccsrc/kernel/gpu/nn/bias_add_grad_gpu_kenel.h index 5c7153a172..c93a050649 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/bias_add_grad_gpu_kenel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/bias_add_grad_gpu_kenel.h @@ -45,7 +45,7 @@ class BiasAddGradGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *dy_addr = GetDeviceAddress(inputs, 0); T *db_addr = GetDeviceAddress(outputs, 0); T *indices_addr = GetDeviceAddress(workspace, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h index 6e218dddef..7bb6aa2a6d 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h @@ -60,7 +60,7 @@ class Conv2dGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h index faae0ae795..b126b542dd 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h @@ -61,7 +61,7 @@ class ConvGradFilterGpuBkwKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h index 325da74486..f7f371067f 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h @@ -61,7 +61,7 @@ class ConvGradInputGpuBkwKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.cc b/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.cc index 937f38137f..0d2a6be9c8 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.cc @@ -76,7 +76,7 @@ void DropoutGpuFwdKernel::InitSizeLists() { } bool DropoutGpuFwdKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.h index 2b0d84a40c..accff17429 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/dropout_gpu_kernel.h @@ -37,7 +37,7 @@ class DropoutGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel_node) override; diff --git a/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.cc b/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.cc index 42c3d279c4..44f603f02d 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.cc @@ -75,7 +75,7 @@ void DropoutGradGpuFwdKernel::InitSizeLists() { } bool DropoutGradGpuFwdKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.h index b59b5d2670..79d4117b58 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/dropout_grad_kernel.h @@ -32,7 +32,7 @@ class DropoutGradGpuFwdKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel_node) override; protected: diff --git a/mindspore/ccsrc/kernel/gpu/nn/flatten_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/flatten_gpu_kernel.h index 975dbd0082..3b0ad8c946 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/flatten_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/flatten_gpu_kernel.h @@ -35,7 +35,7 @@ class FlattenGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input = GetDeviceAddress(inputs, 0); T *output = GetDeviceAddress(outputs, 0); cudaError_t ret = diff --git a/mindspore/ccsrc/kernel/gpu/nn/flatten_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/flatten_grad_gpu_kernel.h index bdae6404dc..0748dc77db 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/flatten_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/flatten_grad_gpu_kernel.h @@ -35,7 +35,7 @@ class FlattenGardGpuBkwKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); T *input = GetDeviceAddress(inputs, 0); T *output = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nn/fused_batch_norm_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/fused_batch_norm_gpu_kernel.h index 3cdf480540..c08b341e78 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/fused_batch_norm_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/fused_batch_norm_gpu_kernel.h @@ -49,7 +49,7 @@ class FusedBatchNormGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); VARIABLE_NOT_USED(stream_ptr); if (is_null_input_) { diff --git a/mindspore/ccsrc/kernel/gpu/nn/fused_batchnorm_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/fused_batchnorm_grad_gpu_kernel.h index 07372ad22d..153b0286b3 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/fused_batchnorm_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/fused_batchnorm_grad_gpu_kernel.h @@ -47,7 +47,7 @@ class FusedBatchNormGradGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(workspace); VARIABLE_NOT_USED(stream_ptr); if (is_null_input_) { diff --git a/mindspore/ccsrc/kernel/gpu/nn/gelu_grad_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/gelu_grad_kernel.h index 7ce6d4d491..6415349012 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/gelu_grad_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/gelu_grad_kernel.h @@ -35,7 +35,7 @@ class GeLUGpuGradKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *dy_addr = GetDeviceAddress(inputs, 0); T *x_addr = GetDeviceAddress(inputs, 1); T *dx_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nn/gelu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/gelu_kernel.h index f0dd37dec4..60968d109b 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/gelu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/gelu_kernel.h @@ -35,7 +35,7 @@ class GeluGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *input_addr = GetDeviceAddress(inputs, 0); T *output_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nn/layer_norm_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/layer_norm_gpu_kernel.h index 29a609d95c..d5ec3ff8f2 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/layer_norm_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/layer_norm_gpu_kernel.h @@ -35,7 +35,7 @@ class LayerNormGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { auto x = GetDeviceAddress(inputs, 0); auto gamma = GetDeviceAddress(inputs, 1); auto beta = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/nn/layer_norm_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/layer_norm_grad_gpu_kernel.h index 72cfc23266..83bdedb9b3 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/layer_norm_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/layer_norm_grad_gpu_kernel.h @@ -35,7 +35,7 @@ class LayerNormGradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { auto x = GetDeviceAddress(inputs, 0); auto dy = GetDeviceAddress(inputs, 1); auto var = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/nn/lstm_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/lstm_gpu_kernel.h index c3e839b9c5..01247f0ed6 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/lstm_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/lstm_gpu_kernel.h @@ -59,7 +59,7 @@ class LstmGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(stream_ptr); auto x_addr = GetDeviceAddress(inputs, 0); auto hx_addr = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_data_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_data_gpu_kernel.h index b12fa3bea5..5591b0c817 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_data_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_data_gpu_kernel.h @@ -61,7 +61,7 @@ class LstmGradDataGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(stream_ptr); auto y_addr = GetDeviceAddress(inputs, 0); auto dy_addr = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_weight_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_weight_gpu_kernel.h index e081b9d070..dd6aae9a00 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_weight_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/lstm_grad_weight_gpu_kernel.h @@ -54,7 +54,7 @@ class LstmGradWeightGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { VARIABLE_NOT_USED(stream_ptr); auto x_addr = GetDeviceAddress(inputs, 0); auto hx_addr = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/nn/momentum_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/momentum_gpu_kernel.h index 2d8afb5780..8452c177db 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/momentum_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/momentum_gpu_kernel.h @@ -34,7 +34,7 @@ class MomentumGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, const std::vector &, - uintptr_t stream_ptr) override { + void *stream_ptr) override { T *variable = GetDeviceAddress(inputs, 0); T *accumulation = GetDeviceAddress(inputs, 1); T *learning_rate = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/nn/pooling_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/pooling_gpu_kernel.h index 2446c22950..faff453775 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/pooling_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/pooling_gpu_kernel.h @@ -59,7 +59,7 @@ class PoolingGpuFwdKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/pooling_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/pooling_grad_gpu_kernel.h index 535f96bbbf..df3454c581 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/pooling_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/pooling_grad_gpu_kernel.h @@ -61,7 +61,7 @@ class PoolingGradGpuFwdKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/relu_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/relu_gpu_kernel.h index d88efd3c7a..4cebc45831 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/relu_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/relu_gpu_kernel.h @@ -43,7 +43,7 @@ class ReLUGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t) override { + const std::vector &outputs, void *) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/relu_grad_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/relu_grad_kernel.h index e93dc31f80..ccc037f6e7 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/relu_grad_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/relu_grad_kernel.h @@ -41,7 +41,7 @@ class ReluGradGpuFwdKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t) override { + const std::vector &outputs, void *) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/rmsprop_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/rmsprop_gpu_kernel.h index d1ca53110b..7eaedfba52 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/rmsprop_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/rmsprop_gpu_kernel.h @@ -35,7 +35,7 @@ class RMSPropGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream) override { + const std::vector &outputs, void *stream) override { if (!use_center_) { T *variable = GetDeviceAddress(inputs, 0); T *mean_square = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/nn/softmax_cross_entropy_with_logits_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/softmax_cross_entropy_with_logits_gpu_kernel.h index 4d50d4753d..6840f0a1eb 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/softmax_cross_entropy_with_logits_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/softmax_cross_entropy_with_logits_gpu_kernel.h @@ -52,7 +52,7 @@ class SoftmaxCrossEntropyWithLogitsGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/softmax_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/softmax_gpu_kernel.h index cc381d7a71..060bc57d56 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/softmax_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/softmax_gpu_kernel.h @@ -52,7 +52,7 @@ class SoftmaxGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/softmax_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/softmax_grad_gpu_kernel.h index a0356c3bc4..003b55c0ed 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/softmax_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/softmax_grad_gpu_kernel.h @@ -51,7 +51,7 @@ class SoftmaxGradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/sparse_softmax_cross_entropy_with_logits_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/sparse_softmax_cross_entropy_with_logits_gpu_kernel.h index d232f7a131..0749172cc6 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/sparse_softmax_cross_entropy_with_logits_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/sparse_softmax_cross_entropy_with_logits_gpu_kernel.h @@ -52,7 +52,7 @@ class SparseSoftmaxCrossEntropyWithLogitsGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/nn/tanh_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/tanh_gpu_kernel.h index 29fb8cab48..7060ad1792 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/tanh_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/tanh_gpu_kernel.h @@ -37,7 +37,7 @@ class TanhGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { auto x_addr = GetDeviceAddress(inputs, 0); auto y_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/nn/tanh_grad_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/tanh_grad_kernel.h index 524dbe372b..b5b52d0acf 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/tanh_grad_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/tanh_grad_kernel.h @@ -37,7 +37,7 @@ class TanhGradKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { auto y_addr = GetDeviceAddress(inputs, 0); auto dy_addr = GetDeviceAddress(inputs, 1); auto dx_addr = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/other/assign_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/other/assign_gpu_kernel.h index 1c1cde4fd4..b41d583a43 100644 --- a/mindspore/ccsrc/kernel/gpu/other/assign_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/other/assign_gpu_kernel.h @@ -33,7 +33,7 @@ class AssignGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { T *var = GetDeviceAddress(inputs, 0); T *value = GetDeviceAddress(inputs, 1); T *output = GetDeviceAddress(outputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_gpu_kernel.h index c1804a5b93..5d2dee3ec7 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_gpu_kernel.h @@ -45,7 +45,7 @@ class BatchNormFold2GpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_grad_gpu_kernel.h index 38adda718c..28a4cf6cd6 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold2_grad_gpu_kernel.h @@ -45,7 +45,7 @@ class BatchNormFold2GradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { if (is_null_input_) { return true; } diff --git a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_gpu_kernel.h index a5a8a10dc0..a90e9b47d7 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_gpu_kernel.h @@ -53,7 +53,7 @@ class BatchNormFoldGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { (void)workspace; auto x = GetDeviceAddress(inputs, 0); auto mean = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_grad_gpu_kernel.h index cc420781da..8cbe5b6927 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/batchnorm_fold_grad_gpu_kernel.h @@ -47,7 +47,7 @@ class BatchNormFoldGradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { (void)workspace; // 'd_batch_mean', 'd_batch_std', 'x', 'batch_mean', 'batch_std', 'current_step' T *d_batch_mean = GetDeviceAddress(inputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/quant/correction_mul_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/correction_mul_gpu_kernel.h index eeab872ab3..38a9532ef5 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/correction_mul_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/correction_mul_gpu_kernel.h @@ -34,7 +34,7 @@ class CorrectionMulGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override { return output_size_list_; } const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { auto *weight = GetDeviceAddress(inputs, 0); auto *gamma = GetDeviceAddress(inputs, 1); auto *running_std = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/quant/correction_mul_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/correction_mul_grad_gpu_kernel.h index 29aeb3be13..3feffa586b 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/correction_mul_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/correction_mul_grad_gpu_kernel.h @@ -35,7 +35,7 @@ class CorrectionMulGradGpuKernel : public GpuKernel { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override { + const std::vector &outputs, void *stream_ptr) override { auto *d_out = GetDeviceAddress(inputs, 0); auto *weight = GetDeviceAddress(inputs, 1); auto *gamma = GetDeviceAddress(inputs, 2); diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.cc b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.cc index ee1cb0d012..ade7c32da0 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.cc @@ -114,7 +114,7 @@ void FakeQuantGpuKernel::InitSizeLists() { } bool FakeQuantGpuKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { float *output = GetDeviceAddress(outputs, 0); float *input = GetDeviceAddress(inputs, 0); float *input_min = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.h index b14268ed62..5a594c615f 100755 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_gpu_kernel.h @@ -32,7 +32,7 @@ class FakeQuantGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel) override; protected: diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.cc b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.cc index 239e55b5b0..7b7e3f1737 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.cc @@ -92,7 +92,7 @@ void FakeQuantGradGpuKernel::InitSizeLists() { } bool FakeQuantGradGpuKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { float *output = GetDeviceAddress(outputs, 0); float *gradient = GetDeviceAddress(inputs, 0); float *input = GetDeviceAddress(inputs, 1); diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.h index cd0f9a4680..04c505d2bd 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_grad_gpu_kernel.h @@ -32,7 +32,7 @@ class FakeQuantGradGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel_node) override; protected: diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.cc b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.cc index c452bb5dd1..083bf7f011 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.cc @@ -118,7 +118,7 @@ void FakeQuantPerChannelGpuKernel::InitSizeLists() { void FakeQuantPerChannelGpuKernel::CalFakeQuantizeForTraining(float *input, float *output, float *input_min, float *input_max, float *d_nudge_min, float *d_nudge_max, - float *d_scale, uintptr_t stream_ptr) { + float *d_scale, void *stream_ptr) { // calculate the input min and max according by the parameter ema and ema_decay. CalMinMaxPerChannel(input, input_min, input_max, input_size_ / sizeof(float), channel_out_, ema_decay_, ema_, reinterpret_cast(stream_ptr)); @@ -139,7 +139,7 @@ void FakeQuantPerChannelGpuKernel::CalFakeQuantizeForTraining(float *input, floa void FakeQuantPerChannelGpuKernel::CalFakeQuantizeForInfer(float *input, float *output, float *input_min, float *input_max, float *d_nudge_min, float *d_nudge_max, - float *d_scale, uintptr_t stream_ptr) { + float *d_scale, void *stream_ptr) { // real launch CalNudgePerChannel(input_min, input_max, quant_min_, quant_max_, d_nudge_min, d_nudge_max, d_scale, channel_out_, reinterpret_cast(stream_ptr)); @@ -149,7 +149,7 @@ void FakeQuantPerChannelGpuKernel::CalFakeQuantizeForInfer(float *input, float * bool FakeQuantPerChannelGpuKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { (void)workspace; float *output = GetDeviceAddress(outputs, 0); float *input = GetDeviceAddress(inputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.h index 8a1bb7293a..bea1a7421f 100755 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_gpu_kernel.h @@ -32,7 +32,7 @@ class FakeQuantPerChannelGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel) override; protected: @@ -40,9 +40,9 @@ class FakeQuantPerChannelGpuKernel : public GpuKernel { private: void CalFakeQuantizeForTraining(float *input, float *output, float *input_min, float *input_max, float *d_nudge_min, - float *d_nudge_max, float *d_scale, uintptr_t stream_ptr); + float *d_nudge_max, float *d_scale, void *stream_ptr); void CalFakeQuantizeForInfer(float *input, float *output, float *input_min, float *input_max, float *d_nudge_min, - float *d_nudge_max, float *d_scale, uintptr_t stream_ptr); + float *d_nudge_max, float *d_scale, void *stream_ptr); size_t input_size_; size_t min_size_; diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.cc b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.cc index f995f81190..88c976285c 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.cc +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.cc @@ -104,7 +104,7 @@ void FakeQuantPerChannelGradGpuKernel::InitSizeLists() { bool FakeQuantPerChannelGradGpuKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { (void)workspace; float *output = GetDeviceAddress(outputs, 0); float *gradient = GetDeviceAddress(inputs, 0); diff --git a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.h index c210f4cc81..fe760d85d2 100644 --- a/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/quant/fake_quant_per_channel_grad_gpu_kernel.h @@ -32,7 +32,7 @@ class FakeQuantPerChannelGradGpuKernel : public GpuKernel { const std::vector &GetOutputSizeList() const override; const std::vector &GetWorkspaceSizeList() const override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; bool Init(const CNodePtr &kernel_node) override; protected: diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.cc b/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.cc index 3cc57fe6d8..dba692606c 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.cc +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.cc @@ -24,17 +24,17 @@ namespace mindspore { namespace kernel { -bool HcomAllBroadCastKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool HcomAllBroadCastKernel::Launch(const std::vector &inputs, + const std::vector & /*workspace*/, + const std::vector & /*outputs*/, void *stream_ptr) { auto context_ptr = MsContext::GetInstance(); MS_EXCEPTION_IF_NULL(context_ptr); if (context_ptr->enable_task_sink()) { return true; } const char *tag = "Hccl-BroadCast"; - auto stream = reinterpret_cast(stream_ptr); hcclResult_t ret = - hcom_broadcast(tag, inputs[0]->addr, hccl_count_, hccl_data_type_list_[0], root_id_, nullptr, stream); + hcom_broadcast(tag, inputs[0]->addr, hccl_count_, hccl_data_type_list_[0], root_id_, nullptr, stream_ptr); if (ret != HCCL_SUCCESS) { MS_LOG(ERROR) << "HcomBroadcastOp : hcom_broadcast fail, return: " << static_cast(ret); return false; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.h b/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.h index d7d02a9451..ca8eba91af 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.h +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_broadcast.h @@ -31,7 +31,7 @@ class HcomAllBroadCastKernel : public HcclKernel { /* Inherit from kernelmod */ bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; private: }; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_gather.cc b/mindspore/ccsrc/kernel/hccl/hcom_all_gather.cc index fde1e3bb12..67cd1001e3 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_gather.cc +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_gather.cc @@ -24,17 +24,16 @@ namespace mindspore { namespace kernel { -bool HcomAllGatherKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool HcomAllGatherKernel::Launch(const std::vector &inputs, const std::vector & /*workspace*/, + const std::vector &outputs, void *stream_ptr) { auto context_ptr = MsContext::GetInstance(); MS_EXCEPTION_IF_NULL(context_ptr); if (context_ptr->enable_task_sink()) { return true; } const char *tag = "Hccl-AllGather"; - auto stream = reinterpret_cast(stream_ptr); hcclResult_t ret = - hcom_all_gather(tag, inputs[0]->addr, outputs[0]->addr, hccl_count_, hccl_data_type_list_[0], nullptr, stream); + hcom_all_gather(tag, inputs[0]->addr, outputs[0]->addr, hccl_count_, hccl_data_type_list_[0], nullptr, stream_ptr); if (ret != HCCL_SUCCESS) { MS_LOG(ERROR) << "HcomAllGatherKernelOp : hcom_all_gather fail, return: " << static_cast(ret); return false; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_gather.h b/mindspore/ccsrc/kernel/hccl/hcom_all_gather.h index f29b5cc0f6..5de2c513cf 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_gather.h +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_gather.h @@ -31,7 +31,7 @@ class HcomAllGatherKernel : public HcclKernel { /* Inherit from kernelmod */ bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; private: }; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.cc b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.cc index a0d96683c2..2bf9823e5d 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.cc +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.cc @@ -24,17 +24,16 @@ namespace mindspore { namespace kernel { -bool HcomAllReduceKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool HcomAllReduceKernel::Launch(const std::vector &inputs, const std::vector & /*workspace*/, + const std::vector &outputs, void *stream_ptr) { auto context_ptr = MsContext::GetInstance(); MS_EXCEPTION_IF_NULL(context_ptr); if (context_ptr->enable_task_sink()) { return true; } const char *tag = "Hccl-AllReduce"; - auto stream = reinterpret_cast(stream_ptr); hcclResult_t ret = hcom_all_reduce(tag, inputs[0]->addr, outputs[0]->addr, hccl_count_, hccl_data_type_list_[0], - op_type_, nullptr, stream); + op_type_, nullptr, stream_ptr); if (ret != HCCL_SUCCESS) { MS_LOG(ERROR) << "HcomAllReduceKernelOp : hcom_all_reduce fail, return: " << static_cast(ret); return false; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.h b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.h index 0a3bdb3284..939abd9de7 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.h +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce.h @@ -30,7 +30,7 @@ class HcomAllReduceKernel : public HcclKernel { /* Inherit from kernelmod */ bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; private: }; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.cc b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.cc index 36341ed6a7..05217108d9 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.cc +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.cc @@ -24,17 +24,17 @@ namespace mindspore { namespace kernel { -bool HcomAllReduceScatterKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool HcomAllReduceScatterKernel::Launch(const std::vector &inputs, + const std::vector & /*workspace*/, + const std::vector &outputs, void *stream_ptr) { auto context_ptr = MsContext::GetInstance(); MS_EXCEPTION_IF_NULL(context_ptr); if (context_ptr->enable_task_sink()) { return true; } const char *tag = "Hccl-ReduceScatter"; - auto stream = reinterpret_cast(stream_ptr); hcclResult_t ret = hcom_reduce_scatter(tag, inputs[0]->addr, outputs[0]->addr, hccl_count_, hccl_data_type_list_[0], - op_type_, nullptr, stream); + op_type_, nullptr, stream_ptr); if (ret != HCCL_SUCCESS) { MS_LOG(ERROR) << "HcomReduceScatterOp : hcom_reduce_scatter fail, return: " << static_cast(ret); return false; diff --git a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.h b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.h index 4c4f821d36..c734b517c6 100644 --- a/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.h +++ b/mindspore/ccsrc/kernel/hccl/hcom_all_reduce_scatter.h @@ -31,7 +31,7 @@ class HcomAllReduceScatterKernel : public HcclKernel { /* Inherit from kernelmod */ bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; private: }; diff --git a/mindspore/ccsrc/kernel/kernel.h b/mindspore/ccsrc/kernel/kernel.h index 4217b56625..684e6cf160 100644 --- a/mindspore/ccsrc/kernel/kernel.h +++ b/mindspore/ccsrc/kernel/kernel.h @@ -124,7 +124,7 @@ class KernelMod { virtual const std::vector &GetOutputSizeList() const = 0; virtual const std::vector &GetWorkspaceSizeList() const = 0; virtual bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) = 0; + const std::vector &outputs, void *stream_ptr) = 0; virtual std::vector GenParameters() { return {}; } virtual ~KernelMod() = default; diff --git a/mindspore/ccsrc/kernel/rts/assign.cc b/mindspore/ccsrc/kernel/rts/assign.cc index b7212c06f6..7f214b6e6f 100644 --- a/mindspore/ccsrc/kernel/rts/assign.cc +++ b/mindspore/ccsrc/kernel/rts/assign.cc @@ -30,10 +30,8 @@ AssignKernel::AssignKernel() {} AssignKernel::~AssignKernel() {} -bool AssignKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { - auto stream = reinterpret_cast(stream_ptr); - +bool AssignKernel::Launch(const std::vector &inputs, const std::vector & /*workspace*/, + const std::vector & /*outputs*/, void *stream_ptr) { if (inputs.size() != 2) { MS_LOG(ERROR) << "inputs size is not two"; return false; @@ -44,7 +42,7 @@ bool AssignKernel::Launch(const std::vector &inputs, const std::vect return true; } rtError_t status = rtMemcpyAsync(inputs[0]->addr, inputs[0]->size, inputs[1]->addr, inputs[1]->size, - RT_MEMCPY_DEVICE_TO_DEVICE, stream); + RT_MEMCPY_DEVICE_TO_DEVICE, stream_ptr); if (status != RT_ERROR_NONE) { MS_LOG(ERROR) << "Assign op rtMemcpyAsync failed!"; return false; diff --git a/mindspore/ccsrc/kernel/rts/assign.h b/mindspore/ccsrc/kernel/rts/assign.h index c4c6014e8a..0e7e52d48f 100644 --- a/mindspore/ccsrc/kernel/rts/assign.h +++ b/mindspore/ccsrc/kernel/rts/assign.h @@ -29,7 +29,7 @@ class AssignKernel : public RtKernel { ~AssignKernel() override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; }; diff --git a/mindspore/ccsrc/kernel/rts/label_goto.cc b/mindspore/ccsrc/kernel/rts/label_goto.cc index b60361bc60..7bcf42a210 100644 --- a/mindspore/ccsrc/kernel/rts/label_goto.cc +++ b/mindspore/ccsrc/kernel/rts/label_goto.cc @@ -45,8 +45,8 @@ bool LabelGotoKernel::Init(const AnfNodePtr &anf_node) { return true; } -bool LabelGotoKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool LabelGotoKernel::Launch(const std::vector & /*inputs*/, const std::vector & /*workspace*/, + const std::vector & /*outputs*/, void * /*stream_ptr*/) { MS_LOG(INFO) << "LabelGotoKernel launch"; return true; } diff --git a/mindspore/ccsrc/kernel/rts/label_goto.h b/mindspore/ccsrc/kernel/rts/label_goto.h index f4356391e6..efccc12d6f 100644 --- a/mindspore/ccsrc/kernel/rts/label_goto.h +++ b/mindspore/ccsrc/kernel/rts/label_goto.h @@ -32,7 +32,7 @@ class LabelGotoKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/label_set.cc b/mindspore/ccsrc/kernel/rts/label_set.cc index 0686cb0219..5aedd012dc 100644 --- a/mindspore/ccsrc/kernel/rts/label_set.cc +++ b/mindspore/ccsrc/kernel/rts/label_set.cc @@ -45,8 +45,8 @@ bool LabelSetKernel::Init(const AnfNodePtr &anf_node) { return true; } -bool LabelSetKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool LabelSetKernel::Launch(const std::vector & /*inputs*/, const std::vector & /*workspace*/, + const std::vector & /*outputs*/, void * /*stream_ptr*/) { MS_LOG(INFO) << "LabelSetKernel launch"; return true; } diff --git a/mindspore/ccsrc/kernel/rts/label_set.h b/mindspore/ccsrc/kernel/rts/label_set.h index 12f0d5b806..d05d81f898 100644 --- a/mindspore/ccsrc/kernel/rts/label_set.h +++ b/mindspore/ccsrc/kernel/rts/label_set.h @@ -32,7 +32,7 @@ class LabelSetKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/label_switch.cc b/mindspore/ccsrc/kernel/rts/label_switch.cc index f428996410..6647ac7eb6 100644 --- a/mindspore/ccsrc/kernel/rts/label_switch.cc +++ b/mindspore/ccsrc/kernel/rts/label_switch.cc @@ -53,8 +53,9 @@ bool LabelSwitchKernel::Init(const AnfNodePtr &anf_node) { return true; } -bool LabelSwitchKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { +bool LabelSwitchKernel::Launch(const std::vector & /*inputs*/, + const std::vector & /*workspace*/, + const std::vector & /*outputs*/, void * /*stream_ptr*/) { MS_LOG(INFO) << "LabelSwitchKernel launch"; return true; } diff --git a/mindspore/ccsrc/kernel/rts/label_switch.h b/mindspore/ccsrc/kernel/rts/label_switch.h index 372dc0cd7a..0accd26afb 100644 --- a/mindspore/ccsrc/kernel/rts/label_switch.h +++ b/mindspore/ccsrc/kernel/rts/label_switch.h @@ -32,7 +32,7 @@ class LabelSwitchKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/memcpy_async.cc b/mindspore/ccsrc/kernel/rts/memcpy_async.cc index 08b040cd15..3d5a7c88ab 100644 --- a/mindspore/ccsrc/kernel/rts/memcpy_async.cc +++ b/mindspore/ccsrc/kernel/rts/memcpy_async.cc @@ -34,9 +34,7 @@ MemCpyAsyncKernel::MemCpyAsyncKernel() {} MemCpyAsyncKernel::~MemCpyAsyncKernel() {} bool MemCpyAsyncKernel::Launch(const std::vector &inputs, const std::vector & /*workspace*/, - const std::vector &outputs, uintptr_t stream_ptr) { - auto stream = reinterpret_cast(stream_ptr); - + const std::vector &outputs, void *stream_ptr) { if (inputs.size() != 1) { MS_LOG(ERROR) << "inputs size is not one"; return false; @@ -51,7 +49,7 @@ bool MemCpyAsyncKernel::Launch(const std::vector &inputs, const std: return true; } rtError_t status = rtMemcpyAsync(outputs[0]->addr, outputs[0]->size, inputs[0]->addr, inputs[0]->size, - RT_MEMCPY_DEVICE_TO_DEVICE, stream); + RT_MEMCPY_DEVICE_TO_DEVICE, stream_ptr); if (status != RT_ERROR_NONE) { MS_LOG(ERROR) << "MemCpyAsync op rtMemcpyAsync failed!"; return false; diff --git a/mindspore/ccsrc/kernel/rts/memcpy_async.h b/mindspore/ccsrc/kernel/rts/memcpy_async.h index b8e39374f8..94bbf1ca1c 100644 --- a/mindspore/ccsrc/kernel/rts/memcpy_async.h +++ b/mindspore/ccsrc/kernel/rts/memcpy_async.h @@ -31,7 +31,7 @@ class MemCpyAsyncKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.cc b/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.cc index 002c252374..ff005f399b 100644 --- a/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.cc +++ b/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.cc @@ -50,10 +50,9 @@ bool ProfilingKernelMod::Init(const AnfNodePtr &anf_node) { return true; } -bool ProfilingKernelMod::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { - MS_LOG(INFO) << "gen task inputs size:" << inputs.size() << ", workspace size:" << workspace.size() - << ", outputs size:" << outputs.size() << ", stream_ptr:" << stream_ptr; +bool ProfilingKernelMod::Launch(const std::vector & /*inputs*/, + const std::vector & /*workspace*/, + const std::vector & /*outputs*/, void * /*stream_ptr*/) { return true; } diff --git a/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.h b/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.h index b9b4a1c05c..f77f3b5c67 100644 --- a/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.h +++ b/mindspore/ccsrc/kernel/rts/profiling_kernel_mod.h @@ -24,7 +24,7 @@ class ProfilingKernelMod : public RtKernel { ProfilingKernelMod() = default; ~ProfilingKernelMod() override = default; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; bool Init(const AnfNodePtr &anf_node) override; diff --git a/mindspore/ccsrc/kernel/rts/recv.cc b/mindspore/ccsrc/kernel/rts/recv.cc index af921083b5..b68380dac8 100644 --- a/mindspore/ccsrc/kernel/rts/recv.cc +++ b/mindspore/ccsrc/kernel/rts/recv.cc @@ -43,10 +43,9 @@ bool RecvKernel::Init(const AnfNodePtr &anf_node) { } bool RecvKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { rtEvent_t stream_event{}; - auto stream = reinterpret_cast(stream_ptr); - auto status = rtStreamWaitEvent(stream, stream_event); + auto status = rtStreamWaitEvent(stream_ptr, stream_event); if (status != RT_ERROR_NONE) { MS_LOG(ERROR) << "Recv rtStreamWaitEvent failed!"; return false; diff --git a/mindspore/ccsrc/kernel/rts/recv.h b/mindspore/ccsrc/kernel/rts/recv.h index 3a32fb9812..68f0b69cc5 100644 --- a/mindspore/ccsrc/kernel/rts/recv.h +++ b/mindspore/ccsrc/kernel/rts/recv.h @@ -31,7 +31,7 @@ class RecvKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/send.cc b/mindspore/ccsrc/kernel/rts/send.cc index abb93f5af4..ebcb53069e 100644 --- a/mindspore/ccsrc/kernel/rts/send.cc +++ b/mindspore/ccsrc/kernel/rts/send.cc @@ -40,10 +40,9 @@ bool SendKernel::Init(const AnfNodePtr &anf_node) { } bool SendKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { rtEvent_t event{}; - auto stream = reinterpret_cast(stream_ptr); - rtError_t status = rtEventRecord(event, stream); + rtError_t status = rtEventRecord(event, stream_ptr); if (status != RT_ERROR_NONE) { MS_LOG(ERROR) << "Send op rtEventRecord failed!"; return false; diff --git a/mindspore/ccsrc/kernel/rts/send.h b/mindspore/ccsrc/kernel/rts/send.h index bbf571b6ab..5c5b7cf09e 100644 --- a/mindspore/ccsrc/kernel/rts/send.h +++ b/mindspore/ccsrc/kernel/rts/send.h @@ -29,7 +29,7 @@ class SendKernel : public RtKernel { ~SendKernel() override; bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/stream_active.cc b/mindspore/ccsrc/kernel/rts/stream_active.cc index ce6d94e4a5..3666dd670f 100644 --- a/mindspore/ccsrc/kernel/rts/stream_active.cc +++ b/mindspore/ccsrc/kernel/rts/stream_active.cc @@ -41,9 +41,8 @@ bool StreamActiveKernel::Init(const AnfNodePtr &anf_node) { } bool StreamActiveKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { MS_LOG(INFO) << "Stream active op launch start"; - auto stream = reinterpret_cast(stream_ptr); if (active_streams_index_.empty()) { MS_LOG(ERROR) << "activeStreamList_ is empty!"; @@ -54,7 +53,7 @@ bool StreamActiveKernel::Launch(const std::vector &inputs, const std rtError_t status; for (auto index : active_streams_index_) { act_stream = kernel::TaskStream::GetInstance()->gen_stream_list()[index]; - status = rtStreamActive(act_stream, stream); + status = rtStreamActive(act_stream, stream_ptr); if (status != RT_ERROR_NONE) { MS_LOG(ERROR) << "Stream active failed!"; return false; diff --git a/mindspore/ccsrc/kernel/rts/stream_active.h b/mindspore/ccsrc/kernel/rts/stream_active.h index 0955de8cef..68c422e7c2 100644 --- a/mindspore/ccsrc/kernel/rts/stream_active.h +++ b/mindspore/ccsrc/kernel/rts/stream_active.h @@ -31,7 +31,7 @@ class StreamActiveKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/rts/stream_switch.cc b/mindspore/ccsrc/kernel/rts/stream_switch.cc index 4fc8f8ef59..9dfb3e8de0 100644 --- a/mindspore/ccsrc/kernel/rts/stream_switch.cc +++ b/mindspore/ccsrc/kernel/rts/stream_switch.cc @@ -51,7 +51,7 @@ bool StreamSwitchKernel::Init(const AnfNodePtr &anf_node) { } bool StreamSwitchKernel::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { + const std::vector &outputs, void *stream_ptr) { MS_LOG(INFO) << "stream switch op launch start"; if (inputs.size() != 2) { MS_LOG(ERROR) << "Stream switch inputs size is " << inputs.size() << ", only support 2"; @@ -59,9 +59,8 @@ bool StreamSwitchKernel::Launch(const std::vector &inputs, const std void *loop_cnt = inputs[0]->addr; void *ites_per_loop = inputs[1]->addr; - auto stream = reinterpret_cast(stream_ptr); rtStream_t true_stream_ = kernel::TaskStream::GetInstance()->gen_stream_list()[true_stream_index_]; - rtError_t status = rtStreamSwitchEx(loop_cnt, cond_, ites_per_loop, true_stream_, stream, data_type_); + rtError_t status = rtStreamSwitchEx(loop_cnt, cond_, ites_per_loop, true_stream_, stream_ptr, data_type_); if (status != RT_ERROR_NONE) { MS_LOG(ERROR) << "Stream switch failed!"; return false; diff --git a/mindspore/ccsrc/kernel/rts/stream_switch.h b/mindspore/ccsrc/kernel/rts/stream_switch.h index 3cc09b7494..4e927f3059 100644 --- a/mindspore/ccsrc/kernel/rts/stream_switch.h +++ b/mindspore/ccsrc/kernel/rts/stream_switch.h @@ -32,7 +32,7 @@ class StreamSwitchKernel : public RtKernel { bool Init(const AnfNodePtr &anf_node) override; bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, uint32_t stream_id) override; diff --git a/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.cc b/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.cc index 6a44ea2cc1..0f377940da 100644 --- a/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.cc +++ b/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.cc @@ -26,8 +26,8 @@ using TbeTaskInfoPtr = std::shared_ptr; using tbe::KernelManager; bool TbeKernelMod::Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) { - if (stream_ptr == 0) { + const std::vector &outputs, void *stream_ptr) { + if (stream_ptr == nullptr) { MS_LOG(ERROR) << "stream_ptr should not be nullptr."; return false; } @@ -55,10 +55,9 @@ bool TbeKernelMod::Launch(const std::vector &inpu [](const AddressPtr &addr) -> void * { return addr->addr; }); } rtL2Ctrl_t *l2ctrl = nullptr; - auto *stream = reinterpret_cast(stream_ptr); const void *stubFunc = reinterpret_cast(func_stub); auto argsSize = static_cast(UlongToUint(sizeof(void *)) * runtimeargs.size()); - if (RT_ERROR_NONE != rtKernelLaunch(stubFunc, blockdim, runtimeargs.data(), argsSize, l2ctrl, stream)) { + if (RT_ERROR_NONE != rtKernelLaunch(stubFunc, blockdim, runtimeargs.data(), argsSize, l2ctrl, stream_ptr)) { MS_LOG(ERROR) << "Call runtime rtKernelLaunch error."; return false; } diff --git a/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.h b/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.h index d7af83e2c9..e0e7ab4646 100644 --- a/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.h +++ b/mindspore/ccsrc/kernel/tbe/tbe_kernel_mod.h @@ -39,7 +39,7 @@ class TbeKernelMod : public AscendKernelMod { const std::vector &GetWorkspaceSizeList() const override { return workspace_size_list_; } bool Launch(const std::vector &inputs, const std::vector &workspace, - const std::vector &outputs, uintptr_t stream_ptr) override; + const std::vector &outputs, void *stream_ptr) override; std::vector GenTask(const std::vector &inputs, const std::vector &workspaces, const std::vector &outputs, uint32_t stream_id) override; std::vector GenParameters() override;