Fix deform_conv2d kernels to use current CUDA stream

[Nueramarcos]

Fixes #9513

Problem

All 6 CUDA kernel launches in deform_conv2d_kernel.cu used <<<blocks, threads>>> with no stream argument, meaning they always ran on the default CUDA stream. This causes silent race conditions when users run deformable convolutions inside a non-default stream (e.g. with torch.cuda.stream(s) or in a multi-stream pipeline).

The same bug affects both the int and int64_t index variants of all three kernels:

• deformable_im2col_kernel (forward)
• deformable_col2im_kernel (backward grad input)
• deformable_col2im_coord_kernel (backward grad offset/mask)

Fix

Add const cudaStream_t stream = at::cuda::getCurrentCUDAStream(); once before each if (use_64bits_indexing) block, then pass stream as the fourth chevron argument to all 6 launches:

// Before
deformable_im2col_kernel<scalar_t, int64_t><<<blocks, threads>>>(...)

// After
const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
deformable_im2col_kernel<scalar_t, int64_t><<<blocks, threads, 0, stream>>>(...)

This matches the existing pattern in roi_pool_kernel.cu and other ops in this repo.

Testing

Stream sanity check (run locally with CUDA):

import torch
from torchvision.ops import deform_conv2d

x = torch.randn(1, 3, 10, 10, device="cuda")
weight = torch.randn(3, 3, 3, 3, device="cuda")
offset = torch.randn(1, 18, 10, 10, device="cuda")

# Verify op works correctly on a non-default stream
s = torch.cuda.Stream()
with torch.cuda.stream(s):
    out = deform_conv2d(x, weight, offset)
s.synchronize()
print("OK")

Full test suite: pytest test/test_ops.py::TestDeformConv2d -xvs

[meta-cla]

Hi @Nueramarcos!

Thank you for your pull request and welcome to our community.

Action Required

In order to merge any pull request (code, docs, etc.), we require contributors to sign our Contributor License Agreement, and we don't seem to have one on file for you.

Process

In order for us to review and merge your suggested changes, please sign at https://code.facebook.com/cla. If you are contributing on behalf of someone else (eg your employer), the individual CLA may not be sufficient and your employer may need to sign the corporate CLA.

Once the CLA is signed, our tooling will perform checks and validations. Afterwards, the pull request will be tagged with CLA signed. The tagging process may take up to 1 hour after signing. Please give it that time before contacting us about it.

If you have received this in error or have any questions, please contact us at cla@meta.com. Thanks!

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/vision/9515

• 📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

❗ 1 Active SEVs

There are 1 currently active SEVs. If your PR is affected, please view them below:

• [ROCm] MI350 CI jobs will have longer queue times due to CI migration

This comment was automatically generated by Dr. CI and updates every 15 minutes.