ENH: (almost) full structured dtype support by seberg · Pull Request #9650 · cupy/cupy

seberg · 2026-01-28T11:36:33Z

~~This is quite far along, but needs cleaning up (and is based on the other PR). So it is a WIP only, just for visibility.~~

Based on gh-9494, this adds the missing support for structured dtypes, that is, allowing casts and multi-field access/assignments.

Major open points:

We need to thread through the dtype much more. This is OK, but we need to refactor things to make use of the NumPy C-API to reduce overheads again.
- It may make sense to do a bit larger refactors around this. EDIT: Actually, I don't think these changes cause any measurable speed regression, some things may be nice but not much needed here.
NumPy doesn't help us align types properly, so we'll need a helper to do so.
- I am curious if we couldn't query the alignof() through the JIT to implement this? (plus memoization).
- This will currently lead to some incorrect paths even for correct inputs in theory (i.e. the comparison compares via a promoted type and that promoted type may be unaligned)
~~It needs tests of course!~~

Smaller points:

I would like this to be interoperable with a plain struct, if such a plain struct can work. That logic exists, but likely doesn't quite work yet.
This PR does not implement subarray support.

To implement this I added a resolve_dtypes() function. This is a concept that NumPy also has, although I made it a bit more raw here. (The divergence has two reasons. First casts as a ufunc don't work well without it, although a cast could maybe be a simple ElementwiseKernel. But also the functionality used in NumPy to actual implement it nicely isn't public, oops.)

Addresses gh-2031 (but one could argue that subarray support is missing)

leofang · 2026-02-03T19:46:22Z

