Skip to content

Fix async error reporting and memory type handling in CCLAdaptor and DeviceAdaptor #266

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
MC952-arch merged 2 commits into from
Oct 15, 2025
Merged

Fix async error reporting and memory type handling in CCLAdaptor and DeviceAdaptor #266

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
49931c2
Ensure that the default type is flagcxMemDevice for deviceMalloc/Free…
MC952-arch Oct 2, 2025
88f7086
Change the input argument flagcxResult_t to flagcxResult_t* in flagcx…
MC952-arch Oct 2, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 9 additions & 6 deletions .github/workflows/test.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,12 +50,15 @@ jobs:
apt update -y
apt-get install clang-format -y
git fetch --all
from_ref=${GITHUB_HEAD_REF}
to_ref=${GITHUB_BASE_REF}

echo "From branch: $from_ref, To branch: $to_ref"

pre-commit run --from-ref "origin/$from_ref" --to-ref "origin/$to_ref"
if [ -n "$GITHUB_HEAD_REF" ] && [ -n "$GITHUB_BASE_REF" ]; then
from_ref="origin/$GITHUB_HEAD_REF"
to_ref="origin/$GITHUB_BASE_REF"
else
from_ref="HEAD^"
to_ref="HEAD"
fi
echo "From reference: $from_ref; To reference: $to_ref"
pre-commit run --from-ref "$from_ref" --to-ref "$to_ref"
continue-on-error: false

- name: Check the current working directory
Expand Down
15 changes: 9 additions & 6 deletions .github/workflows/torch-api-test.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,12 +47,15 @@ jobs:
run: |
cd /__w/FlagCX/FlagCX
git fetch --all
from_ref=${GITHUB_HEAD_REF}
to_ref=${GITHUB_BASE_REF}

echo "From branch: $from_ref, To branch: $to_ref"

pre-commit run --from-ref "origin/$from_ref" --to-ref "origin/$to_ref"
if [ -n "$GITHUB_HEAD_REF" ] && [ -n "$GITHUB_BASE_REF" ]; then
from_ref="origin/$GITHUB_HEAD_REF"
to_ref="origin/$GITHUB_BASE_REF"
else
from_ref="HEAD^"
to_ref="HEAD"
fi
echo "From reference: $from_ref; To reference: $to_ref"
pre-commit run --from-ref "$from_ref" --to-ref "$to_ref"
continue-on-error: false

- name: Run `make` to build the project
Expand Down
15 changes: 9 additions & 6 deletions .github/workflows/unit-test.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,12 +50,15 @@ jobs:
apt update -y
apt-get install clang-format -y
git fetch --all
from_ref=${GITHUB_HEAD_REF}
to_ref=${GITHUB_BASE_REF}

echo "From branch: $from_ref, To branch: $to_ref"

pre-commit run --from-ref "origin/$from_ref" --to-ref "origin/$to_ref"
if [ -n "$GITHUB_HEAD_REF" ] && [ -n "$GITHUB_BASE_REF" ]; then
from_ref="origin/$GITHUB_HEAD_REF"
to_ref="origin/$GITHUB_BASE_REF"
else
from_ref="HEAD^"
to_ref="HEAD"
fi
echo "From reference: $from_ref; To reference: $to_ref"
pre-commit run --from-ref "$from_ref" --to-ref "$to_ref"
continue-on-error: false

- name: Build Google Test
Expand Down
2 changes: 1 addition & 1 deletion flagcx/adaptor/ccl/bootstrap_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -80,7 +80,7 @@ flagcxResult_t bootstrapAdaptorCommUserRank(const flagcxInnerComm_t comm,

// TODO: unsupported
flagcxResult_t bootstrapAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return flagcxNotSupported;
}

Expand Down
5 changes: 2 additions & 3 deletions flagcx/adaptor/ccl/cncl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -106,10 +106,9 @@ flagcxResult_t cnclAdaptorCommUserRank(const flagcxInnerComm_t comm,
return (flagcxResult_t)c2f_ret_map[cnclGetCommRank(rank, comm->base)];
}

