GH-47677: [C++][GPU] Allow building with CUDA 13 by pitrou · Pull Request #48259 · apache/arrow

pitrou · 2025-11-26T08:48:51Z

What changes are included in this PR?

Add compatibility fix for CUDA 13 in C++ CUDA tests
Add CI builds with CUDA 13
Disable Numba interop tests because of a regression in Numba-CUDA: see Revert #536 "perf: remove context threading in various pointer abstractions" NVIDIA/numba-cuda#611

Are these changes tested?

Yes, on CI.

Are there any user-facing changes?

No.

GitHub Issue: [C++][GPU] Support for CUDA 13? #47677

pitrou · 2025-11-26T08:49:02Z

@github-actions crossbow submit cuda

github-actions · 2025-11-26T08:49:15Z

⚠️ GitHub issue #47677 has been automatically assigned in GitHub to PR creator.

github-actions · 2025-11-26T08:51:19Z

Revision: 60f974c

Submitted crossbow builds: ursacomputing/crossbow @ actions-c2a5e1506b

Task	Status
test-cuda-cpp-ubuntu-22.04-cuda-11.7.1
test-cuda-python-ubuntu-22.04-cuda-11.7.1

pitrou · 2025-11-26T09:24:32Z

@github-actions crossbow submit cuda

github-actions · 2025-11-26T09:26:47Z

Revision: 42610fe

Submitted crossbow builds: ursacomputing/crossbow @ actions-719d93fe63

Task	Status
test-cuda-cpp-ubuntu-22.04-cuda-11.7.1
test-cuda-cpp-ubuntu-24.04-cuda-13.0.2
test-cuda-python-ubuntu-22.04-cuda-11.7.1
test-cuda-python-ubuntu-24.04-cuda-13.0.2

pitrou · 2025-11-26T09:58:06Z

So, the Numba interop tests fail with:

arrow-dev/lib/python3.12/site-packages/pyarrow/tests/test_cuda_numba_interop.py:233: 
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 

>   ???
E   TypeError: MemoryPointer.__init__() got multiple values for argument 'pointer'

Did Numba change its CUDA interaction APIs again? @gmarkall

pitrou · 2025-11-26T09:59:05Z

This is presumably happening in

arrow/python/pyarrow/_cuda.pyx

Lines 458 to 465 in 39c865e

    
               def to_numba(self): 
        
                   """Return numba memory pointer of CudaBuffer instance. 
        
                   """ 
        
                   import ctypes 
        
                   from numba.cuda.cudadrv.driver import MemoryPointer 
        
                   return MemoryPointer(self.context.to_numba(), 
        
                                        pointer=ctypes.c_void_p(self.address), 
        
                                        size=self.size)

pitrou · 2025-11-26T10:10:48Z

@github-actions crossbow submit cudapython*

github-actions · 2025-11-26T10:13:08Z

Revision: ed01642

Submitted crossbow builds: ursacomputing/crossbow @ actions-bc22ec352a

Task	Status
test-cuda-python-ubuntu-22.04-cuda-11.7.1
test-cuda-python-ubuntu-24.04-cuda-13.0.2

pitrou · 2025-11-26T10:42:53Z

Ok, given that the tests now fail for other reasons, I think we might just disable them on CI and let interested parties contribute.

pitrou · 2025-11-26T13:50:19Z

@github-actions crossbow submit cudapython*

github-actions · 2025-11-26T13:52:29Z

Revision: 5080baf

Submitted crossbow builds: ursacomputing/crossbow @ actions-d5735696f4

Task	Status
test-cuda-python-ubuntu-22.04-cuda-11.7.1
test-cuda-python-ubuntu-24.04-cuda-13.0.2

pitrou · 2025-11-26T13:56:48Z

cpp/src/arrow/gpu/cuda_test.cc

+#if CUDA_VERSION >= 13000
+    RETURN_NOT_OK(StatusFromCuda(cuCtxCreate_v4(&ctx, /*ctxCreateParams=*/nullptr,
+                                                /*flags=*/0, device_->handle())));
+#else
    RETURN_NOT_OK(StatusFromCuda(cuCtxCreate(&ctx, /*flags=*/0, device_->handle())));
+#endif


@gmarkall Does this fix look ok?

pitrou · 2025-11-26T14:14:46Z

@github-actions crossbow submit cudapython*

github-actions · 2025-11-26T14:17:01Z

Revision: 9423f81

Submitted crossbow builds: ursacomputing/crossbow @ actions-c59a2b88bb

Task	Status
test-cuda-python-ubuntu-22.04-cuda-11.7.1
test-cuda-python-ubuntu-24.04-cuda-13.0.2

gmarkall · 2025-11-26T14:44:45Z

