Pull requests / #2327

#2327 qwen35: widen F16 by hand so upload_f32 builds on CUDA 12.1

closed · @nekomario28 · 0 comments · View on GitHub

AMD / HIPNVIDIA / CUDA

Description

Fixes Niko1221/Strata#2322.

`upload_f32` runs on the host and widens F16 with `__half2float(__ushort_as_half(...))`. `__ushort_as_half` is a `__device__` intrinsic, so on CUDA toolkits that do not have a host form (12.1 and 12.4 per the report) `nvcc` refuses it:

```
src/qwen35/q36_engine.cu:809: calling a __device__ function "__ushort_as_half(unsigned short)" from a __host__ function "upload_f32" is not allowed
```

This widens the bits the way the neighbouring `bf16_to_f32` already does for BF16, so the host path stops naming a device intrinsic and the file compiles on every toolkit.

## Testing

- `scripts/syntax_check.sh src/qwen35/q36_engine.cu` (HIP compat shim) passes; it passed before too, so it is a regression guard here, not a reproduction of the CUDA-only error.
- The conversion itself was checked against a reference for every 16-bit input: `f16_to_f32` matches NumPy's `float16` -> `float32` bit-for-bit across all 65536 values (0 mismatches; NaN treated as a class). That covers the subnormal branch.
- Not run: a real CUDA build (no NVIDIA toolchain here), which is the build the report is about. AGENTS.md asks for builds of touched backends; this is the boundary.

Related on strata.com

Editorial links to help you install, pick models, or read release notes — not part of the upstream thread.