internal review #6

Chao1Han · 2024-09-23T08:37:53Z

Fixes #ISSUE_NUMBER

zhangxiaoli73 · 2024-09-30T06:27:53Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+      auto currentTimepoint = std::chrono::steady_clock::now();
+      auto timeElapsed = std::chrono::duration_cast<std::chrono::milliseconds>(
+          currentTimepoint - workStartTime_);
+      std::chrono::milliseconds opTimeout = std::chrono::milliseconds(60000);


where do you use it?

zhangxiaoli73 · 2024-10-10T01:31:54Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

@@ -67,17 +67,15 @@ ccl::reduction getXcclReduceOp(const ReduceOp& reduceOp, at::Tensor& input) {
    return xcclOps.at(reduceOp);
  } catch (const std::out_of_range&) {
    switch (reduceOp) {


No need to switch

zhangxiaoli73 · 2024-10-12T01:41:22Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+    : Work(rank, opType, "profilingTitle", inputs),
+      device_(device),
+      workStartTime_(std::chrono::steady_clock::now()) {
+  unsigned char enable_timing = 0;


If you always set it as 0, then we don't need to keep it, right?

Yes, Defining this variable serves as a form of annotation, informing reviewers and users that 0 represents the state of enable_timing, which is meaningful.

zhangxiaoli73 · 2024-10-12T01:42:34Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+            "Work ran for ",
+            timeElapsed.count(),
+            " milliseconds before timing out.");
+        TORCH_CHECK(false, exceptionMsg)


TORCH_CHECK(false, exceptionMsg); abort();

zhangxiaoli73 · 2024-10-12T01:44:51Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+} // namespace
+
+static std::mutex xcclCommDevIdxMapMutex;
+static std::unordered_map<std::shared_ptr<xcclComm_t>, int> xcclCommDevIdxMap;


Those static variables are not used in your code. Please check.

zhangxiaoli73 · 2024-10-12T01:46:03Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+  blockingWait_ = getCvarBool(TORCH_XCCL_BLOCKING_WAIT, false);
+  init();
+
+  // Intel oneCCL requires passing CCL_LOCAL_RANK and CCL_LOCAL_SIZE for non-MPI


More comment for why we use LOCAL_RANK and LOCAL_WORLD_SIZE.

zhangxiaoli73 · 2024-10-12T01:46:21Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+std::shared_ptr<xcclComm_t> ProcessGroupXCCL::getXCCLComm(
+    const std::string& deviceKey,
+    at::Device& device) {
+  if (deviceKey.empty()) {


C10_THROW_ERROR_WITH

zhangxiaoli73 · 2024-10-12T01:48:12Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+    devXCCLCommMap_.emplace(deviceKey, XCCLComm);
+  }
+
+  xcclStreamsMap_.emplace(deviceKey, std::move(stream));


so xcclEventsMap does not needed?

zhangxiaoli73 · 2024-10-12T01:49:06Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+    PreProcess pre,
+    PostProcess post,
+    OpType opType) {
+  using traits = function_traits<Fn>;


which collective need attribute as a must?

Yes, allgather meet build error

zhangxiaoli73 · 2024-10-12T01:50:03Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+  for (const auto i : c10::irange(inputs.size())) {
+    c10::xpu::XPUCachingAllocator::recordStream(
+        inputs[i].storage().data_ptr(), stream);
+    fn(inputs[i], outputs[i], attr, *comm, stream);


add comment for output record stream.

zhangxiaoli73 · 2024-10-12T01:52:29Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.hpp

+          false, "ProcessGroupXCCL::WorkXCCL::isSuccess not implemented");
+    }
+
+    void abort() override {


abort here?

zhangxiaoli73 · 2024-10-12T01:52:38Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.hpp

+
+    bool isCompleted() override;
+
+    bool isSuccess() const override {


zhangxiaoli73 · 2024-10-13T03:17:33Z

torch/csrc/distributed/c10d/ProcessGroupNCCL.cpp

      case ReduceOp::BXOR:
-        C10_THROW_ERROR(ValueError, "Cannot use ReduceOp.BXOR with NCCL");
+        C10_THROW_ERROR(


don't change NCCL now.

zhangxiaoli73 · 2024-10-13T12:38:08Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+
+  c10::impl::VirtualGuardImpl impl(device.type());
+  c10::Stream stream = impl.getStream(device);
+  sycl::queue& q = c10::xpu::XPUStream(stream).queue();


It's a big bug to use current stream as communication stream.

zhangxiaoli73 · 2024-10-17T06:19:39Z

torch/csrc/distributed/c10d/ProcessGroupXCCL.cpp

+    int rank,
+    OpType opType,
+    const std::optional<std::vector<at::Tensor>>& inputs)
+    : Work(rank, opType, "profilingTitle", inputs),


Need change

zhangxiaoli73 · 2024-10-21T01:10:16Z

torch/csrc/distributed/c10d/ProcessGroup.hpp

@@ -126,6 +131,13 @@ class TORCH_API ProcessGroup : public torch::CustomClassHolder {
    return backendType_;
  };

+  inline bool backendSupportsSequenceNumbers(BackendType backendType) {
+    if (backendType == BackendType::GLOO || backendType == BackendType::NCCL ||
+        backendType == BackendType::XCCL || backendType == BackendType::UCC)


Do you make sure that we need to support this sequence number?

Sequence number used by RECORD_PARAM_COMMS. so we need it

zhangxiaoli73 · 2024-10-21T01:15:04Z

torch/testing/_internal/common_distributed.py

@@ -180,7 +181,8 @@ def skip_if_lt_x_gpu(x):
    def decorator(func):
        @wraps(func)
        def wrapper(*args, **kwargs):
-            if torch.cuda.is_available() and torch.cuda.device_count() >= x:
+            if (torch.cuda.is_available() and torch.cuda.device_count() >= x) or \


Don't use if for accelerator related check

See pytorch#140725 (comment) Running `torch.mps.synchronize()` after metal kernel resulted in infinite wait inside `[_MTLCommandBuffer waitUntilCompleted]` ``` (lldb) bt * thread #1, queue = 'com.apple.main-thread', stop reason = signal SIGSTOP * frame #0: 0x00000001aa919084 Metal`pthread_cond_wait + 12 frame #1: 0x00000001aa78b1b4 Metal`-[_MTLCommandBuffer waitUntilCompleted] + 84 frame #2: 0x00000001032bf358 libtorch_python.dylib`torch::mps::MPSModule_deviceSynchronize(_object*, _object*) + 40 frame #3: 0x0000000100e94c20 Python`cfunction_vectorcall_NOARGS + 100 frame #4: 0x0000000100e389b8 Python`PyObject_Vectorcall + 92 frame #5: 0x0000000100f61e38 Python`_PyEval_EvalFrameDefault + 19040 frame #6: 0x0000000100f5d180 Python`PyEval_EvalCode + 200 frame #7: 0x0000000100fcd1a4 Python`run_eval_code_obj + 104 frame #8: 0x0000000100fccbe4 Python`run_mod + 168 frame #9: 0x0000000100fcb518 Python`pyrun_file + 164 frame #10: 0x0000000100fca854 Python`_PyRun_SimpleFileObject + 256 frame #11: 0x0000000100fca4e8 Python`_PyRun_AnyFileObject + 80 frame #12: 0x0000000100ff2028 Python`pymain_run_file_obj + 164 frame #13: 0x0000000100ff1ce4 Python`pymain_run_file + 72 frame pytorch#14: 0x0000000100ff0f74 Python`Py_RunMain + 988 frame pytorch#15: 0x0000000100ff1564 Python`pymain_main + 304 frame pytorch#16: 0x0000000100ff1604 Python`Py_BytesMain + 40 frame pytorch#17: 0x000000019f630274 dyld`start + 2840 ``` Pull Request resolved: pytorch#141296 Approved by: https://github.com/huydhn

…143550) # Motivation Fix pytorch#143543 # Solution We should raise python exception instead of aborting... # Additional Context without this PR: ```python >>> import torch >>> torch.accelerator.current_stream(torch.accelerator.device_count()) terminate called after throwing an instance of 'c10::Error' what(): device is out of range, device is 2, total number of device is 2. Exception raised from check_device_index at /home/dvrogozh/git/pytorch/pytorch/c10/xpu/XPUFunctions.h:36 (most recent call first): frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0xac (0x7f30707eb95c in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libc10.so) frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xf3 (0x7f307078fc57 in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libc10.so) frame #2: <unknown function> + 0x19a3e (0x7f3070c2ba3e in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libc10_xpu.so) frame #3: c10::xpu::getCurrentXPUStream(signed char) + 0x2f (0x7f3070c2c83f in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libc10_xpu.so) frame #4: <unknown function> + 0x1ca35 (0x7f3070c2ea35 in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libc10_xpu.so) frame #5: <unknown function> + 0x653f15 (0x7f3083391f15 in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libtorch_python.so) frame #6: <unknown function> + 0x39e5f2 (0x7f30830dc5f2 in /home/dvrogozh/git/pytorch/pytorch/torch/lib/libtorch_python.so) <omitting python frames> frame pytorch#20: <unknown function> + 0x29d90 (0x7f308b19bd90 in /lib/x86_64-linux-gnu/libc.so.6) frame pytorch#21: __libc_start_main + 0x80 (0x7f308b19be40 in /lib/x86_64-linux-gnu/libc.so.6) Aborted (core dumped) ``` with this PR: ```python >>> import torch >>> torch.accelerator.current_stream(torch.accelerator.device_count()) Traceback (most recent call last): File "<stdin>", line 1, in <module> File "/home/pt-gpu/4T-4652/guangyey/stock-pytorch/torch/accelerator/__init__.py", line 123, in current_stream return torch._C._accelerator_getStream(device_index) RuntimeError: The device index is out of range. It must be in [0, 2), but got 2. ``` Pull Request resolved: pytorch#143550 Approved by: https://github.com/EikanWang, https://github.com/dvrogozh, https://github.com/albanD

…pytorch#144120) (pytorch#146372) Summary: # Summary ### Sticky points Cuda-graph rng handling has changed / deviated from original implementation. We will be left with a dangling 'offset' val and confusing naming due to BC ## Dependencies - Flash PR: Dao-AILab/flash-attention#1419 ### Other Points - The BC linter is complaining about losing generate.py and its functions which is not real BC surface cc albanD imported-using-ghimport Test Plan: Imported from OSS Building in dev `buck build @//mode/dev-nosan -c fbcode.nvcc_arch=h100a //caffe2:ATen-cu --show-full-output ` I and Nming the .so I do see that the flash symbols are correctly named: ``` 0000000001c3dfb0 t pytorch_flash::run_mha_bwd(pytorch_flash::Flash_bwd_params&, CUstream_st*)::$_0::operator()() const::{lambda()#1}::operator()() const::{lambda()#1}::operator()() const::{lambda()#7}::operator()() const 0000000001c36080 t pytorch_flash::run_mha_fwd(pytorch_flash::Flash_fwd_params&, CUstream_st*, bool)::$_0::operator()() const::{lambda()#2}::operator()() const::{lambda()#1}::operator()() const::{lambda()#6}::operator()() const 0000000001c360e0 t pytorch_flash::run_mha_fwd(pytorch_flash::Flash_fwd_params&, CUstream_st*, bool)::$_0::operator()() const::{lambda()#2}::operator()() const::{lambda()#1}::operator()() const::{lambda()#7}::operator()() const 0000000001c35fc0 t pytorch_flash::run_mha_fwd(pytorch_flash::Flash_fwd_params&, CUstream_st*, bool)::$_0::operator()() const::{lambda()#1}::operator()() const::{lambda()#1}::operator()() const::{lambda()#6}::operator()() const 0000000001c36020 t pytorch_flash::run_mha_fwd(pytorch_flash::Flash_fwd_params&, CUstream_st*, bool)::$_0::operator()() const::{lambda()#1}::operator()() const::{lambda()#1}::operator()() const::{lambda()#7}::operator()() const ``` Reviewed By: vkuzo Differential Revision: D68502879 Pulled By: drisspg Pull Request resolved: pytorch#146372 Approved by: https://github.com/jbschlosser

Chao1Han added 8 commits September 20, 2024 03:02

Xccl process group for Pytorch

652da01

Merge remote-tracking branch 'upstream/main' into xccl-bak

0cb0016

Align latest

a71d69a

hidden env

af6f03c

refine findccl code

88bea25

Add comments for build xccl

f6ea934

refine workxccl

1226e3b

refine timeout

d62e0be

zhangxiaoli73 reviewed Sep 30, 2024

View reviewed changes

Chao1Han added 5 commits September 30, 2024 06:33

rm head

714de2a

update

0923781

minor fix

31d092d

rm duplicate code and refine cmake

cbea299

update cmake

ef261c6

zhangxiaoli73 reviewed Oct 10, 2024

View reviewed changes

Chao1Han added 5 commits October 10, 2024 02:12

hidden xccl specific

6c648cd

fix ci fail

e621fe6

rm vir fun and modify tensor check

3f225d9

Merge branch 'xccl-bak' into xccl-bak2

1138a4a

refine collective, getcomm

8e5e78a

zhangxiaoli73 reviewed Oct 12, 2024

View reviewed changes

Chao1Han added 5 commits October 12, 2024 02:23

accept comments

1267963

rm attr

3d55b85

add default ccl root dir

f69059a

update

bed720c

update

fd44abe

zhangxiaoli73 reviewed Oct 13, 2024

View reviewed changes

code refine

d12b922

zhangxiaoli73 reviewed Oct 13, 2024

View reviewed changes

Chao1Han added 6 commits October 14, 2024 02:38

minor fix

b57e812

update

5968f0f

update

edba8aa

Refine specific code

56a5e7f

accept comments

a062f9f

Merge branch 'xccl-bak' into xccl-bak2

ae90994

zhangxiaoli73 reviewed Oct 17, 2024

View reviewed changes

Chao1Han added 3 commits October 17, 2024 08:17

rm header and refine profilehead

ab04fc0

add get_device_count

4ee49fb

add backendSupportsSequenceNumbers

1a2c9c2

zhangxiaoli73 reviewed Oct 21, 2024

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

internal review #6

internal review #6

Chao1Han commented Sep 23, 2024

zhangxiaoli73 Sep 30, 2024

zhangxiaoli73 Oct 10, 2024

zhangxiaoli73 Oct 12, 2024

Chao1Han Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

Chao1Han Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

Chao1Han Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 12, 2024

zhangxiaoli73 Oct 13, 2024

zhangxiaoli73 Oct 13, 2024

zhangxiaoli73 Oct 17, 2024

zhangxiaoli73 Oct 21, 2024

Chao1Han Oct 21, 2024

zhangxiaoli73 Oct 21, 2024


		bool isCompleted() override;

		bool isSuccess() const override {

internal review #6

Are you sure you want to change the base?

internal review #6

Conversation

Chao1Han commented Sep 23, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment