[Engine] Lower nvidia architecture floor to allow older maxwell cards.

I simply change the code to allow architecture => 52 and it worked.

Thanks for the issue. To keep the tracker readable, every item needs an area prefix in the title and the parts of the issue form filled in. This one is missing:

Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other], then a short description. Example: "[Engine] decode stops after 2 minutes".

Edit the issue (title and description) and this note goes away on its own. If it is still missing parts in 5 days it will be closed, which is not a rejection: fix it and comment /reopen.

Details: CONTRIBUTING.md

Thanks. Which card is it and which Strata version? Maxwell (sm_52) has no fast half-precision path and the experimental CUDA 12 engine is built for sm_60 and newer, so changing the number alone may start but run slowly or give wrong output. If you tested it, please tell us the model, the speed and whether the answer text looked right.

Thanks — following up on the questions above: which Maxwell card is it (model and VRAM) and which Strata version did you patch? If you ran it, the model, the decode speed and whether the answer text looked right would tell us whether lowering the floor is safe. A diff of the change you made would also help. If you can, add [Engine] to the title so it routes to the engine queue.

Thanks, this follows the format now.

I was able to get Strata (v0.1.44.1) building and running cleanly on Maxwell (sm_52, tested on dual Quadro M6000 24GB on Ubuntu 26.04.1 LTS with NVIDIA driver 580.178.04).

To get past the initial build gating, I only had to lower the minimum architecture checks down to 52 in two files:

setup.py (line 1092):

def sm60_card(arch) -> bool:
    return 52 <= int(arch) <= 70

CMakeLists.txt (line 184):
if(_base LESS 52)
message(FATAL_ERROR "CMAKE_CUDA_ARCHITECTURES '${CMAKE_CUDA_ARCHITECTURES}': even the experimental build "
"needs compute capability 5.2 (sm_52) or newer")
endif()

Build command
Bash
cmake -B build -G Ninja
-DSTRATA_EXPERIMENTAL_SM60=ON
-DCMAKE_CUDA_ARCHITECTURES=52
cmake --build build -j
Verification & Performance
Model: Qwen3.8-Flash-Next-GSQ-RCO-IQ1

Hardware: 2x Quadro M6000 24GB (48GB total VRAM)

Stability: No runtime kernel or PTX crashes occurred during inference—execution and emulation fallback paths handled execution cleanly on the modern driver stack.

Generation speed: Averaging around ~22 t/s decode tokens.