// TODO: change params's type from flagcxResult_t to flagcxResult_t*
flagcxResult_t cnclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
asyncError = c2f_ret_map[cnclGetCommAsyncError(comm->base)];
flagcxResult_t *asyncError) {
*asyncError = c2f_ret_map[cnclGetCommAsyncError(comm->base)];
return flagcxSuccess;
}

Expand Down
4 changes: 2 additions & 2 deletions flagcx/adaptor/ccl/dunccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,9 +67,9 @@ flagcxResult_t duncclAdaptorCommUserRank(const flagcxInnerComm_t comm,
}

flagcxResult_t duncclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return (flagcxResult_t)ncclCommGetAsyncError(comm->base,
(ncclResult_t *)&asyncError);
(ncclResult_t *)asyncError);
}

// TODO: unsupported
Expand Down
2 changes: 1 addition & 1 deletion flagcx/adaptor/ccl/gloo_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -111,7 +111,7 @@ flagcxResult_t glooAdaptorCommUserRank(const flagcxInnerComm_t comm,

// TODO: unsupported
flagcxResult_t glooAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return flagcxNotSupported;
}

Expand Down
2 changes: 1 addition & 1 deletion flagcx/adaptor/ccl/hccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -137,7 +137,7 @@ flagcxResult_t hcclAdaptorCommUserRank(const flagcxInnerComm_t comm,

// TODO: unsupported
flagcxResult_t hcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return flagcxNotSupported;
}

Expand Down
4 changes: 2 additions & 2 deletions flagcx/adaptor/ccl/ixnccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,9 +67,9 @@ flagcxResult_t ixncclAdaptorCommUserRank(const flagcxInnerComm_t comm,
}

flagcxResult_t ixncclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return (flagcxResult_t)ncclCommGetAsyncError(comm->base,
(ncclResult_t *)&asyncError);
(ncclResult_t *)asyncError);
}

// TODO: unsupported
Expand Down
4 changes: 2 additions & 2 deletions flagcx/adaptor/ccl/mccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,9 +71,9 @@ flagcxResult_t mcclAdaptorCommUserRank(const flagcxInnerComm_t comm,
}

flagcxResult_t mcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return (flagcxResult_t)mcclCommGetAsyncError(comm->base,
(mcclResult_t *)&asyncError);
(mcclResult_t *)asyncError);
}

// TODO: unsupported
Expand Down
2 changes: 1 addition & 1 deletion flagcx/adaptor/ccl/mpi_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,7 +132,7 @@ flagcxResult_t mpiAdaptorCommUserRank(const flagcxInnerComm_t comm, int *rank) {
}

flagcxResult_t mpiAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return flagcxNotSupported;
}

Expand Down
4 changes: 2 additions & 2 deletions flagcx/adaptor/ccl/musa_mccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,9 +67,9 @@ flagcxResult_t mcclAdaptorCommUserRank(const flagcxInnerComm_t comm,
}

flagcxResult_t mcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return (flagcxResult_t)mcclCommGetAsyncError(comm->base,
(mcclResult_t *)&asyncError);
(mcclResult_t *)asyncError);
}

// TODO: unsupported
Expand Down
4 changes: 2 additions & 2 deletions flagcx/adaptor/ccl/nccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,9 +66,9 @@ flagcxResult_t ncclAdaptorCommUserRank(const flagcxInnerComm_t comm,
}

flagcxResult_t ncclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return (flagcxResult_t)ncclCommGetAsyncError(comm->base,
(ncclResult_t *)&asyncError);
(ncclResult_t *)asyncError);
}

