Skip to content

XCCL Backend support #1

Open
Chao1Han wants to merge 6 commits into
mainfrom
xccl_sycl
Open

XCCL Backend support #1
Chao1Han wants to merge 6 commits into
mainfrom
xccl_sycl

Conversation

@Chao1Han

@Chao1Han Chao1Han commented Nov 3, 2025

Copy link
Copy Markdown
Owner

RFC: #2
Build
export USE_XCCL=ON
export USE_NCCLX=OFF
export USE_NCCL=OFF
export USE_GLOO=OFF
export USE_SYSTEM_LIBS=1
conda install fmt -y
source oneCCL2/env/setvars.sh // direct use oneCCL2 C API
python setup.py install

Comment thread comms/torchcomms/device/XpuApi.cpp Outdated
prop->maxThreadsDim[2] = max_work_item_sizes[2];

// Max grid size - not directly available in SYCL, use reasonable defaults
prop->maxGridSize[0] = 2147483647;

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

then, should we submit request to sycl?

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In the current NCCL implementation, getDeviceProperties isn’t actually used anywhere; it merely fills in some values to maintain consistency with NCCL’s format. In fact, within PyTorch, the DeviceProperties functions for XPU and CUDA are entirely different. Once this function becomes actively used in the future, I’ll coordinate with guangye to discuss how to modify it.

int device = ::c10::xpu::current_device();
sycl::device& dev = ::c10::xpu::get_raw_device(device);

*devPtr = sycl::malloc_device(size, dev, ctx);

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

So it does its own memory management? Is there any cache strategy?

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just a wrapper of sycl malloc api, used by barrier(create and cache tiny buffer allreduce)


xpu_result_t DefaultXpuApi::eventCreate(xpuEvent_t& event) {
try {
event = ::at::xpu::XPUEvent(false); // No timing

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We can expect timing enable in basekit 2025.2, right?

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes.

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

then, should we enable it?

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

which in next fun eventCreateWithFlags

Comment thread comms/torchcomms/xccl/TorchWorkXCCL.hpp Outdated
class TorchCommWindowXCCL;

// Forward declaration for test class
namespace test {

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

why it's needed?

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don’t think it’s necessary, but since other backends have reserved this interface, it’s likely intended for future use.

//
// NOTE: finalize MUST return without holding references to any work object,
// otherwise it may leak object and cause side effects.
stream_work_queues_.clear();

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

what does those two queues for?

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Store the works queued on each stream to allow the background cleanup thread to traverse the queue of every stream.

}

onecclResult_t DefaultXcclApi::commDeregister(onecclComm_t comm, void *handle) {
// return onecclCommDeregister(comm, handle);

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please check if we have submit feature request to oneCCL. If not, then a ticket is needed.

@zhangxiaoli73

Copy link
Copy Markdown

I don't get where is the profiler enabled.

checkAndAbortIfTimedOutOrError();
ensureTensorContiguous(tensor);

tracing_->recordEventWithInputOutput("all_reduce", rank_, {tensor}, {tensor});

Copy link
Copy Markdown
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@zhangxiaoli73 here is profiler register. https://github.com/meta-pytorch/torchcomms/blob/f0a411256447066c9b384727ad785ce8dc8e3ca0/comms/torchcomms/TorchCommTracing.cpp#L31 call RECORD_PARAM_COMMS_DATA, same profiler helper with processgroup

Chao1Han pushed a commit that referenced this pull request Feb 27, 2026
Summary:
Pull Request resolved: meta-pytorch#326

add test building/running support

Test Plan:
```
ctest --build-and-test ./ ./build --build-generator "Ninja" --build-options -DBUILD_TESTS=ON --test-command ctest --output-on-failure

    Start 1: TorchCommOptionsTest.EnvToValueImplBool
1/3 Test #1: TorchCommOptionsTest.EnvToValueImplBool ...   Passed    0.31 sec
    Start 2: TorchCommOptionsTest.StringToBool
2/3 Test #2: TorchCommOptionsTest.StringToBool .........   Passed    0.32 sec
    Start 3: TorchCommFactoryTest
3/3 Test #3: TorchCommFactoryTest ......................   Passed    0.32 sec

100% tests passed, 0 tests failed out of 3

Total Test time (real) =   0.95 sec
```

imported-using-ghimport

Reviewed By: d4l3k

Differential Revision: D91159066

Pulled By: dolpm

fbshipit-source-id: 77758d1d2d82535001e03cdfd13be299adece2aa
Chao1Han pushed a commit that referenced this pull request Feb 27, 2026
Summary:
Pull Request resolved: meta-pytorch#326

add test building/running support

Test Plan:
```
ctest --build-and-test ./ ./build --build-generator "Ninja" --build-options -DBUILD_TESTS=ON --test-command ctest --output-on-failure

    Start 1: TorchCommOptionsTest.EnvToValueImplBool
1/3 Test #1: TorchCommOptionsTest.EnvToValueImplBool ...   Passed    0.31 sec
    Start 2: TorchCommOptionsTest.StringToBool
2/3 Test #2: TorchCommOptionsTest.StringToBool .........   Passed    0.32 sec
    Start 3: TorchCommFactoryTest
3/3 Test #3: TorchCommFactoryTest ......................   Passed    0.32 sec

100% tests passed, 0 tests failed out of 3

Total Test time (real) =   0.95 sec
```

imported-using-ghimport

Reviewed By: d4l3k

Differential Revision: D91159066

Pulled By: dolpm

fbshipit-source-id: 77758d1d2d82535001e03cdfd13be299adece2aa
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants