Fix cuPointerGetAttribute(s) returning 0 for CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL

[fnrizzi]

cuPointerGetAttribute and cuPointerGetAttributes always return 0 for CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL, regardless of which device the pointer was actually allocated on.
The root cause is I think a write/read field mismatch in
_HelperCUpointer_attribute (cuda_bindings/cuda/bindings/_lib/utils.pxi):

• __cinit__ allocates storage for DEVICE_ORDINAL in the signed _int
  member and points _cptr at it, so the driver writes the ordinal there.
• pyObj() reads DEVICE_ORDINAL back from the unsigned _uint member,
  which is never written and stays 0.
  Because the value is read from the wrong member, the ordinal is always 0.

Fix

Read DEVICE_ORDINAL from _int in pyObj() so the read path matches the write path in __cinit__.

Test

Adds a regression test that allocates on each visible GPU and asserts the
ordinal reported by both the singular and plural bindings matches the
allocation device (ground truth from cudaPointerGetAttributes).
Existing tests missed this because they only asserted call success / that
singular and plural agree (both share the buggy read path and both returned 0).
The test needs >= 2 GPUs and skips otherwise, since on device 0 a wrong 0 is
indistinguishable from the correct value.

Below are also minimal reproducers in c and python . the c one works correctly , the python does not and motivated this MR.

/*
 * nvcc -o minimal minimal.c -lcuda
 * ./minimal
 */

#include <cuda.h>
#include <cuda_runtime_api.h>
#include <stdio.h>

int main(void) {
    cuInit(0);

    int ndev = 0;
    cudaGetDeviceCount(&ndev);
    printf("visible GPUs: %d\n", ndev);

    for (int dev = 0; dev < ndev; ++dev) {
        cudaSetDevice(dev);

        void *ptr = NULL;
        cudaMalloc(&ptr, 256);

        /* runtime */
        struct cudaPointerAttributes attr;
        cudaPointerGetAttributes(&attr, ptr);
        int rt = attr.device;

        /* driver, singular */
        int singular = -1;
        cuPointerGetAttribute(&singular, CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL,
                              (CUdeviceptr)ptr);

        /* driver, plural */
        int plural = -1;
        CUpointer_attribute attrs[1] = {CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL};
        void *data[1] = {&plural};
        cuPointerGetAttributes(1, attrs, data, (CUdeviceptr)ptr);

        printf("device %d: rt=%d singular=%d plural=%d\n", dev, rt, singular, plural);

        cudaFree(ptr);
    }
    return 0;
}

from cuda.bindings import runtime as cbr
from cuda.bindings import driver as cbd

ORD = cbd.CUpointer_attribute.CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL
_, ndev = cbr.cudaGetDeviceCount()
for dev in range(ndev):
    cbr.cudaSetDevice(dev)
    _, ptr = cbr.cudaMalloc(256)
    rt = cbr.cudaPointerGetAttributes(ptr)[1].device
    singular = cbd.cuPointerGetAttribute(ORD, cbd.CUdeviceptr(ptr))[1]
    plural = cbd.cuPointerGetAttributes(1, [ORD], cbd.CUdeviceptr(ptr))[1][0]
    print(f"device {dev}: rt={rt} singular={singular} plural={plural}")
    cbr.cudaFree(ptr)

when i run these, before the fix of this MR, I get:

bin/nvcc -o exe minimal.c -lcuda && ./exe
visible GPUs: 4
device 0: rt=0 singular=0 plural=0
device 1: rt=1 singular=1 plural=1
device 2: rt=2 singular=2 plural=2
device 3: rt=3 singular=3 plural=3

and

python minimal.py
device 0: rt=0 singular=0 plural=0
device 1: rt=1 singular=0 plural=0
device 2: rt=2 singular=0 plural=0
device 3: rt=3 singular=0 plural=0

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.