From 389ada4fac2258a5833fb5c9ab26a730933d8496 Mon Sep 17 00:00:00 2001
From: Rob Taylor <rob.taylor@chipflow.io>
Date: Wed, 4 Mar 2026 18:54:45 +0000
Subject: [PATCH] Fix CUDA/HIP build: use i32 typedef instead of bare int for
 arrival_state_offset
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

The bindgen-generated Rust FFI bindings don't recognize bare C `int` —
they need `i32` (typedef for int32_t from ulib/types.hpp) to map
correctly to Rust's i32.

Co-developed-by: Claude Code v2.1.50 (claude-opus-4-6)
---
 csrc/kernel_v1.cu       | 4 ++--
 csrc/kernel_v1.hip.cpp  | 4 ++--
 csrc/kernel_v1_impl.cuh | 4 ++--
 3 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/csrc/kernel_v1.cu b/csrc/kernel_v1.cu
index a89d15e..bf47c1d 100644
--- a/csrc/kernel_v1.cu
+++ b/csrc/kernel_v1.cu
@@ -25,7 +25,7 @@ void simulate_v1_noninteractive_simple_scan_cuda(
   usize num_cycles,
   usize state_size,
   u32 *states_noninteractive,
-  int arrival_state_offset
+  i32 arrival_state_offset
   )
 {
   const u32 *timing_constraints = nullptr;
@@ -59,7 +59,7 @@ void simulate_v1_noninteractive_timed_cuda(
   u32 *states_noninteractive,
   const u32 *timing_constraints,
   u8 *event_buffer,
-  int arrival_state_offset
+  i32 arrival_state_offset
   )
 {
   void *arg_ptrs[12] = {
diff --git a/csrc/kernel_v1.hip.cpp b/csrc/kernel_v1.hip.cpp
index a3fa933..a941a94 100644
--- a/csrc/kernel_v1.hip.cpp
+++ b/csrc/kernel_v1.hip.cpp
@@ -42,7 +42,7 @@ void simulate_v1_noninteractive_simple_scan_hip(
   usize num_cycles,
   usize state_size,
   u32 *states_noninteractive,
-  int arrival_state_offset
+  i32 arrival_state_offset
   )
 {
   validate_warp_size();
@@ -79,7 +79,7 @@ void simulate_v1_noninteractive_timed_hip(
   u32 *states_noninteractive,
   const u32 *timing_constraints,
   u8 *event_buffer,
-  int arrival_state_offset
+  i32 arrival_state_offset
   )
 {
   validate_warp_size();
diff --git a/csrc/kernel_v1_impl.cuh b/csrc/kernel_v1_impl.cuh
index 2a4dfb9..b9a26c1 100644
--- a/csrc/kernel_v1_impl.cuh
+++ b/csrc/kernel_v1_impl.cuh
@@ -52,7 +52,7 @@ __device__ void simulate_block_v1(
   u32 clock_period_ps,
   EventBuffer *__restrict__ event_buffer,
   u32 cycle_i,
-  int arrival_state_offset  // offset in output_state for arrival data (0 = disabled)
+  i32 arrival_state_offset  // offset in output_state for arrival data (0 = disabled)
   )
 {
   int script_pi = 0;
@@ -574,7 +574,7 @@ __global__ void simulate_v1_noninteractive_simple_scan(
   u32 *__restrict__ states_noninteractive,
   const u32 *__restrict__ timing_constraints,
   EventBuffer *__restrict__ event_buffer,
-  int arrival_state_offset
+  i32 arrival_state_offset
   )
 {
   assert(num_blocks == gridDim.x);