flagcxResult_t ncclAdaptorMemAlloc(void **ptr, size_t size) {
Expand Down
4 changes: 2 additions & 2 deletions flagcx/adaptor/ccl/rccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,9 +66,9 @@ flagcxResult_t rcclAdaptorCommUserRank(const flagcxInnerComm_t comm,
}

flagcxResult_t rcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return (flagcxResult_t)ncclCommGetAsyncError(comm->base,
(ncclResult_t *)&asyncError);
(ncclResult_t *)asyncError);
}

// TODO: unsupported
Expand Down
2 changes: 1 addition & 1 deletion flagcx/adaptor/ccl/xccl_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -131,7 +131,7 @@ flagcxResult_t xcclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }

// Unsupported
flagcxResult_t xcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
flagcxResult_t asyncError) {
flagcxResult_t *asyncError) {
return flagcxNotSupported;
}

Expand Down
12 changes: 6 additions & 6 deletions flagcx/adaptor/device/cuda_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,14 +45,14 @@ flagcxResult_t cudaAdaptorDeviceMalloc(void **ptr, size_t size,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaHostAlloc(ptr, size, cudaHostAllocMapped));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
} else {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
}
Comment thread
MC952-arch marked this conversation as resolved.
return flagcxSuccess;
}
Expand All @@ -61,14 +61,14 @@ flagcxResult_t cudaAdaptorDeviceFree(void *ptr, flagcxMemType_t type,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaFreeHost(ptr));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
} else {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
}
Comment on lines +66 to 72

@gemini-code-assist gemini-code-assist Bot Oct 2, 2025

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

medium

For improved clarity and robustness, it's better to explicitly check for flagcxMemDevice rather than relying on an else block. This makes the code easier to understand and safer if new flagcxMemType_t values are added in the future.

  } else if (type == flagcxMemDevice) {
    if (stream == NULL) {
      DEVCHECK(cudaFree(ptr));
    } else {
      DEVCHECK(cudaFreeAsync(ptr, stream->base));
    }
  }

return flagcxSuccess;
}
Expand Down
12 changes: 6 additions & 6 deletions flagcx/adaptor/device/ducuda_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,14 +45,14 @@ flagcxResult_t ducudaAdaptorDeviceMalloc(void **ptr, size_t size,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaMallocHost(ptr, size));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
} else {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
}
Comment thread
MC952-arch marked this conversation as resolved.
return flagcxSuccess;
}
Expand All @@ -61,14 +61,14 @@ flagcxResult_t ducudaAdaptorDeviceFree(void *ptr, flagcxMemType_t type,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaFreeHost(ptr));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
} else {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
}
Comment on lines +66 to 72

@gemini-code-assist gemini-code-assist Bot Oct 2, 2025

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

medium

For improved clarity and robustness, it's better to explicitly check for flagcxMemDevice rather than relying on an else block. This makes the code easier to understand and safer if new flagcxMemType_t values are added in the future.

Suggested change
} else {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
}
} else if (type == flagcxMemDevice) {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
}

