[BUG] CUDA error 719 in debug-mode tile adjoint with an unlaunched half-vector kernel

[shi-eric]

Bug description

CUDA error 719 occurs while running the tile struct tests in debug mode. The failure can be reproduced with a module containing two kernels:

• A differentiable tile kernel that maps and reduces a custom struct.
• An unlaunched tile kernel whose struct contains a wp.vec3h field.

The forward pass returns the expected loss. Tape.backward() fails when the kernels are compiled in debug mode.

Reproduction

import warp as wp

wp.config.mode = "debug"
wp.config.verify_cuda = True

TILE_M = wp.constant(8)
TILE_DIM = 64


@wp.struct
class TileMapStruct:
    x: wp.float32
    y: wp.vec3


@wp.struct
class HalfVectorStruct:
    value: wp.vec3h


@wp.func
def tile_map_struct_scale(value: TileMapStruct) -> TileMapStruct:
    result = TileMapStruct()
    result.x = value.x + wp.float32(1.0)
    result.y = value.y * wp.float32(2.0)
    return result


@wp.func
def tile_map_struct_add(a: TileMapStruct, b: TileMapStruct) -> TileMapStruct:
    result = TileMapStruct()
    result.x = a.x + b.x
    result.y = a.y + b.y
    return result


@wp.func
def tile_map_struct_sum(value: TileMapStruct) -> float:
    return value.x + value.y[0] + value.y[1] + value.y[2]


@wp.kernel(enable_backward=False)
def tile_half_vector_kernel(input: wp.array[HalfVectorStruct], output: wp.array[HalfVectorStruct]):
    values = wp.tile_load(input, shape=TILE_M)
    wp.tile_store(output, values)


@wp.kernel
def tile_map_struct_grad_kernel(input: wp.array[TileMapStruct], loss: wp.array[float]):
    i = wp.tid()
    values = wp.tile_load(input, shape=TILE_M, offset=i * TILE_M)
    scaled = wp.tile_map(tile_map_struct_scale, values)

    bias = TileMapStruct()
    bias.x = wp.float32(10.0)
    bias.y = wp.vec3(1.0, 2.0, 3.0)
    biased = wp.tile_map(tile_map_struct_add, scaled, bias)

    components = wp.tile_map(tile_map_struct_sum, biased)
    wp.tile_store(loss, wp.tile_sum(components), offset=i)


device = "cuda:0"
data = []

for i in range(TILE_M):
    value = TileMapStruct()
    value.x = float(i)
    value.y = wp.vec3(float(i), float(i + 1), float(i + 2))
    data.append(value)

input_wp = wp.array(data, dtype=TileMapStruct, requires_grad=True, device=device)
loss_wp = wp.zeros(1, dtype=float, requires_grad=True, device=device)

with wp.Tape() as tape:
    wp.launch_tiled(
        tile_map_struct_grad_kernel,
        dim=[1],
        inputs=[input_wp],
        outputs=[loss_wp],
        block_dim=TILE_DIM,
        device=device,
    )

print("loss:", loss_wp.numpy())

tape.backward(loss_wp)

print("gradient x:", input_wp.grad.numpy()["x"])
print("gradient y:", input_wp.grad.numpy()["y"])

Run it from a Warp checkout with a fresh kernel cache:

WARP_CACHE_PATH=/tmp/warp-tile-half-repro uv run repro.py

The forward pass completes:

loss: [380.]

The backward launch then fails:

Warp CUDA error 719: unspecified launch failure
(in function wp_cuda_context_check, .../warp/native/warp.cu:2924)

Warp Error: Error launching kernel:
tile_map_struct_grad_kernel on device cuda:0:
CUDA error detected: 719

Traceback (most recent call last):
  ...
  tape.backward(loss_wp)
  ...
RuntimeError: CUDA error detected: 719

Without wp.config.verify_cuda = True, the asynchronous error can appear at a later synchronizing operation instead. The original repository test reported it during a device-to-host copy:

test_tile_map_custom_struct_cuda_0 ... ERROR
Warp CUDA error 719: unspecified launch failure
(in function wp_memcpy_d2h, .../warp/native/warp.cu:1097)

Later CUDA operations in the same process also fail after the first error.

A successful run produces:

loss: [380.]
gradient x: [1. 1. 1. 1. 1. 1. 1. 1.]
gradient y: [[2. 2. 2.]
 [2. 2. 2.]
 [2. 2. 2.]
 [2. 2. 2.]
 [2. 2. 2.]
 [2. 2. 2.]
 [2. 2. 2.]
 [2. 2. 2.]]

Controls

Fresh processes on commit 18866f004d3a2e17ee0a953569f1f55ba2648cc4 produced these results:

Mode	Half-vector definitions present	CUDA stack limit	Result
Debug	Yes	Default, measured at 1,024 bytes	CUDA error 719 during Tape.backward()
Release	Yes	Default	Expected loss and gradients
Debug	No	Default	Expected loss and gradients
Debug	Yes	Requested and measured at 1,536 bytes	Expected loss and gradients

The 1,536-byte value is a tested passing value, not a measured minimum.

tile_half_vector_kernel is never launched. Removing HalfVectorStruct and tile_half_vector_kernel from the module makes the debug-mode run pass.

Changing block_dim from 64 to 8 does not prevent the failure at the default stack limit.

System information

• Commit: 18866f004d3a2e17ee0a953569f1f55ba2648cc4
• Warp: 1.17.0.dev4
• Python: 3.12.13
• NumPy: 2.5.0
• OS: Linux 6.8.0-1054-gke, x86_64, glibc 2.35
• GPU: NVIDIA RTX PRO 6000 Blackwell Server Edition MIG 1g.24gb
• Compute capability: SM 120
• GPU memory: 24 GiB
• NVIDIA driver: 595.58.03
• CUDA driver API: 13.2
• CUDA toolkit and NVRTC: 13.0

[shi-eric]

The fix for this issue will be expanded to cover the wp.bfloat16 case as well.