Maxwell sm_52 floor for the opt-in old-card build (#3218)

setup: Maxwell gets a warning with the build recipe, not an env switch (

Merge f-027-gh-eng-s (#3291 #3955 #3317 #4036 #2893 #3218 #3200: serv…

Thank you. Fixed in 0.2.7: Maxwell gets a warning with the build recipe, not an env switch. We could check it with unit tests and builds, but not on your hardware or with your workload. Please try 0.2.7 and tell us how it goes.

Release: https://github.com/Niko1221/Strata/releases/tag/v0.2.7

Tested v0.2.7 with -DSTRATA_EXPERIMENTAL_SM60=ON -DCMAKE_CUDA_ARCHITECTURES=52 on Ubuntu 26.04.1 LTS with dual Quadro M6000 24GB.

The CMake configure warning triggered cleanly as intended:

Building for sm_52 (Maxwell): community-tested on one reporter's Quadro M6000 only, no native half-precision arithmetic...

However, compilation halted at [1/220] on src/prefill/moe_ws_down.cu due to fma.rn.f16x2 and add.rn.f16x2:

ptxas error : Feature 'f16 arithemetic and compare instructions' requires .target sm_53 or higherCause
Maxwell sm_52 lacks native hardware FP16 math instructions (native f16x2 instructions require compute capability 5.3 or 6.0+). The direct PTX inline assembly calls to fma.rn.f16x2 and add.rn.f16x2 cause ptxas to reject the module when targeting sm_52.

Working Fix
In src/prefill/moe_ws_down.cu (lines 43–65), guarding those functions with #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530 and promoting them through FP32 allows ptxas to assemble cleanly without errors:

Code snippet
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530
__device__ __forceinline__ uint32_t hfma2(const uint32_t a, const uint32_t b, const uint32_t c) {
    float fa_lo = __half2float(__ushort_as_half((unsigned short)(a & 0xFFFFu)));
    float fa_hi = __half2float(__ushort_as_half((unsigned short)(a >> 16)));
    float fb_lo = __half2float(__ushort_as_half((unsigned short)(b & 0xFFFFu)));
    float fb_hi = __half2float(__ushort_as_half((unsigned short)(b >> 16)));
    float fc_lo = __half2float(__ushort_as_half((unsigned short)(c & 0xFFFFu)));
    float fc_hi = __half2float(__ushort_as_half((unsigned short)(c >> 16)));
    
    unsigned short res_lo = __half_as_ushort(__float2half(fmaf(fa_lo, fb_lo, fc_lo)));
    unsigned short res_hi = __half_as_ushort(__float2half(fmaf(fa_hi, fb_hi, fc_hi)));
    return (uint32_t)res_lo | ((uint32_t)res_hi << 16);
}
__device__ __forceinline__ uint32_t hadd2(const uint32_t a, const uint32_t b) {
    float fa_lo = __half2float(__ushort_as_half((unsigned short)(a & 0xFFFFu)));
    float fa_hi = __half2float(__ushort_as_half((unsigned short)(a >> 16)));
    float fb_lo = __half2float(__ushort_as_half((unsigned short)(b & 0xFFFFu)));
    float fb_hi = __half2float(__ushort_as_half((unsigned short)(b >> 16)));
    
    unsigned short res_lo = __half_as_ushort(__float2half(fa_lo + fb_lo));
    unsigned short res_hi = __half_as_ushort(__float2half(fa_hi + fb_hi));
    return (uint32_t)res_lo | ((uint32_t)res_hi << 16);
}
__device__ __forceinline__ float h2sum(const uint32_t h) {
    float lo = __half2float(__ushort_as_half((unsigned short)(h & 0xFFFFu)));
    float hi = __half2float(__ushort_as_half((unsigned short)(h >> 16)));
    return lo + hi;
}
__device__ __forceinline__ float f16_to_f32(const uint32_t bits16) {
    return __half2float(__ushort_as_half((unsigned short)(bits16 & 0xFFFFu)));
}
#else
__device__ __forceinline__ uint32_t hfma2(const uint32_t a, const uint32_t b, const uint32_t c) {
    uint32_t d;
    asm("fma.rn.f16x2 %0, %1, %2, %3;" : "=r"(d) : "r"(a), "r"(b), "r"(c));
    return d;
}
__device__ __forceinline__ uint32_t hadd2(const uint32_t a, const uint32_t b) {
    uint32_t d;
    asm("add.rn.f16x2 %0, %1, %2;" : "=r"(d) : "r"(a), "r"(b));
    return d;
}

<img width="903" height="560" alt="Image" src="https://github.com/user-attachments/assets/81cce368-d8e8-40fb-a147-fe33a3fdd98a" />
<img width="906" height="884" alt="Image" src="https://github.com/user-attachments/assets/59beee97-80be-4322-9787-aea740379baf" />
<img width="919" height="507" alt="Image" src="https://github.com/user-attachments/assets/ff9de4a9-f980-4aac-867f-a15875f9b258" />

__device__ __forceinline__ float h2sum(const uint32_t h) {
    float f;
    asm("{\n\t.reg .f16 lo, hi;\n\t.reg .f32 a, b;\n\tmov.b32 {lo, hi}, %1;\n\t"
        "cvt.f32.f16 a, lo;\n\tcvt.f32.f16 b, hi;\n\tadd.f32 %0, a, b;\n\t}"
        : "=f"(f) : "r"(h));
    return f;
}
__device__ __forceinline__ float f16_to_f32(const uint32_t bits16) {
    float f;
    asm("{\n\t.reg .b16 h;\n\tcvt.u16.u32 h, %1;\n\tcvt.f32.f16 %0, h;\n\t}" : "=f"(f) : "r"(bits16));
    return f;
}
#endif

A correction to our reply of 10-27: this is not fixed. Your 0.2.7 log shows the sm_52 build stopping at src/prefill/moe_ws_down.cu, because fma.rn.f16x2 needs sm_53 or newer. The 0.2.7 change did not cover that kernel. It is on our list for the next update. Thank you for testing it again.

[Other] track the local Maxwell (sm_52) build in the fork (not upstre…

Navigation Menu

Repository navigation

[Engine] Lower nvidia architecture floor to allow older maxwell cards. #3218

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

[Engine] Lower nvidia architecture floor to allow older maxwell cards. #3218

Description

Activity

github-actions commented on Oct 19, 2026

Niko1221 commented on Oct 21, 2026

Niko1221 commented on Oct 22, 2026

github-actions commented on Oct 25, 2026

jmuffins commented on Oct 25, 2026

Niko1221 commented on Oct 27, 2026

jmuffins commented on Oct 29, 2026

jmuffins commented on Oct 29, 2026

Niko1221 commented on Oct 29, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions