You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mxnet.apache.org by GitBox <gi...@apache.org> on 2021/06/04 08:34:12 UTC
[GitHub] [incubator-mxnet] Kh4L opened a new pull request #20331: [WIP] Add async GPU depency Engine
Kh4L opened a new pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331
## Description ##
Changes how the GPU operations are synced in the `ThreadedEngine`.
GPU operators needed to be completed so the output variables (either `NDArray` or Resource) could be tagged as readable. The engine workers needed to `cudaStreamSynchronize` on the GPU stream to know when the variable is ready to read. This prevented an optimal GPU kernels overlapping with CPU operations.
This PR introduces a sync mechanism which leverages `cudaEvent` to avoid host synchronization between GPU operations.
The GPU writing operators tag the variables with a `cudaEvent` based sync object, and the operators reading the variables sync or wait on these objects.
### Details ###
// TODO: add details
Co-author: @ptrendx
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
szha commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r664135001
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
Review comment:
We already have `MXNET_ENGINE_TYPE` which decides whether `NaiveEngine` (i.e. synchronous scheduler) is used, so unless there's a strong reason I'd prefer not to introduce new environment variables here.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
szha commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-947884668
@Kh4L glad that the CI passed. @ptrendx @barry-jin could you take a look and see if there's any other concern to be addressed?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-947283143
Jenkins CI successfully triggered : [windows-gpu, centos-cpu, unix-cpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r658482251
##########
File path: include/mxnet/base.h
##########
@@ -348,10 +348,10 @@ struct RunContext {
* \brief the auxiliary stream of the device, can be nullptr or Stream<gpu>* in GPU mode
*/
void *aux_stream;
- /*!
- * \brief indicator of whether this execution is run in bulk mode
+ /*!
+ * \brief pointer to the cuda event pool used by the dependecy engine
Review comment:
"dependency"
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-945325359
@mxnet-bot run ci [unix-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha commented on a change in pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
szha commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r737056772
##########
File path: src/engine/threaded_engine.cc
##########
@@ -533,5 +546,206 @@ void ThreadedEngine::OnCompleteStatic(Engine* engine, void* opr_block_, const dm
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine* engine, void* opr_block, const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+static inline bool IsEngineAsync() {
+ std::string type = dmlc::GetEnv("MXNET_ENGINE_TYPE", std::string(""));
+ std::string async_engine_tag("Async");
+ auto tag_pos = type.find(async_engine_tag);
+ return tag_pos != std::string::npos;
+}
+
+void ThreadedEngine::OnStartCPU(Engine* engine, void* opr_block, const dmlc::Error* error) {
+ static bool use_new_dep_engine = IsEngineAsync();
+ if (!use_new_dep_engine) {
+ return;
+ }
Review comment:
Why not test it in the CI now? @Kh4L
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
szha commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r664135001
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
Review comment:
We already have `MXNET_ENGINE_TYPE` which decides whether `NaiveEngine` (i.e. synchronous scheduler) is used, so unless there's a strong reason I'd prefer not to introduce new environment variables here.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-868213909
I will review this PR next week.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sandeep-krishnamurthy commented on pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
sandeep-krishnamurthy commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-875125514
@apeforest @eric-haibin-lin can you please help review this PR?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r701324371
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaEventSynchronize(*ev));
+ }
+}
+
+void ThreadedEngine::OnStartGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& writer : sync_obj.writer_event) {
+ if (writer.event.expired()) {
+ sync_obj.writer_event.clear();
+ break;
+ }
+ if (writer.stream != worker_stream->stream_) {
+ // if there is already a reader on the same stream as us,
+ // it already synced with that writer and we can rely on
+ // the ongoing sync
+ bool found = false;
+ for (const auto& reader : reader_events) {
+ if (reader.stream == worker_stream->stream_) {
+ found = true;
+ break;
+ }
+ }
+ if (!found) {
+ AddEventHelper(&event_per_stream,
+ writer);
+ }
+ }
+ }
+ }
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // check for expired events and delete them
+ auto &reader_events = sync_obj.reader_events;
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ // if there are some readers, we wait for them
+ for (auto& cuda_event : reader_events) {
+ if (worker_stream->stream_ != cuda_event.stream) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
Review comment:
There can only be 1 writer event active (vs multiple readers).
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-860489315
@mxnet-bot run ci [centos-cpu, centos-gpu, edge, unix-cpu, unix-gpu, windows-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
szha commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-865476458
I plan to take a closer look this week. could @ptrendx also take a look?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r702268292
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaEventSynchronize(*ev));
+ }
+}
+
+void ThreadedEngine::OnStartGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& writer : sync_obj.writer_event) {
+ if (writer.event.expired()) {
+ sync_obj.writer_event.clear();
+ break;
+ }
+ if (writer.stream != worker_stream->stream_) {
+ // if there is already a reader on the same stream as us,
+ // it already synced with that writer and we can rely on
+ // the ongoing sync
+ bool found = false;
+ for (const auto& reader : reader_events) {
+ if (reader.stream == worker_stream->stream_) {
+ found = true;
+ break;
+ }
+ }
+ if (!found) {
+ AddEventHelper(&event_per_stream,
+ writer);
+ }
+ }
+ }
+ }
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // check for expired events and delete them
+ auto &reader_events = sync_obj.reader_events;
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ // if there are some readers, we wait for them
+ for (auto& cuda_event : reader_events) {
+ if (worker_stream->stream_ != cuda_event.stream) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
Review comment:
There can be only one event in the `writer_event` vector, `clear` here could be replaced by `pop_back`.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r658483058
##########
File path: include/mxnet/base.h
##########
@@ -348,10 +348,10 @@ struct RunContext {
* \brief the auxiliary stream of the device, can be nullptr or Stream<gpu>* in GPU mode
*/
void *aux_stream;
- /*!
- * \brief indicator of whether this execution is run in bulk mode
+ /*!
+ * \brief pointer to the cuda event pool used by the dependecy engine
*/
- bool is_bulk;
+ void *event_pool = nullptr;
Review comment:
Leaving this comment here just as a way to explain why the `is_bulk` is removed. `is_bulk` was a way to mark operations as not needing synchronization when being a part of a single engine bulk. This PR moves the handling of the synchronization to the engine itself, so this is no longer useful.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-947283086
@mxnet-bot run ci [centos-cpu, unix-cpu, windows-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha merged pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
szha merged pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-854464891
Hey @Kh4L , Thanks for submitting the PR
All tests are already queued to run once. If tests fail, you can trigger one or more tests again with the following commands:
- To trigger all jobs: @mxnet-bot run ci [all]
- To trigger specific jobs: @mxnet-bot run ci [job1, job2]
***
**CI supported jobs**: [unix-cpu, centos-gpu, sanity, windows-cpu, miscellaneous, centos-cpu, unix-gpu, windows-gpu, edge, website, clang]
***
_Note_:
Only following 3 categories can trigger CI :PR Author, MXNet Committer, Jenkins Admin.
All CI tests must pass before the PR can be merged.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-941956394
Jenkins CI successfully triggered : [sanity]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r701326230
##########
File path: src/engine/threaded_engine_perdevice.cc
##########
@@ -103,7 +118,11 @@ class ThreadedEnginePerDevice : public ThreadedEngine {
MSHADOW_CATCH_ERROR(mshadow::SetDevice<gpu>(ctx.dev_id));
#endif
}
- this->ExecuteOprBlock(RunContext{ctx, nullptr, nullptr, false}, opr_block);
+ CallbackOnStart on_start = this->CreateOnStart(ThreadedEngine::OnStartStatic,
+ opr_block);
+ CallbackOnComplete callback = this->CreateCallback(ThreadedEngine::OnCompleteStatic,
+ opr_block);
Review comment:
Mechanism introduced in this PR is in addition to the dependency tracking that already existed, and the OnCompleteStatic contains dependency update code (so it has to be called from the GPU version as well).
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-860489396
Jenkins CI successfully triggered : [unix-cpu, windows-gpu, edge, unix-gpu, centos-gpu, centos-cpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-884377070
@sandeep-krishnamurthy what's the review status?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sandeep-krishnamurthy commented on pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
sandeep-krishnamurthy commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-875125646
@Zha0q1 @mseth10 can you please help review this PR?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-918026503
@mxnet-bot run ci [unix-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-941956394
Jenkins CI successfully triggered : [sanity]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] TristonC commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
TristonC commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-861919393
@szha @sandeep-krishnamurthy Do you have someone in your team or elsewhere for a code review?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-918016478
@mxnet-bot run ci [website, unix-cpu, unix-gpu, centos-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r701325238
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
Review comment:
Because the pool is per stream so for 2 events from the same pool the one with the larger pool_index had to be recorded later (and synchronizing with that event will synchronize with both).
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r701324009
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaEventSynchronize(*ev));
+ }
+}
+
+void ThreadedEngine::OnStartGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& writer : sync_obj.writer_event) {
+ if (writer.event.expired()) {
+ sync_obj.writer_event.clear();
+ break;
+ }
+ if (writer.stream != worker_stream->stream_) {
+ // if there is already a reader on the same stream as us,
+ // it already synced with that writer and we can rely on
+ // the ongoing sync
+ bool found = false;
+ for (const auto& reader : reader_events) {
+ if (reader.stream == worker_stream->stream_) {
+ found = true;
+ break;
+ }
+ }
+ if (!found) {
+ AddEventHelper(&event_per_stream,
+ writer);
+ }
+ }
+ }
+ }
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // check for expired events and delete them
+ auto &reader_events = sync_obj.reader_events;
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ // if there are some readers, we wait for them
+ for (auto& cuda_event : reader_events) {
+ if (worker_stream->stream_ != cuda_event.stream) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ if (worker_stream->stream_ != sync_obj.writer_event[0].stream) {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaStreamWaitEvent(worker_stream->stream_, *ev, 0));
+ }
+}
+
+void ThreadedEngine::OnCompleteGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+
+ if (!use_new_dep_engine) {
+ worker_stream->Wait();
+ ThreadedEngine::OnCompleteStatic(engine, info->opr_block, error);
+ GPUWorkerSyncInfo::Delete(info);
+ return;
+ }
+
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ auto* event_pool = static_cast<CUDAEventPool*>(info->event_pool);
+ auto[event, event_pool_idx] = event_pool->GetNextEvent();
+ auto ev = event.lock();
+ MSHADOW_CUDA_CALL(cudaEventRecord(*ev, worker_stream->stream_));
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // If some reader event is already recorded on the same stream,
+ // we want to replace ourselves by it
+ int i;
+ for (i = 0; i < sync_obj.reader_events.size(); ++i) {
+ auto stream = sync_obj.reader_events[i].stream;
+ if (stream == worker_stream->stream_) {
+ sync_obj.reader_events[i].event = event;
+ sync_obj.reader_events[i].pool_index = event_pool_idx;
+ break;
+ }
+ }
+ if (i == sync_obj.reader_events.size()) {
+ sync_obj.reader_events.push_back({event, worker_stream->stream_, event_pool_idx});
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ sync_obj.reader_events.clear();
+ sync_obj.writer_event.clear();
+ sync_obj.writer_event.push_back({event, worker_stream->stream_, event_pool_idx});
Review comment:
Right - since if you are the writer to the variable, you had to wait for every previous reader/writer already there is no reason to keep those events anymore.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-941956362
@mxnet-bot run ci [sanity]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-941956362
@mxnet-bot run ci [sanity]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] barry-jin commented on a change in pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
barry-jin commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r733144257
##########
File path: src/engine/threaded_engine.cc
##########
@@ -533,5 +546,206 @@ void ThreadedEngine::OnCompleteStatic(Engine* engine, void* opr_block_, const dm
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine* engine, void* opr_block, const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+static inline bool IsEngineAsync() {
+ std::string type = dmlc::GetEnv("MXNET_ENGINE_TYPE", std::string(""));
+ std::string async_engine_tag("Async");
+ auto tag_pos = type.find(async_engine_tag);
+ return tag_pos != std::string::npos;
+}
+
+void ThreadedEngine::OnStartCPU(Engine* engine, void* opr_block, const dmlc::Error* error) {
+ static bool use_new_dep_engine = IsEngineAsync();
+ if (!use_new_dep_engine) {
+ return;
+ }
Review comment:
I think this part is not tested in CI because I didn't see any changes to use Async tag inside test functions. https://github.com/apache/incubator-mxnet/blob/master/ci/docker/runtime_functions.sh
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] szha commented on pull request #20331: [WIP] Add async GPU depency Engine
Posted by GitBox <gi...@apache.org>.
szha commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-865476458
I plan to take a closer look this week. could @ptrendx also take a look?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-918016570
Jenkins CI successfully triggered : [unix-gpu, unix-cpu, website, centos-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-918026579
Jenkins CI successfully triggered : [unix-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] Kh4L commented on a change in pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
Kh4L commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r735377849
##########
File path: ci/jenkins/Jenkinsfile_unix_gpu
##########
@@ -49,7 +49,8 @@ core_logic: {
custom_steps.test_unix_cpp_package_gpu('gpu'),
// TODO(szha): fix and reenable the hanging issue. tracked in #18098
// custom_steps.test_unix_distributed_kvstore_gpu('gpu'),
- custom_steps.test_unix_byteps_gpu('gpu'),
+ // TODO(spanev): reenable when byteps is updated with the new dep engine API
Review comment:
Good idea, tracked here https://github.com/apache/incubator-mxnet/issues/20697
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a change in pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
ptrendx commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r734715666
##########
File path: ci/jenkins/Jenkinsfile_unix_gpu
##########
@@ -49,7 +49,8 @@ core_logic: {
custom_steps.test_unix_cpp_package_gpu('gpu'),
// TODO(szha): fix and reenable the hanging issue. tracked in #18098
// custom_steps.test_unix_distributed_kvstore_gpu('gpu'),
- custom_steps.test_unix_byteps_gpu('gpu'),
+ // TODO(spanev): reenable when byteps is updated with the new dep engine API
Review comment:
Could you add an issue to track this?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] barry-jin commented on a change in pull request #20331: [WIP] Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
barry-jin commented on a change in pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#discussion_r701285974
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaEventSynchronize(*ev));
+ }
+}
+
+void ThreadedEngine::OnStartGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& writer : sync_obj.writer_event) {
+ if (writer.event.expired()) {
+ sync_obj.writer_event.clear();
+ break;
+ }
+ if (writer.stream != worker_stream->stream_) {
+ // if there is already a reader on the same stream as us,
+ // it already synced with that writer and we can rely on
+ // the ongoing sync
+ bool found = false;
+ for (const auto& reader : reader_events) {
+ if (reader.stream == worker_stream->stream_) {
+ found = true;
+ break;
+ }
+ }
+ if (!found) {
+ AddEventHelper(&event_per_stream,
+ writer);
+ }
+ }
+ }
+ }
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // check for expired events and delete them
+ auto &reader_events = sync_obj.reader_events;
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ // if there are some readers, we wait for them
+ for (auto& cuda_event : reader_events) {
+ if (worker_stream->stream_ != cuda_event.stream) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
Review comment:
Could you elaborate why should all the writer events be removed when the first writer event expires.
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
+ }
+ } else {
+ (*events_per_stream).emplace(event_stream, cuda_event);
+ }
+}
+
+void ThreadedEngine::OnStartCPU(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& cuda_event : reader_events) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaEventSynchronize(*ev));
+ }
+}
+
+void ThreadedEngine::OnStartGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+ if (!use_new_dep_engine) {
+ return;
+ }
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ std::unordered_map<cudaStream_t, EventInfo> event_per_stream;
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ auto &reader_events = sync_obj.reader_events;
+ // check for expired events and delete them
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ for (auto& writer : sync_obj.writer_event) {
+ if (writer.event.expired()) {
+ sync_obj.writer_event.clear();
+ break;
+ }
+ if (writer.stream != worker_stream->stream_) {
+ // if there is already a reader on the same stream as us,
+ // it already synced with that writer and we can rely on
+ // the ongoing sync
+ bool found = false;
+ for (const auto& reader : reader_events) {
+ if (reader.stream == worker_stream->stream_) {
+ found = true;
+ break;
+ }
+ }
+ if (!found) {
+ AddEventHelper(&event_per_stream,
+ writer);
+ }
+ }
+ }
+ }
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // check for expired events and delete them
+ auto &reader_events = sync_obj.reader_events;
+ reader_events.erase(std::remove_if(reader_events.begin(), reader_events.end(),
+ [&](const EventInfo e_i) {
+ return e_i.event.expired();
+ }), reader_events.end());
+ // if there are some readers, we wait for them
+ for (auto& cuda_event : reader_events) {
+ if (worker_stream->stream_ != cuda_event.stream) {
+ AddEventHelper(&event_per_stream, cuda_event);
+ }
+ }
+ if (!sync_obj.writer_event.empty()) {
+ if (sync_obj.writer_event[0].event.expired()) {
+ sync_obj.writer_event.clear();
+ } else {
+ if (worker_stream->stream_ != sync_obj.writer_event[0].stream) {
+ AddEventHelper(&event_per_stream, sync_obj.writer_event[0]);
+ }
+ }
+ }
+ }
+ for (auto event : event_per_stream) {
+ auto ev = event.second.event.lock();
+ MSHADOW_CUDA_CALL(cudaStreamWaitEvent(worker_stream->stream_, *ev, 0));
+ }
+}
+
+void ThreadedEngine::OnCompleteGPU(Engine *engine, void *sync_info,
+ const dmlc::Error* error) {
+ auto *info = reinterpret_cast<GPUWorkerSyncInfo *>(sync_info);
+ CHECK(info->stream != nullptr);
+
+ auto *worker_stream = reinterpret_cast<mshadow::Stream<gpu> *>(info->stream);
+ static bool use_new_dep_engine = dmlc::GetEnv("MXNET_ASYNC_GPU_ENGINE", false);
+
+ if (!use_new_dep_engine) {
+ worker_stream->Wait();
+ ThreadedEngine::OnCompleteStatic(engine, info->opr_block, error);
+ GPUWorkerSyncInfo::Delete(info);
+ return;
+ }
+
+ ThreadedOpr *threaded_opr = static_cast<OprBlock*>(info->opr_block)->opr;
+ auto* event_pool = static_cast<CUDAEventPool*>(info->event_pool);
+ auto[event, event_pool_idx] = event_pool->GetNextEvent();
+ auto ev = event.lock();
+ MSHADOW_CUDA_CALL(cudaEventRecord(*ev, worker_stream->stream_));
+ for (auto* read_var : threaded_opr->const_vars) {
+ auto &sync_obj = read_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ // If some reader event is already recorded on the same stream,
+ // we want to replace ourselves by it
+ int i;
+ for (i = 0; i < sync_obj.reader_events.size(); ++i) {
+ auto stream = sync_obj.reader_events[i].stream;
+ if (stream == worker_stream->stream_) {
+ sync_obj.reader_events[i].event = event;
+ sync_obj.reader_events[i].pool_index = event_pool_idx;
+ break;
+ }
+ }
+ if (i == sync_obj.reader_events.size()) {
+ sync_obj.reader_events.push_back({event, worker_stream->stream_, event_pool_idx});
+ }
+ }
+
+ for (auto* write_var : threaded_opr->mutable_vars) {
+ auto &sync_obj = write_var->sync_object;
+ std::lock_guard<std::mutex> l(sync_obj.mutex);
+ sync_obj.reader_events.clear();
+ sync_obj.writer_event.clear();
+ sync_obj.writer_event.push_back({event, worker_stream->stream_, event_pool_idx});
Review comment:
This probably answers my last question. But what's the purpose of clearing all the reader and writer events? Is it because current event should synchronize with the cleared reader and writer events?
##########
File path: src/engine/threaded_engine.cc
##########
@@ -523,5 +543,206 @@ void ThreadedEngine::OnCompleteStatic(Engine *engine, void *opr_block_,
OprBlock::Delete(opr_block);
}
+void ThreadedEngine::OnStartStatic(Engine *engine, void *opr_block,
+ const dmlc::Error* error) {
+ // no-op
+}
+
+#if MXNET_USE_CUDA
+static inline void AddEventHelper(
+ std::unordered_map<cudaStream_t, EventInfo>* events_per_stream,
+ const EventInfo& cuda_event) {
+ auto event_stream = cuda_event.stream;
+ if (events_per_stream->count(event_stream) > 0) {
+ if ((*events_per_stream)[event_stream].pool_index < cuda_event.pool_index) {
+ (*events_per_stream)[event_stream] = cuda_event;
Review comment:
Why should the event be replaced by others with larger pool_index?
##########
File path: src/engine/threaded_engine_perdevice.cc
##########
@@ -103,7 +118,11 @@ class ThreadedEnginePerDevice : public ThreadedEngine {
MSHADOW_CATCH_ERROR(mshadow::SetDevice<gpu>(ctx.dev_id));
#endif
}
- this->ExecuteOprBlock(RunContext{ctx, nullptr, nullptr, false}, opr_block);
+ CallbackOnStart on_start = this->CreateOnStart(ThreadedEngine::OnStartStatic,
+ opr_block);
+ CallbackOnComplete callback = this->CreateCallback(ThreadedEngine::OnCompleteStatic,
+ opr_block);
Review comment:
Why use OnStartStatic and OnCompleteStatic here instead of creating GPU onstart and oncomplete for GPU case, creating static-onstart and static-oncomplete for CPU case.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #20331: Add async GPU dependency Engine
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #20331:
URL: https://github.com/apache/incubator-mxnet/pull/20331#issuecomment-945325414
Jenkins CI successfully triggered : [unix-gpu]
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org