[wasmi_c_api] Error on wasm_func_call, argument vec mismatch on C vs Rust

[dimitris-aspetakis]

Hello!

I'm totally new to this project, and I've just recently started using the c_api (v51.1) to call some webassembly functions.

Here is the snippet that gives me trouble.

    // Prepare the VC function's parameters:
    //   lower_bound(u64), upper_bound(u64),
    //   grid_x(u32),  grid_y(u32),  grid_z(u32),
    //   block_x(u32), block_y(u32), block_z(u32),
    //   original_kernel_args[...]
    size_t vc_param_count = 8 + (n_args - 2);

    printf("Parameter count: %zu\n", vc_param_count);

    wasm_val_t *vc_params = malloc(sizeof(wasm_val_t) * vc_param_count);
    if (!vc_params) {
      fprintf(
          stderr,
          "INTERPOSER ERROR: Failed to allocate memory for VC parameters\n");
      free(new_args);
      return cudaErrorLaunchFailure;
    }

    size_t param_idx = 0;

    // Parameters 1-2: Lower and upper bounds of global PTX state space (u64).
    vc_params[param_idx++] = (wasm_val_t)WASM_I64_VAL(bounds[0]);
    vc_params[param_idx++] = (wasm_val_t)WASM_I64_VAL(bounds[1]);

    // Parameters 3-5: Grid dimensions (u32).
    vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(gridDim.x);
    vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(gridDim.y);
    vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(gridDim.z);

    // Parameters 6-8: Block dimensions (u32).
    vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(blockDim.x);
    vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(blockDim.y);
    vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(blockDim.z);

    // Get VC function type signature for parameter type mapping.
    wasm_functype_t *vc_func_type = wasm_func_type(kernel_metadata->vc_func);
    const wasm_valtype_vec_t *vc_param_types = NULL;
    if (vc_func_type) {
      vc_param_types = wasm_functype_params(vc_func_type);
    }

    // Add original kernel parameters (excluding the last 2 bounds we added).
    // Use the VC function type signature for precise type casting.
    for (size_t i = 0; i < n_args - 2; i++) {
      size_t vc_param_index = 8 + i; // Skip the 8 standard parameters.
      wasm_valkind_t expected_kind = WASM_I64; // Default fallback.

      if (vc_param_types && vc_param_index < vc_param_types->size) {
        expected_kind = wasm_valtype_kind(vc_param_types->data[vc_param_index]);
      } else {
        fprintf(stderr,
                "INTERPOSER WARNING: Could not get VC function type info for "
                "parameter %zu\n",
                i);
      }

      // Cast kernel parameter based on VC function's expected type.
      switch (expected_kind) {
      case WASM_I32:
        printf("    -> Converting to WASM_I32: %d\n", *(int32_t *)args[i]);
        vc_params[param_idx++] = (wasm_val_t)WASM_I32_VAL(*(int32_t *)args[i]);
        break;
      case WASM_I64:
        printf("    -> Converting to WASM_I64: %ld\n", *(int64_t *)args[i]);
        vc_params[param_idx++] = (wasm_val_t)WASM_I64_VAL(*(int64_t *)args[i]);
        break;
      case WASM_F32:
        printf("    -> Converting to WASM_F32: %f\n", *(float *)args[i]);
        vc_params[param_idx++] = (wasm_val_t)WASM_F32_VAL(*(float *)args[i]);
        break;
      case WASM_F64:
        printf("    -> Converting to WASM_F64: %f\n", *(double *)args[i]);
        vc_params[param_idx++] = (wasm_val_t)WASM_F64_VAL(*(double *)args[i]);
        break;
      default:
        fprintf(stderr,
                "INTERPOSER WARNING: Unexpected VC function parameter type %d "
                "for arg %zu, using i64\n",
                expected_kind, i);
        vc_params[param_idx++] = (wasm_val_t)WASM_I64_VAL(*(int64_t *)args[i]);
        break;
      }
    }

    // Clean up the function type.
    if (vc_func_type) {
      wasm_functype_delete(vc_func_type);
    }

    // Prepare parameter and result vectors.
    wasm_val_vec_t params = {vc_param_count, vc_params};
    wasm_val_t result = WASM_INIT_VAL;
    wasm_val_vec_t result_vec = {1, &result};

    // Print all parameter types and values for debugging
    printf("INTERPOSER: VC function parameters:\n");
    for (size_t i = 0; i < vc_param_count; i++) {
      switch (vc_params[i].kind) {
      case WASM_I32:
        printf("  Param %zu: WASM_I32 = %d\n", i, vc_params[i].of.i32);
        break;
      case WASM_I64:
        printf("  Param %zu: WASM_I64 = %ld\n", i, vc_params[i].of.i64);
        break;
      case WASM_F32:
        printf("  Param %zu: WASM_F32 = %f\n", i, vc_params[i].of.f32);
        break;
      case WASM_F64:
        printf("  Param %zu: WASM_F64 = %f\n", i, vc_params[i].of.f64);
        break;
      default:
        printf("  Param %zu: Unknown type %d\n", i, vc_params[i].kind);
        break;
      }
    }

    // Call the VC function.
    wasm_trap_t *trap =
        wasm_func_call(kernel_metadata->vc_func, &params, &result_vec);

    if (trap) {
      fprintf(stderr, "INTERPOSER ERROR: VC function execution trapped\n");
      wasm_trap_delete(trap);
      free(vc_params);
      free(new_args);
      return cudaErrorLaunchFailure;
    }

    // Check the result - non-zero indicates bounds violation.
    if (result.of.i32 != 0) {
      printf("INTERPOSER WARNING: VC function detected bounds violation\n");
      _exit(-1);
    } else {
      printf("INTERPOSER: VC function verification passed!\n");
    }

    free(vc_params);
  } else {
    printf("INTERPOSER WARNING: No VC function available for this kernel\n");
  }

  return cudaLaunchKernel_original(func, gridDim, blockDim, new_args, sharedMem,
                                   stream);