+
+// Hack to replace std::false_type and std::true_type
+template<bool Value>
+struct my_integral_constant {


FYI std::integral_constant is available through cupy/workaround.h

leofang · 2026-02-03T20:12:16Z

    buf.format = dtype_format[dtype]
+
+
+def make_gpu_aligned_dtype(


Two questions:

How does this relate to GPU, other than that we ensure alignment? I assume make_aligned_dtype is fine too?

How about moving this to cuda.core? We need this function for libraries lower than CuPy in the stack (Add cuda.core.utils.get_aligned_dtype NVIDIA/cuda-python#734).

Moving sounds good. CPU alignment is smaller than GPU one often. While GPU is (typically?) itemsize aligned. An example that always works is complex numbers, which CUDA uses the full itemsize alignment for copies, I think, while for the CPU the compiler will only require the floats alignment.

FWIW, I am not sure that encoding a custom larger alignment is super useful. It seems mainly needed for nested structs which dould maybe be solve directly by passing field alignment constraints offset (that are only used to compute offsets).

Yeah for 1 I mainly wanted to clarify for the function name. It seems calling it make_aligned_dtype() (removing "gpu") is fine. There is nothing GPU-specific.

For 2, my suggestion is

move it to cuda.core (I can take over this if you like, or if you have time to check it there it's also fine 🙂 )

try-import it in CuPy, if cuda.core is new enough to offer it, use it

otherwise, use a private copy vendored in CuPy (and we can remove it in 1-2 major releases)

Sorry yeah, I'll work on moving it. Thought I may have an idea about how to signal alignment, but likely this is pretty good (and we only embed if users asked for somdthing special).

(I suppose we cannot move the C++ type construction, might befun one day, but not a priority.)

seberg · 2026-02-21T18:09:02Z

Ok, this grew to 1k lines, so maybe should split it up. But while it still needs some polishing, I think this should be good for review.

The basic changes are:

Introduce a header for field based struct view (and one for unstructured, we don't have std::array, I think)
Create a helper cupy.make_aligned_dtype, we need this internally as well.
Create Python code to build the C type from the structured dtype (i.e. get_typename()), this is fairly complicated!
- Currently, it builds a struct if possible. This may serve no real purpose, so maybe we should just remove it (and maybe add field names via .at<"name"> if that works.)
We need a resolve_dtypes function, this is needed to figure out the right loop dtypes for comparisons.
CScalars may have arbitrary size now and in a few places we needed to propagate them better.
get_typename(..., type_headers) is now type_decls (changed everywhere). Dependencies matter now, so it needs a topological sort. Feels a bit overpowered, but also kinda OK...

I am sure there will be places with some holes (if just due to an .item()), and e.g. we won't notice if there should be a ComplexWarning due to a nested cast, but that seems fine...

The main questions here may well be around __cuda_alignment__ or so.

mergify · 2026-02-28T17:56:19Z

This pull request is now in conflicts. Could you fix it @seberg? 🙏

seberg · 2026-03-03T13:07:14Z

/test mini

leofang

I am half-way done!

leofang · 2026-03-07T18:25:38Z

Sign in to view

        # is the case, this always passes.
        # with pytest.raises(ValueError):
        #     func(ByteSwapped(base))


Q: Is this comment invalid now?

Yeah, I hadn't fixed that yet here. Let's see, it might not be much more than a 2 line change to make it work well enough.

OK, I added support for this, but with slightly extra logic to work with the necessary padding more reliably (NumPy doesn't do this well, but also doesn't pad by default, I need to think about what API to add in NumPy to improve this).

The asarray() path still doesn't raise, because CAI doesn't strictly check the dtype compatibility right now. (tried enabling it, but I forgot about datetime)

leofang · 2026-03-07T18:27:02Z

+class TestFieldCasting:
+    @testing.numpy_cupy_array_equal()
+    def test_casting_simple(self, xp):
+        a = xp.array([(1, 2.), (3, 4.)], dtype="i8,f8")


Q: if we use "i4,f8" instead, will we get to test the padding behavior?

I am not auto-padding for xp.array() right now, so it would just fail on this test. There are some tests below (mainly using complex types, because they are always float aligned in NumPy/C and itemsize aligned in CUDA) that test the padding paths.

Ah. What if it's np.dtype([('f1', np.int32), ('f2', np.float64)], align=True) instead? Would it also fail?

In principle yes, float64 CPU alignment can be 4 bytes. That said, it may be that it isn't on any hardware we support in practice (I am not sure it is on any 64bit platforms).

leofang · 2026-03-07T18:28:44Z

+        # When the data can be defined as a normal struct, then
+        # we allow access via `.data`.
+        # TODO/NOTE(seberg): Is this even desired?!
+        kernel = cupy.ElementwiseKernel(
+            'T in', 'T out', 'out.data.f1 = -in.data.f0;',


Q: How do we determine the field names on the device side?

They come directly from the dtype field names, if the dtype can be represented by a C-struct, then I generate the typedef struct {} and use it for .data to make them available.
We could allow something like out.get<field> but it seems C++20 or so and we don't actually need field names for our kernels, it is only interesting for custom kernels.

Commenting also because we should decide whether we want this at all. We could just delete the .data attribute (as we don't use it).
So this is a convenience feature for custom kernels avoiding careful dtype checking + reinterpret casts in the kernel.

We could just delete the .data attribute (as we don't use it).

So instead of

'out.data.f1 = -in.data.f0;'

we write

'out.f1 = -in.f0;'

? If so, it sounds reasonable to me.

This C++ feature seems very internal-ish and I expect only the team (+ contributors) would be using it within the codebase for now, so I am not too worried about it. I am aware some users reaching into CArray stuff, but it's really undocumented and offers no guarantee.

Ah, right now it is out.data.f1, I could maybe subclass the structured dtype to achieve writing out.f1. I.e. .data is the normal struct. I had a subclassing solution initially, but then thought if we don't need it much, subclassing may add a bit much complexity.

The alternative would be out.get<f1>() = -in.get<f0>() if C++20 is available.

would be using it within the codebase for now, so I am not too worried about it.

So maybe it is just a bit "experimental" which seems fine. But not sure if we are likely to need it internally (maybe in cupyx).

yeah let's keep it c++17 compatible for now

leofang · 2026-03-13T05:05:55Z

+    return alignment
+
+
+def _build_struct_typename(dtype, type_decls):


Q: Should we use functool.cache or the internal _utils.cache here?

Unfortunately, caching needs to use the identity, so I didn't add it yet. But yeah, maybe I should (it needs to be a custom caching, unfortunately, which may mean cache misses are pretty common), since this is going to be slowish.

leofang · 2026-03-13T22:49:14Z

(Let me merge with main to catch up with recent CI cache update)

leofang · 2026-03-13T22:49:28Z

/test mini

leofang · 2026-03-14T20:49:12Z

/test mini

mergify · 2026-05-07T04:30:59Z

This pull request is now in conflicts. Could you fix it @seberg? 🙏

…fields)

The full thing is too suggestive IMO. E.g. recurse=... would have to actually recurse in that case, but currently doesn't.

seberg · 2026-05-12T14:23:56Z

Erik had given me some claude red-team tasks. I applied most of them (the largest diff being that it suggested to re-organize the header to allow more than 250 fields; not really vital but sounds good).
There were also two issues with the "expose .data as an actual C-struct" with field names part. But since I removed that, those went away.

A few of the suggestion seemed wrong, or way to niche to do something about (biggest niche one is support for dtype([]), I also leaned to not early reject "i4,>i4" although also happy to do -- it's super niche either way).

/test mini

eriknw

LGTM, primarily from a user perspective.

Thanks for applying the appropriate updates from my Claude's suggestions. I agree with your judgement about which to ignore.

The one other comment I left is more about technical compliance with CAI v3 that I think we should be aware of (and these comments leave a paper trail for). Probably fine as is though, so I'm approving the PR 👍

seberg · 2026-05-12T20:17:15Z

This causes a full rebuild, so unfortunately needs at least one more iteration.

/test mini

seberg · 2026-05-13T10:59:09Z

(Was fine before except one test not finishing yet, so skipping for the doc change.)

/test skip

seberg · 2026-05-13T11:36:30Z

Let me put this in then, thanks for the reviews! I think make_aligned_dtype is fine, the least sure I am is about the __cuda_alignment__ metadata, but it does make nesting work well and doesn't seem like a bad choice.

The doc addition is rather trivial and renders fine but we can follow-up of course.

seberg · 2026-05-13T11:38:11Z

Ah, silly me, I can't actually merge, need someone to do a final approve as the one got stale :).

leofang

A few observations from reading through the PR.

leofang · 2026-05-13T19:58:46Z

+template<size_t Size>
+class UnstructuredVoid {
+public:
+  char data[Size];


UnstructuredVoid<Size> inherits the alignment of char (i.e. 1). That means CArray<UnstructuredVoid<8>> will also get alignment 1, so device accesses to multi-byte voids won't be naturally aligned. Consider adding alignas(...) — e.g. the largest power-of-two ≤ min(Size, 16) — to match the alignment CUDA would give an equivalent-sized trivial type.

It would be nice to use a higher alignment when possible, but in principle there is nothing stopping you from having a ?,V4 dtype or so, where the V4 is random data that has no alignment guarantees.
Maybe we'll have to either allow users to override it explicitly, or use actual runtime alignment for the kernel building (I like the latter if it works).

seberg · 2026-05-14T06:46:06Z

/test mini

leofang · 2026-05-14T13:45:49Z

/test mini

leofang · 2026-05-14T23:02:57Z

/test mini,windows

seberg · 2026-05-15T07:38:38Z

/test cuda120,windows

ENH: (almost) full structured dtype support

seberg mentioned this pull request Jan 28, 2026

ENH: Explorer a path to string (and structured dtype) support for cupy #7462

Closed

seberg force-pushed the full-structured branch 2 times, most recently from c10ee65 to cc1fa6b Compare February 3, 2026 19:19

leofang reviewed Feb 3, 2026

View reviewed changes

leofang mentioned this pull request Feb 18, 2026

WIP: Add cuda.core.utils.make_aligned_dtype NVIDIA/cuda-python#1636

Draft

2 tasks

seberg force-pushed the full-structured branch from cc1fa6b to 9d9fbd7 Compare February 20, 2026 16:51

seberg changed the title ~~WIP: (almost) full structured dtype support~~ ENH: (almost) full structured dtype support Feb 21, 2026

seberg marked this pull request as ready for review February 21, 2026 18:13

seberg requested a review from a team as a code owner February 21, 2026 18:13

seberg force-pushed the full-structured branch from cac0e08 to f81d2b4 Compare March 2, 2026 13:34

leofang assigned asi1024 and leofang Mar 4, 2026

leofang reviewed Mar 11, 2026

View reviewed changes

leofang reviewed Mar 13, 2026

View reviewed changes

seberg force-pushed the full-structured branch from 9e3d8cd to 4169667 Compare March 13, 2026 11:42

leofang added cat:feature New features/APIs to-be-backported Pull-requests to be backported to stable branch prio:high labels Mar 14, 2026

ikrommyd mentioned this pull request Apr 13, 2026

Implement ndarray.byteswap() #9860

Merged

leofang self-requested a review April 27, 2026 15:10

leofang assigned eriknw May 7, 2026

kmaehashi requested a review from asi1024 May 8, 2026 07:42

seberg added 8 commits May 12, 2026 15:07

Add a simple assignment test

7e11a51

Allow unstructured void comparison

6f6021c

fixup fill

5aac205

pre-seed alignment dict

4edf206

Remove __device__ as suggested by claude

4f45b13

Other small fixups

7612059

Use suggested rewrite of index_sequence for Field header (allow many …

c965bb2

…fields)

Fixup the "nice error" thing (not the full thing)

cdee2a2

The full thing is too suggestive IMO. E.g. recurse=... would have to actually recurse in that case, but currently doesn't.

eriknw approved these changes May 12, 2026

View reviewed changes

Comment thread cupy/_core/core.pyx Resolved

Explicitly allow "descr" not being set as unstructured void in CAI

8c2a0c5

DOC: Add make_aligned_dtype to the docs

67e7755

leofang reviewed May 13, 2026

View reviewed changes

Address some review comments (full init, use enum for flags)

bc34ef1

leofang approved these changes May 14, 2026

View reviewed changes

seberg enabled auto-merge May 15, 2026 09:14

seberg merged commit 4db704c into cupy:main May 15, 2026
54 checks passed

chainer-ci mentioned this pull request May 15, 2026

[backport] ENH: (almost) full structured dtype support #9927

Merged

seberg added a commit that referenced this pull request May 15, 2026

Merge pull request #9650 from seberg/full-structured

5c88198

ENH: (almost) full structured dtype support

seberg deleted the full-structured branch May 17, 2026 11:39

kmaehashi added this to the v15.0.0a1 milestone May 21, 2026

kmaehashi mentioned this pull request May 30, 2026

[v14] Bump version to v14.1.1 #9971

Merged

		return alignment


		def _build_struct_typename(dtype, type_decls):

Conversation

seberg commented Jan 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

seberg commented Feb 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

mergify Bot commented Feb 28, 2026

Uh oh!

seberg commented Mar 3, 2026

Uh oh!

leofang left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

This comment was marked as outdated.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

seberg Mar 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

leofang commented Mar 13, 2026

Uh oh!

leofang commented Mar 13, 2026

Uh oh!

leofang commented Mar 14, 2026

Uh oh!

mergify Bot commented May 7, 2026

Uh oh!

seberg commented May 12, 2026

Uh oh!

eriknw left a comment

Choose a reason for hiding this comment

seberg commented Jan 28, 2026 •

edited

Loading

seberg commented Feb 21, 2026 •

edited

Loading

seberg Mar 11, 2026 •

edited

Loading