ballot and ffs question

[zjin-lcf]

The graph coloring example in CUDA shows the functions __ballot and __ffs.

https://userweb.cs.txstate.edu/~burtscher/research/ECL-GC/ECL-GC_12.cu

Do they have SYCL equivalents ?

Thanks

[Michoumichmich]

For __ffs you can take a loot at sycl::ctz. For ballot you could try the following, but what would happen if you have more than 32 work-items in a sub-group/warp ?

    uint32_t ballot(sycl::sub_group sg, int predicate) {
        size_t id = sg.get_local_linear_id();
        uint32_t local_val = (predicate ? 1u : 0u) << id;
        return sycl::reduce_over_group(sg, local_val, sycl::plus<>());
    }

[zjin-lcf]

Thank you for the solutions! I will give them a try. The SYCL subgroup extension might mention 'ballot' at some time.

[Michoumichmich]

You should be able to assert sycl::popcount(ballot(it.get_sub_group(), 1)) == it.get_sub_group().get_local_range().size(). The version with sycl::bit_and should work too but it fails to compile with _Z33__spirv_GroupNonUniformBitwiseAndjjj so I don't know

[zjin-lcf]

I find that reduce_over_group() and plus<> are not available in the oneAPI toolkit 2021.3.