cpp/src/arrow/gpu/cuda_test.cc

  Result<CUcontext> NonPrimaryRawContext() {
    CUcontext ctx;
+#if CUDA_VERSION >= 13000
+    RETURN_NOT_OK(StatusFromCuda(cuCtxCreate_v4(&ctx, /*ctxCreateParams=*/nullptr,


I'm not sure cuCtxCreate_v4 is designated as a public API. For CUDA 13, cuCtxCreate is #defined as cuCtxCreate_v4, so I think I'd be inclined to write this as:

Suggested change

RETURN_NOT_OK(StatusFromCuda(cuCtxCreate_v4(&ctx, /*ctxCreateParams=*/nullptr,

RETURN_NOT_OK(StatusFromCuda(cuCtxCreate(&ctx, /*ctxCreateParams=*/nullptr,

It seems it is public:
https://docs.nvidia.com/cuda/archive/12.6.3/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX_1gd84cbb0ad9470d66dc55e0830d56ef4d

(unrelated, but using a #define isn't very nice for non-C languages that would want to access the driver API through FFI)

That is from the 12.6 API and this was removed in the 13.0 API: https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__CTX.html#group__CUDA__CTX (but not the ABI).

Well, weirdly it succeeds on the 13.0.2 CUDA build... But, yes, I can just switch to calling the macro.

The way it's defined in the cuda.h header is that cuCtxCreate is a macro defined as cuCtxCreate_v4:

#define cuCtxCreate cuCtxCreate_v4 ... CUresult CUDAAPI cuCtxCreate(CUcontext *pctx, CUctxCreateParams *ctxCreateParams, unsigned int flags, CUdevice dev);

This was an API breaking change in CUDA 13. If you want to generically support CUDA 12 and CUDA 13, I'd recommend using the cuGetProcAddress API (https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__DRIVER__ENTRY__POINT.html#group__CUDA__DRIVER__ENTRY__POINT_1gcae5adad00590572ab35b2508c2d6e0d) which is similar to using dlsym and allows you to work against the ABI which is guaranteed to never be broken.

pitrou · 2025-11-26T14:47:49Z

@github-actions crossbow submit cudacpp*

github-actions · 2025-11-26T15:31:46Z

Revision: 9ea7b75

Submitted crossbow builds: ursacomputing/crossbow @ actions-03152c375f

Task	Status
test-cuda-cpp-ubuntu-22.04-cuda-11.7.1
test-cuda-cpp-ubuntu-24.04-cuda-13.0.2

gmarkall · 2025-11-26T15:58:09Z

Did Numba change its CUDA interaction APIs again? @gmarkall

That was a change made in NVIDIA/numba-cuda#536 and it did change a public API. I had failed to foresee that there would be an effect outside of Numba-CUDA. This change has only just made it into a published release in the last few days, so let me assess what the best way forward is and post an update here.

PR to restore the public-facing surface of these APIs to what it was before is here: NVIDIA/numba-cuda#610 - my intention is that if this is merged, to make another release shortly afterwards so that the latest version is consistent with earlier versions.

NVIDIA#536)" This reverts commit 9a56516. This changed the public API of `MemoryPointer` and related classes, and the context that they held was used by Arrow (see apache/arrow#48259 (comment)): > Numba interop tests fail with: ``` arrow-dev/lib/python3.12/site-packages/pyarrow/tests/test_cuda_numba_interop.py:233: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ > ??? E TypeError: MemoryPointer.__init__() got multiple values for argument 'pointer' ``` This commit reverts the change, as it was intended to improve performance without changing functionality, but has had a functional change as a side effect. Following the merge of this PR, we should be able to remove some of the `@require_context` decorators with some more targeted changes.

pitrou · 2025-11-27T12:30:35Z

@raulcd @jorisvandenbossche What do you think? Should we wait for Numba-Cuda to publish a new release or should we just merge this PR with a skip as currently done?

raulcd · 2025-11-27T12:53:21Z

Should we wait for Numba-Cuda to publish a new release or should we just merge this PR with a skip as currently done?

Given that we are going to lose access to the cuda runners for an unspecified amount of time starting this weekend I'd rather merge with the skip and create a follow up issue to remove the skip once we have runners again and can validate.

pitrou · 2025-11-27T13:55:35Z

@raulcd Do you want to give this a review?

pitrou · 2025-11-27T13:57:33Z

@github-actions crossbow submit cuda

github-actions · 2025-11-27T13:59:53Z

Revision: 00fccbb

Submitted crossbow builds: ursacomputing/crossbow @ actions-23548261b9

Task	Status
test-cuda-cpp-ubuntu-22.04-cuda-11.7.1
test-cuda-cpp-ubuntu-24.04-cuda-13.0.2
test-cuda-python-ubuntu-22.04-cuda-11.7.1
test-cuda-python-ubuntu-24.04-cuda-13.0.2

gmarkall · 2025-11-27T14:04:13Z

@raulcd @jorisvandenbossche What do you think? Should we wait for Numba-Cuda to publish a new release or should we just merge this PR with a skip as currently done?

FWIW, I'm hoping to resolve the Numba-CUDA issue and publish a new release today.

pitrou · 2025-11-27T14:12:46Z

FWIW, I'm hoping to resolve the Numba-CUDA issue and publish a new release today.

Actually testing Numba interop may still require GH-47371 to be fixed, though.

gmarkall · 2025-11-27T15:58:13Z

Thanks for the pointer - I'll move onto that issue shortly afterwards.

python/pyarrow/_cuda.pyx

…ctions" (#611) This reverts commit 9a56516. This changed the public API of `MemoryPointer` and related classes, and the context that they held was used by Arrow (see apache/arrow#48259 (comment)): > Numba interop tests fail with: ``` arrow-dev/lib/python3.12/site-packages/pyarrow/tests/test_cuda_numba_interop.py:233: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ > ??? E TypeError: MemoryPointer.__init__() got multiple values for argument 'pointer' ``` This commit reverts the change, as it was intended to improve performance without changing functionality, but has had a functional change as a side effect. Following the merge of this PR, we should be able to remove some of the `@require_context` decorators with some more targeted changes.

raulcd

Thanks @pitrou .

Thanks @gmarkall for the pointer to the API change.

conbench-apache-arrow · 2025-11-28T16:52:30Z

After merging your PR, Conbench analyzed the 4 benchmarking runs that have been run so far on merge-commit ab4a096.

There were no benchmark performance regressions. 🎉

The full Conbench report has more details.

github-actions bot added Component: C++ awaiting review Awaiting review labels Nov 26, 2025

pitrou force-pushed the gh47677-cuda13 branch from 60f974c to 42610fe Compare November 26, 2025 09:24

github-actions bot added the Component: Python label Nov 26, 2025

pitrou mentioned this pull request Nov 26, 2025

[Python][GPU] Numba interop tests broken by Numba API changes #48265

Open

pitrou commented Nov 26, 2025

View reviewed changes

github-actions bot added awaiting committer review Awaiting committer review and removed awaiting review Awaiting review labels Nov 26, 2025

pitrou force-pushed the gh47677-cuda13 branch from 5080baf to 9423f81 Compare November 26, 2025 14:10

gmarkall reviewed Nov 26, 2025

View reviewed changes

pitrou marked this pull request as ready for review November 26, 2025 14:48

pitrou requested review from assignUser, jonkeane, kou and raulcd as code owners November 26, 2025 14:48

gmarkall mentioned this pull request Nov 26, 2025

Revert #536 "perf: remove context threading in various pointer abstractions" NVIDIA/numba-cuda#611

Merged

pitrou added 4 commits November 27, 2025 14:57

apacheGH-47677: [C++][GPU] Allow building with CUDA 13

971a5b8

Try workaround for newer Numba

6196bb4

Skip Numba interop tests

a8b5668

Use cuCtxCreate

00fccbb

pitrou force-pushed the gh47677-cuda13 branch from 9ea7b75 to 00fccbb Compare November 27, 2025 13:57

raulcd reviewed Nov 27, 2025

View reviewed changes

python/pyarrow/_cuda.pyx Show resolved Hide resolved

github-actions bot added awaiting changes Awaiting changes and removed awaiting committer review Awaiting committer review labels Nov 27, 2025

raulcd approved these changes Nov 28, 2025

View reviewed changes

github-actions bot added awaiting merge Awaiting merge and removed awaiting changes Awaiting changes labels Nov 28, 2025

raulcd mentioned this pull request Nov 28, 2025

[Python][GPU] Remove cuda numba interop test skip once cuda-numba is released #48281

Closed

raulcd merged commit ab4a096 into apache:main Nov 28, 2025
49 checks passed

raulcd removed the awaiting merge Awaiting merge label Nov 28, 2025

raulcd mentioned this pull request Nov 28, 2025

[C++][GPU] Support for CUDA 13? #47677

Closed

pitrou deleted the gh47677-cuda13 branch November 28, 2025 17:13

	RETURN_NOT_OK(StatusFromCuda(cuCtxCreate_v4(&ctx, /ctxCreateParams=/nullptr,
	RETURN_NOT_OK(StatusFromCuda(cuCtxCreate(&ctx, /ctxCreateParams=/nullptr,

Conversation

pitrou commented Nov 26, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What changes are included in this PR?

Are these changes tested?

Are there any user-facing changes?

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

pitrou Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

gmarkall Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

pitrou Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

pitrou Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

kkraus14 Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

pitrou Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

kkraus14 Nov 26, 2025

Choose a reason for hiding this comment

Uh oh!

pitrou commented Nov 26, 2025

Uh oh!

github-actions bot commented Nov 26, 2025

Uh oh!

gmarkall commented Nov 26, 2025

Uh oh!

pitrou commented Nov 27, 2025

Uh oh!

raulcd commented Nov 27, 2025

Uh oh!

pitrou commented Nov 27, 2025

Uh oh!

pitrou commented Nov 27, 2025

Uh oh!

github-actions bot commented Nov 27, 2025

Uh oh!

gmarkall commented Nov 27, 2025

Uh oh!

pitrou commented Nov 27, 2025

Uh oh!

gmarkall commented Nov 27, 2025

Uh oh!

Uh oh!

raulcd left a comment

pitrou commented Nov 26, 2025 •

edited

Loading