You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "LeshengJin (via GitHub)" <gi...@apache.org> on 2023/09/18 19:30:49 UTC
[GitHub] [tvm] LeshengJin opened a new pull request, #15776: [Disco] Integrate RCCL
LeshengJin opened a new pull request, #15776:
URL: https://github.com/apache/tvm/pull/15776
This PR integrates RCCL for amd multi-GPU parallelism.
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] LeshengJin commented on a diff in pull request #15776: [Disco] Integrate RCCL
Posted by "LeshengJin (via GitHub)" <gi...@apache.org>.
LeshengJin commented on code in PR #15776:
URL: https://github.com/apache/tvm/pull/15776#discussion_r1330653754
##########
src/runtime/disco/ccl/utils.h:
##########
@@ -19,22 +19,21 @@
#ifndef TVM_RUNTIME_DISCO_NCCL_UTILS_H_
#define TVM_RUNTIME_DISCO_NCCL_UTILS_H_
-#include <nccl.h>
#include <tvm/runtime/data_type.h>
#include <tvm/runtime/disco/session.h>
#include "../utils.h"
namespace tvm {
namespace runtime {
-namespace nccl {
+namespace ccl {
-#define NCCL_CALL(cmd) \
- do { \
- ncclResult_t r = cmd; \
- if (r != ncclSuccess) { \
- LOG(FATAL) << "NCCLErrror: " << ncclGetErrorString(r); \
- } \
+#define NCCL_CALL(cmd) \
Review Comment:
```suggestion
#define CCL_CALL(cmd) \
```
Do you think it is necessary to do a renaming of functions in `utils.h`?
##########
src/runtime/disco/ccl/utils.h:
##########
@@ -19,22 +19,21 @@
#ifndef TVM_RUNTIME_DISCO_NCCL_UTILS_H_
#define TVM_RUNTIME_DISCO_NCCL_UTILS_H_
-#include <nccl.h>
#include <tvm/runtime/data_type.h>
#include <tvm/runtime/disco/session.h>
#include "../utils.h"
namespace tvm {
namespace runtime {
-namespace nccl {
+namespace ccl {
-#define NCCL_CALL(cmd) \
- do { \
- ncclResult_t r = cmd; \
- if (r != ncclSuccess) { \
- LOG(FATAL) << "NCCLErrror: " << ncclGetErrorString(r); \
- } \
+#define NCCL_CALL(cmd) \
Review Comment:
```suggestion
#define CCL_CALL(cmd) \
```
Do you think it is necessary to do a renaming of functions in `utils.h`? @junrushao
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao commented on a diff in pull request #15776: [Disco] Integrate RCCL
Posted by "junrushao (via GitHub)" <gi...@apache.org>.
junrushao commented on code in PR #15776:
URL: https://github.com/apache/tvm/pull/15776#discussion_r1330656037
##########
tests/python/disco/test_rccl.py:
##########
@@ -0,0 +1,376 @@
+# Licensed to the Apache Software Foundation (ASF) under one
Review Comment:
you may parameterize the original `test_nccl.py`
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] LeshengJin commented on a diff in pull request #15776: [Disco] Integrate RCCL
Posted by "LeshengJin (via GitHub)" <gi...@apache.org>.
LeshengJin commented on code in PR #15776:
URL: https://github.com/apache/tvm/pull/15776#discussion_r1330654090
##########
src/runtime/disco/ccl/utils.h:
##########
@@ -19,22 +19,21 @@
#ifndef TVM_RUNTIME_DISCO_NCCL_UTILS_H_
#define TVM_RUNTIME_DISCO_NCCL_UTILS_H_
-#include <nccl.h>
#include <tvm/runtime/data_type.h>
#include <tvm/runtime/disco/session.h>
#include "../utils.h"
namespace tvm {
namespace runtime {
-namespace nccl {
+namespace ccl {
-#define NCCL_CALL(cmd) \
- do { \
- ncclResult_t r = cmd; \
- if (r != ncclSuccess) { \
- LOG(FATAL) << "NCCLErrror: " << ncclGetErrorString(r); \
- } \
+#define NCCL_CALL(cmd) \
+ do { \
+ ncclResult_t r = cmd; \
+ if (r != ncclSuccess) { \
+ LOG(FATAL) << CCLString "Errror: " << ncclGetErrorString(r); \
+ } \
} while (0)
inline ncclDataType_t AsNCCLDataType(runtime::DataType dtype) {
Review Comment:
```suggestion
inline ncclDataType_t AsCCLDataType(runtime::DataType dtype) {
```
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao commented on a diff in pull request #15776: [Disco] Integrate RCCL
Posted by "junrushao (via GitHub)" <gi...@apache.org>.
junrushao commented on code in PR #15776:
URL: https://github.com/apache/tvm/pull/15776#discussion_r1330808771
##########
src/runtime/disco/ccl/ccl.cc:
##########
@@ -16,9 +16,48 @@
* specific language governing permissions and limitations
* under the License.
*/
+#ifndef TVM_NCCL_RCCL_SWITCH
+#define TVM_NCCL_RCCL_SWITCH 0 // 0: NCCL, 1: RCCL
+#endif
+
+#if TVM_NCCL_RCCL_SWITCH == 0
#include <cuda_runtime_api.h>
-#include <dlpack/dlpack.h>
#include <nccl.h>
+
+#include "../../cuda/cuda_common.h"
+
+using runtimeStream_t = cudaStream_t;
+
+#define runtimeCall CUDA_CALL
+#define SetDevice cudaSetDevice
+#define StreamCreate cudaStreamCreate
+#define StreamSync cudaStreamSynchronize
+#define StreamDestroy cudaStreamDestroy
+#define CommDestroy ncclCommDestroy
+#define RuntimeDeviceType DLDeviceType::kDLCUDA
+#define runtimeString "cuda"
+#define CCLString "nccl"
Review Comment:
if possible, avoid `#define` for constant values, and we should follow google c++ naming style:
```suggestion
#define TVM_DISCO_DEVICE_CALL CUDA_CALL
#define TVM_DISCO_DEVICE_SET_DEVICE cudaSetDevice
#define TVM_DISCO_DEVICE_STREAM_CREATE cudaStreamCreate
#define TVM_DISCO_DEVICE_STREAM_SYNC cudaStreamSynchronize
#define TVM_DISCO_DEVICE_STREAM_DESTROY cudaStreamDestroy
#define TVM_DISCO_DEVICE_NAME "cuda"
#define TVM_DISCO_CCL_DESTROY ncclCommDestroy
#define TVM_DISCO_CCL_NAME "nccl"
const constexpr DLDeviceType = DLDeviceType::kDLCUDA;
```
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao commented on a diff in pull request #15776: [Disco] Integrate RCCL
Posted by "junrushao (via GitHub)" <gi...@apache.org>.
junrushao commented on code in PR #15776:
URL: https://github.com/apache/tvm/pull/15776#discussion_r1330804996
##########
tests/python/disco/test_ccl.py:
##########
@@ -28,22 +28,51 @@
from tvm.runtime import disco as di
from tvm.runtime.relax_vm import VirtualMachine
from tvm.script import relax as R
+from tvm import get_global_func
_all_session_kinds = [di.ThreadedSession, di.ProcessSession]
-
-
-@pytest.mark.parametrize("session_kind", _all_session_kinds)
-def test_init(session_kind):
+_ccl = get_global_func("runtime.disco.compiled_ccl")()
+_test_inputs = [(sess_kind, _ccl) for sess_kind in _all_session_kinds]
+
+
+def create_device_target(ccl):
+ if ccl == "nccl":
Review Comment:
can you use `target.from_device("cuda:0")` instead?
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] jinhongyii merged pull request #15776: [Disco] Integrate RCCL
Posted by "jinhongyii (via GitHub)" <gi...@apache.org>.
jinhongyii merged PR #15776:
URL: https://github.com/apache/tvm/pull/15776
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao commented on a diff in pull request #15776: [Disco] Integrate RCCL
Posted by "junrushao (via GitHub)" <gi...@apache.org>.
junrushao commented on code in PR #15776:
URL: https://github.com/apache/tvm/pull/15776#discussion_r1330805517
##########
tests/python/disco/test_ccl.py:
##########
@@ -28,22 +28,51 @@
from tvm.runtime import disco as di
from tvm.runtime.relax_vm import VirtualMachine
from tvm.script import relax as R
+from tvm import get_global_func
_all_session_kinds = [di.ThreadedSession, di.ProcessSession]
-
-
-@pytest.mark.parametrize("session_kind", _all_session_kinds)
-def test_init(session_kind):
+_ccl = get_global_func("runtime.disco.compiled_ccl")()
+_test_inputs = [(sess_kind, _ccl) for sess_kind in _all_session_kinds]
Review Comment:
we should test "sess_kind x ccl" instead
--
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@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org