With some print-debugging from inside wasmi I saw that the wasm_val_vec_t of the function parameters seems to contain only half of the arguments I'm trying to pass, with each second being lost. I'll first post my debug-print in func.rs:197 (inside the wasm_func_call implementation):

    let (wt_params, wt_results) = prepare_params_and_results(
        &mut dst,
        params.iter().map(|i| {
            let tmp: Val = i.to_val();
            println!("Arg type: {:?}", tmp.ty());
            tmp
        }),
        results.len(),
    );

And now, I'll post my output. The Param # lines come from my C code, and the Arg type ... lines from my wasmi print-debugging.

INTERPOSER: Calling VC function for verification...
Parameter count: 11
    -> Converting to WASM_F32: 2.000000
    -> Converting to WASM_I64: 137309628399616
    -> Converting to WASM_I64: 137309628400128
INTERPOSER: VC function parameters:
  Param 0: WASM_I64 = 0
  Param 1: WASM_I64 = -1
  Param 2: WASM_I32 = 1
  Param 3: WASM_I32 = 1
  Param 4: WASM_I32 = 1
  Param 5: WASM_I32 = 4
  Param 6: WASM_I32 = 1
  Param 7: WASM_I32 = 1
  Param 8: WASM_F32 = 2.000000
  Param 9: WASM_I64 = 137309628399616
  Param 10: WASM_I64 = 137309628400128
Arg type: I64
Arg type: I32
Arg type: I32
Arg type: I32
Arg type: F32
Arg type: I64

thread '<unnamed>' panicked at crates/c_api/src/val.rs:146:17:
internal error: entered unreachable code: `wasm_val_t`: cannot contain non-function reference values
stack backtrace:
   0: __rustc::rust_begin_unwind
             at /rustc/29483883eed69d5fb4db01964cdf2af4d86e9cb2/library/std/src/panicking.rs:697:5
   1: core::panicking::panic_fmt
             at /rustc/29483883eed69d5fb4db01964cdf2af4d86e9cb2/library/core/src/panicking.rs:75:14
   2: wasm_func_call
   3: cudaLaunchKernel
   4: _Z19__device_stub__axpyfPfS_
   5: main
   6: <unknown>
   7: __libc_start_main
   8: _start

I'm sorry for the long question, I felt like hitting a road block and couldn't find concrete examples of what I'm trying to do using wasmi's c_api — I really hope this isn't a silly error on my part (I didn't make this a github issue because I tend to believe that's the case here).

[dimitris-aspetakis]

Actually (now that I have a clearer mind and I'm not rage-debugging) I think I found something:

Here is the struct definition in C:

typedef struct wasm_val_t {
  wasm_valkind_t kind;
  union {
    int32_t i32;
    int64_t i64;
    float32_t f32;
    float64_t f64;
    struct wasm_ref_t* ref;
  } of;
} wasm_val_t;

And here in Rust:

pub struct wasm_val_t {
    /// The kind of the Wasm value.
    pub kind: wasm_valkind_t,
    /// The underlying data of the Wasm value classified by `kind`.
    pub of: wasm_val_union,
}

Importantly, wasm_val_union includes a u128 field, making it 16 bytes instead of the 8 in the C version:

pub union wasm_val_union {
    /// A Wasm 32-bit signed integer.
    pub i32: i32,
    /// A Wasm 64-bit signed integer.
    pub i64: i64,
    /// A Wasm 32-bit unsigned integer.
    pub u32: u32,
    /// A Wasm 64-bit unsigned integer.
    pub u64: u64,
    /// A Wasm 32-bit float.
    pub f32: f32,
    /// A Wasm 64-bit float.
    pub f64: f64,
    /// A Wasm `v128` value.
    pub v128: u128,
    /// A Wasm referenced object.
    pub ref_: *mut wasm_ref_t,
}

…and hence the difference in pointer arithmetic.

Now I believe it's closer to an actual issue, and not an error on my side, so I filed a bug report (#1669).