return flagcxSuccess;
}
Expand Down
12 changes: 6 additions & 6 deletions flagcx/adaptor/device/hip_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,14 +45,14 @@ flagcxResult_t hipAdaptorDeviceMalloc(void **ptr, size_t size,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(hipHostMalloc(ptr, size));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(hipMallocManaged(ptr, size, hipMemAttachGlobal));
} else {
if (stream == NULL) {
DEVCHECK(hipMalloc(ptr, size));
} else {
DEVCHECK(hipMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(hipMallocManaged(ptr, size, hipMemAttachGlobal));
}
Comment thread
MC952-arch marked this conversation as resolved.
return flagcxSuccess;
}
Expand All @@ -61,14 +61,14 @@ flagcxResult_t hipAdaptorDeviceFree(void *ptr, flagcxMemType_t type,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(hipFreeHost(ptr));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(hipFree(ptr));
} else {
if (stream == NULL) {
DEVCHECK(hipFree(ptr));
} else {
DEVCHECK(hipFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(hipFree(ptr));
}
Comment on lines +66 to 72

@gemini-code-assist gemini-code-assist Bot Oct 2, 2025

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

medium

For improved clarity and robustness, it's better to explicitly check for flagcxMemDevice rather than relying on an else block. This makes the code easier to understand and safer if new flagcxMemType_t values are added in the future.

Suggested change
} else {
if (stream == NULL) {
DEVCHECK(hipFree(ptr));
} else {
DEVCHECK(hipFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(hipFree(ptr));
}
} else if (type == flagcxMemDevice) {
if (stream == NULL) {
DEVCHECK(hipFree(ptr));
} else {
DEVCHECK(hipFreeAsync(ptr, stream->base));
}
}

return flagcxSuccess;
}
Expand Down
12 changes: 6 additions & 6 deletions flagcx/adaptor/device/ixcuda_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,14 +45,14 @@ flagcxResult_t ixcudaAdaptorDeviceMalloc(void **ptr, size_t size,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaMallocHost(ptr, size));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
} else {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
}
Comment on lines +50 to 56

@gemini-code-assist gemini-code-assist Bot Oct 2, 2025

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

medium

For improved clarity and robustness, it's better to explicitly check for flagcxMemDevice rather than relying on an else block. This makes the code easier to understand and safer if new flagcxMemType_t values are added in the future.

Suggested change
} else {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
}
} else if (type == flagcxMemDevice) {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
}

return flagcxSuccess;
}
Expand All @@ -61,14 +61,14 @@ flagcxResult_t ixcudaAdaptorDeviceFree(void *ptr, flagcxMemType_t type,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaFreeHost(ptr));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
} else {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
}
Comment on lines +66 to 72

@gemini-code-assist gemini-code-assist Bot Oct 2, 2025

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

medium

For improved clarity and robustness, it's better to explicitly check for flagcxMemDevice rather than relying on an else block. This makes the code easier to understand and safer if new flagcxMemType_t values are added in the future.

Suggested change
} else {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
}
} else if (type == flagcxMemDevice) {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
}

return flagcxSuccess;
}
Expand Down
12 changes: 6 additions & 6 deletions flagcx/adaptor/device/kunlunxin_adaptor.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,16 +47,16 @@ flagcxResult_t kunlunAdaptorDeviceMalloc(void **ptr, size_t size,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaMallocHost(ptr, size));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
} else {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
// The underlying interface here is synchronous, not an asynchronous
// implementation.
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
}
Comment on lines +52 to 60

@gemini-code-assist gemini-code-assist Bot Oct 2, 2025

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

medium

For improved clarity and robustness, it's better to explicitly check for flagcxMemDevice rather than relying on an else block. This makes the code easier to understand and safer if new flagcxMemType_t values are added in the future.

Suggested change
} else {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
// The underlying interface here is synchronous, not an asynchronous
// implementation.
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaMallocManaged(ptr, size, cudaMemAttachGlobal));
}
} else if (type == flagcxMemDevice) {
if (stream == NULL) {
DEVCHECK(cudaMalloc(ptr, size));
} else {
// The underlying interface here is synchronous, not an asynchronous
// implementation.
DEVCHECK(cudaMallocAsync(ptr, size, stream->base));
}
}

return flagcxSuccess;
}
Expand All @@ -65,16 +65,16 @@ flagcxResult_t kunlunAdaptorDeviceFree(void *ptr, flagcxMemType_t type,
flagcxStream_t stream) {
if (type == flagcxMemHost) {
DEVCHECK(cudaFreeHost(ptr));
} else if (type == flagcxMemDevice) {
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
} else {
if (stream == NULL) {
DEVCHECK(cudaFree(ptr));
} else {
// The underlying interface here is synchronous, not an asynchronous
// implementation.
DEVCHECK(cudaFreeAsync(ptr, stream->base));
}
} else if (type == flagcxMemManaged) {
DEVCHECK(cudaFree(ptr));
}
Comment thread
MC952-arch marked this conversation as resolved.
return flagcxSuccess;
}
Expand Down
Loading