Hacktoberfest 2026: the issues maintainers tagged for October, open and beginner-friendly. Browse Hacktoberfest issues

[Engine] SM70/Volta build fails: q36_engine.cu calls __device__ __half_as_ushort from host function verify_q8

Open Beginner friendly
  • Is something wrong?
#3,524 2 comments 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 3 days

Nobody has claimed this yet.

Assessment

Difficulty
2/5
Estimated time
1-3 hours
Newbie friendliness
74/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Active
Tech stack
cpp
Domain
backend, build-system

Research direction

Start in src/qwen35/q36_engine.cu at the host-side verify_q8() debug output around line 1659, and check how that translation unit handles host and device half conversions. Replace the host call to __half_as_ushort() with a host-safe way to read the half value's bits. Build the strata and strata_qwen35 targets with STRATA_EXPERIMENTAL_SM60=ON and CUDA architecture 70; done means the build succeeds.

Written by the indexing model from the issue text.

Description

Summary

Building the 0.1.45+ engine with -DSTRATA_EXPERIMENTAL_SM60=ON (Volta / SM70, e.g. Tesla V100) fails: src/qwen35/q36_engine.cu calls the device-only intrinsic __half_as_ushort() from a host function.

Environment
Strata main @ 9de04cd4 (Control 0.1.46.3, engine 0.1.46.3)
OS Linux (Ubuntu 24.04) x86_64
GPU Tesla V100-SXM2-16GB (SM70, driver 580.178.04)
CUDA 12.0.140 (build-cuda12, CMAKE_CUDA_ARCHITECTURES=70)
Host compiler g++-12 (12.4.0)
CMake option STRATA_EXPERIMENTAL_SM60=ON

Error

src/qwen35/q36_engine.cu(1659): error: calling a __device__ function("__half_as_ushort(    ::__half)") from a __host__ function("verify_q8") is not allowed

1 error detected in the compilation of "src/qwen35/q36_engine.cu".
ninja: build stopped: subcommand failed.

setup.sh --update retries the build once, hits the same error, then falls back to the previously installed engine (could not compile the updated engine: starting the installed one), so the update silently no-ops while claiming success.

Cause

verify_q8() is a host function (it only runs under STRATA_Q36_VDEBUG for debug output). Line 1659 casts the fp16 value of the CPU butterfly sum to its bit pattern:

std::fprintf(stderr, "VDEBUG block %zu: ds new %08x old %08x, qs bytes %s, cpu butterfly sum %.9g (half %04x)\n",
             blk, ds_a, ds_b, std::memcmp(&a[o + 4], &b[o + 4], 32) == 0 ? "equal" : "DIFFER", (double) t[0],
             (unsigned) __half_as_ushort(__float2half(t[0])));

__half_as_ushort is declared __device__ in cuda_fp16.h; on this arch/header combination nvcc does not accept the host call. It is not guarded by #ifdef __CUDA_ARCH__, so the SM60/SM70 path cannot build at all.

Fix

Read the bit pattern on the host directly — same value, no device intrinsic:

unsigned short t0_bits;
{ const __half t0h = __float2half(t[0]); std::memcpy(&t0_bits, &t0h, 2); }
std::fprintf(stderr, "... (half %04x)\n", ..., (unsigned) t0_bits);

(or __half2ushort_rn / union pun; anything host-side). I verified this builds and runs:

  • cmake --build build-cuda12 --target strata strata_qwen35 → OK
  • ./update.sh → engine compiled: .../engine-cuda12/strata, Strata is updated (engine 0.1.46.3)
  • Model starts, /health 200, 262144 ctx, chat completions work

Suggestion

Guard the debug cast with #if defined(__CUDA_ARCH__) or use the host-safe conversion above, so the SM60 experimental path keeps compiling. Secondarily: when the engine build fails, update.sh exiting 0 with "Strata is updated" is misleading — a non-zero exit code (or a clear "engine NOT updated" banner) would save users from thinking they upgraded when they did not.

Dominant language
C++
Stars
11.6k
Forks
1k
No merged PRs in 30d

Getting set up

This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

More from Niko1221/Strata

All issues in Niko1221/Strata

Similar issues

More C++ issues

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.