cursed.cpp

__device__ uint32_t scopeViolation(bool* secretlyAlwaysTrue,
                                   bool* secretlyAlwaysFalse,
                                   uint32_t x) {
    asm volatile("{");
    asm volatile(".reg .b32 tmp;");
    asm volatile("mov.b32 tmp, 76;"); // <- Write to unused variable.
    asm volatile("}");
    uint32_t output;
    if (*secretlyAlwaysFalse) {
        asm volatile(".reg .b32 tmp;"); // <- Declarations count, even if unreachable.
        asm volatile("mov.b32 tmp, 42;"); // <- Never runs
        asm volatile("{");
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wptx-warnings"
        asm volatile(".reg .b32 tmp;"); // <- Inaccessible local variable.
#pragma clang diagnostic pop
        asm volatile("}");
    }
    if (*secretlyAlwaysTrue) {
        asm volatile("mov.b32 tmp, 4;"); // <- Actually runs.
    }

    asm volatile("add.u32 tmp, %0, tmp;" :: "r"(x));
    if (*secretlyAlwaysTrue) {
        asm volatile("add.u32 tmp, %0, tmp;" :: "r"(x));
    }
    asm volatile("mov.b32 %0, tmp;" : "=r"(output));

    // Overall effect is `return x * 2 + 4;`.
    return output;
}