.version 8.0
.target sm_70
.address_size 64

.entry clean_kernel(
    .param .u64 param_base
)
{
    .reg .u32 %r<10>;
    .reg .u64 %rd<10>;

    // Load base address
    ld.param.u64 %rd0, [param_base];

    // SAFE: Compute address using constant offset
    // This is clean: %rd3 depends on %rd0 (param) and 128 (constant)
    add.u64 %rd3, %rd0, 128;

    // Store to computed address (should PASS)
    mov.u32 %r5, 0xCAFEBABE;
    st.global.u32 [%rd3], %r5;

    ret;
}
