Skip to content

Commit 2ba4bfc

Browse files
juenglinjuenglin
committed
fix(bindings): add channel-format kind coverage for normalized and YUV formats
1 parent 51a8f4e commit 2ba4bfc

2 files changed

Lines changed: 227 additions & 1 deletion

File tree

‎cuda_bindings/cuda/bindings/_lib/cyruntime/cyruntime.pxi‎

Lines changed: 156 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -259,6 +259,26 @@ cdef cudaError_t getDescInfo(const cudaChannelFormatDesc* d, int *numberOfChanne
259259
elif d[0].f in (cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized1010102,):
260260
if (d[0].x != 10) or (d[0].y != 10) or (d[0].z != 10) or (d[0].w != 2):
261261
return cudaErrorInvalidChannelDescriptor
262+
elif d[0].f in (cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed422,
263+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed444,):
264+
if (d[0].x != 8) or (d[0].y != 8) or (d[0].z != 8) or (d[0].w != 8):
265+
return cudaErrorInvalidChannelDescriptor
266+
elif d[0].f in (cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar420,
267+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar422,
268+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar444,
269+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar420,
270+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar422,
271+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar444,):
272+
if (d[0].x != 8) or (d[0].y != 8) or (d[0].z != 8) or (d[0].w != 0):
273+
return cudaErrorInvalidChannelDescriptor
274+
elif d[0].f in (cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar420,
275+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar422,
276+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar444,
277+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar420,
278+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar422,
279+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar444,):
280+
if (d[0].x != 16) or (d[0].y != 16) or (d[0].z != 16) or (d[0].w != 0):
281+
return cudaErrorInvalidChannelDescriptor
262282
else:
263283
return cudaErrorInvalidChannelDescriptor
264284

@@ -414,12 +434,104 @@ cdef cudaError_t getDescInfo(const cudaChannelFormatDesc* d, int *numberOfChanne
414434
elif case_desc(d, 10, 10, 10, 2, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized1010102):
415435
numberOfChannels[0] = 4
416436
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT_101010_2
437+
# 2a: normalized 8/16-bit kinds — validated above but previously had no case_desc arm
438+
elif case_desc(d, 8, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized8X1):
439+
numberOfChannels[0] = 1
440+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT8X1
441+
elif case_desc(d, 8, 8, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized8X2):
442+
numberOfChannels[0] = 2
443+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT8X2
444+
elif case_desc(d, 8, 8, 8, 8, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized8X4):
445+
numberOfChannels[0] = 4
446+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT8X4
447+
elif case_desc(d, 8, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized8X1):
448+
numberOfChannels[0] = 1
449+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_SNORM_INT8X1
450+
elif case_desc(d, 8, 8, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized8X2):
451+
numberOfChannels[0] = 2
452+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_SNORM_INT8X2
453+
elif case_desc(d, 8, 8, 8, 8, cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized8X4):
454+
numberOfChannels[0] = 4
455+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_SNORM_INT8X4
456+
elif case_desc(d, 16, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized16X1):
457+
numberOfChannels[0] = 1
458+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT16X1
459+
elif case_desc(d, 16, 16, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized16X2):
460+
numberOfChannels[0] = 2
461+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT16X2
462+
elif case_desc(d, 16, 16, 16, 16, cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized16X4):
463+
numberOfChannels[0] = 4
464+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UNORM_INT16X4
465+
elif case_desc(d, 16, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized16X1):
466+
numberOfChannels[0] = 1
467+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_SNORM_INT16X1
468+
elif case_desc(d, 16, 16, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized16X2):
469+
numberOfChannels[0] = 2
470+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_SNORM_INT16X2
471+
elif case_desc(d, 16, 16, 16, 16, cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized16X4):
472+
numberOfChannels[0] = 4
473+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_SNORM_INT16X4
474+
# 2b: CUDA 13.3/13.4 packed and multi-planar YUV kinds
475+
elif case_desc(d, 8, 8, 8, 8, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed422):
476+
numberOfChannels[0] = 4
477+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_PACKED_422
478+
elif case_desc(d, 8, 8, 8, 8, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed444):
479+
numberOfChannels[0] = 4
480+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_PACKED_444
481+
elif case_desc(d, 8, 8, 8, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar420):
482+
numberOfChannels[0] = 3
483+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_SEMIPLANAR_420
484+
elif case_desc(d, 16, 16, 16, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar420):
485+
numberOfChannels[0] = 3
486+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT16_SEMIPLANAR_420
487+
elif case_desc(d, 8, 8, 8, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar422):
488+
numberOfChannels[0] = 3
489+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_SEMIPLANAR_422
490+
elif case_desc(d, 16, 16, 16, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar422):
491+
numberOfChannels[0] = 3
492+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT16_SEMIPLANAR_422
493+
elif case_desc(d, 8, 8, 8, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar444):
494+
numberOfChannels[0] = 3
495+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_SEMIPLANAR_444
496+
elif case_desc(d, 16, 16, 16, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar444):
497+
numberOfChannels[0] = 3
498+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT16_SEMIPLANAR_444
499+
elif case_desc(d, 8, 8, 8, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar420):
500+
numberOfChannels[0] = 3
501+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_PLANAR_420
502+
elif case_desc(d, 16, 16, 16, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar420):
503+
numberOfChannels[0] = 3
504+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT16_PLANAR_420
505+
elif case_desc(d, 8, 8, 8, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar422):
506+
numberOfChannels[0] = 3
507+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_PLANAR_422
508+
elif case_desc(d, 16, 16, 16, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar422):
509+
numberOfChannels[0] = 3
510+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT16_PLANAR_422
511+
elif case_desc(d, 8, 8, 8, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar444):
512+
numberOfChannels[0] = 3
513+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT8_PLANAR_444
514+
elif case_desc(d, 16, 16, 16, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar444):
515+
numberOfChannels[0] = 3
516+
format[0] = cydriver.CUarray_format_enum.CU_AD_FORMAT_UINT16_PLANAR_444
417517
else:
418518
return cudaErrorInvalidChannelDescriptor
419519

420520
if d[0].f in (cudaChannelFormatKind.cudaChannelFormatKindNV12,
421521
cudaChannelFormatKind.cudaChannelFormatKindUnsignedBlockCompressed6H,
422-
cudaChannelFormatKind.cudaChannelFormatKindSignedBlockCompressed6H,):
522+
cudaChannelFormatKind.cudaChannelFormatKindSignedBlockCompressed6H,
523+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar420,
524+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar420,
525+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar422,
526+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar422,
527+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar444,
528+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar444,
529+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar420,
530+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar420,
531+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar422,
532+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar422,
533+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar444,
534+
cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar444,):
423535
if numberOfChannels[0] != 3:
424536
return cudaErrorInvalidChannelDescriptor
425537
else:
@@ -536,6 +648,49 @@ cdef cudaError_t getChannelFormatDescFromDriverDesc(cudaChannelFormatDesc* pRunt
536648
channel_size = 8
537649
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UNORM_INT_101010_2:
538650
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized1010102
651+
# 2b: CUDA 13.3/13.4 packed and multi-planar YUV formats
652+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_PACKED_422:
653+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed422
654+
channel_size = 8
655+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_PACKED_444:
656+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed444
657+
channel_size = 8
658+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_SEMIPLANAR_420:
659+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar420
660+
channel_size = 8
661+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT16_SEMIPLANAR_420:
662+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar420
663+
channel_size = 16
664+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_SEMIPLANAR_422:
665+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar422
666+
channel_size = 8
667+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT16_SEMIPLANAR_422:
668+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar422
669+
channel_size = 16
670+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_SEMIPLANAR_444:
671+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar444
672+
channel_size = 8
673+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT16_SEMIPLANAR_444:
674+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar444
675+
channel_size = 16
676+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_PLANAR_420:
677+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar420
678+
channel_size = 8
679+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT16_PLANAR_420:
680+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar420
681+
channel_size = 16
682+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_PLANAR_422:
683+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar422
684+
channel_size = 8
685+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT16_PLANAR_422:
686+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar422
687+
channel_size = 16
688+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT8_PLANAR_444:
689+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar444
690+
channel_size = 8
691+
elif pDriverDesc[0].Format == cydriver.CU_AD_FORMAT_UINT16_PLANAR_444:
692+
pRuntimeDesc[0].f = cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar444
693+
channel_size = 16
539694
else:
540695
return cudaErrorInvalidChannelDescriptor
541696

‎cuda_bindings/tests/test_cudart.py‎

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1985,3 +1985,74 @@ def test_cudaGraphConditionalHandleCreate_v2():
19851985

19861986
(err,) = cudart.cudaGraphDestroy(graph)
19871987
assertSuccess(err)
1988+
1989+
1990+
@pytest.mark.agent_authored(model="claude-sonnet-4-6")
1991+
def test_ffi_phase2_normalized_channel_kinds_round_trip():
1992+
"""Normalized 8/16-bit channel kinds passed validity but hit the else-branch
1993+
in the case_desc chain and returned cudaErrorInvalidChannelDescriptor.
1994+
Verify they no longer error at the binding layer (CUDA may still reject
1995+
the format on a given device, which is fine)."""
1996+
pairs = [
1997+
(8, 0, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized8X1),
1998+
(8, 8, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized8X2),
1999+
(8, 8, 8, 8, cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized8X4),
2000+
(8, 0, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized8X1),
2001+
(8, 8, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized8X2),
2002+
(8, 8, 8, 8, cudart.cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized8X4),
2003+
(16, 0, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized16X1),
2004+
(16, 16, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized16X2),
2005+
(16, 16, 16, 16, cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsignedNormalized16X4),
2006+
(16, 0, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized16X1),
2007+
(16, 16, 0, 0, cudart.cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized16X2),
2008+
(16, 16, 16, 16, cudart.cudaChannelFormatKind.cudaChannelFormatKindSignedNormalized16X4),
2009+
]
2010+
for x, y, z, w, kind in pairs:
2011+
desc = cudart.cudaChannelFormatDesc()
2012+
desc.x, desc.y, desc.z, desc.w, desc.f = x, y, z, w, kind
2013+
err, arr = cudart.cudaMallocArray(desc, 4, 1, 0)
2014+
if err == cudart.cudaError_t.cudaSuccess:
2015+
(free_err,) = cudart.cudaFreeArray(arr)
2016+
assert free_err == cudart.cudaError_t.cudaSuccess
2017+
else:
2018+
assert err != cudart.cudaError_t.cudaErrorInvalidChannelDescriptor, (
2019+
f"Binding layer rejected {kind.name} -- case_desc arm may be missing"
2020+
)
2021+
2022+
2023+
@pytest.mark.agent_authored(model="claude-sonnet-4-6")
2024+
def test_ffi_phase2_new_yuv_channel_kinds_no_binding_error():
2025+
"""CUDA 13.3/13.4 packed and multi-planar YUV kinds must not be rejected
2026+
by the binding layer with cudaErrorInvalidChannelDescriptor."""
2027+
three_ch_8 = [
2028+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar420,
2029+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar422,
2030+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8SemiPlanar444,
2031+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar420,
2032+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar422,
2033+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Planar444,
2034+
]
2035+
three_ch_16 = [
2036+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar420,
2037+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar422,
2038+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned16SemiPlanar444,
2039+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar420,
2040+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar422,
2041+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned16Planar444,
2042+
]
2043+
four_ch_8 = [
2044+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed422,
2045+
cudart.cudaChannelFormatKind.cudaChannelFormatKindUnsigned8Packed444,
2046+
]
2047+
for bits, w_bit, kinds in ((8, 0, three_ch_8), (16, 0, three_ch_16), (8, 8, four_ch_8)):
2048+
for kind in kinds:
2049+
desc = cudart.cudaChannelFormatDesc()
2050+
desc.x, desc.y, desc.z, desc.w, desc.f = bits, bits, bits, w_bit, kind
2051+
err, arr = cudart.cudaMallocArray(desc, 4, 1, 0)
2052+
if err == cudart.cudaError_t.cudaSuccess:
2053+
(free_err,) = cudart.cudaFreeArray(arr)
2054+
assert free_err == cudart.cudaError_t.cudaSuccess
2055+
else:
2056+
assert err != cudart.cudaError_t.cudaErrorInvalidChannelDescriptor, (
2057+
f"Binding layer rejected {kind.name} -- arm may be missing"
2058+
)

0 commit comments

Comments
 (0)