porting CUDA __activemask()

[jinz2014]

What is the SYCL function for CUDA's __activemask() ?

__activemask() returns a bit mask of currently active warp lanes. The nth bit of the result is 1 if the nth warp lane is active

[AlexeySachkov]

Hi @jinz2014,

In the core SYCL 2020 specification there are no functions available which allow you to operate with masks of active work-items within a group, nor there are group functions which can be called in a non-uniform control-flow (where only a subset of work-items within a group encountered a group function call).

We do have some extensions for sub-group masks and non-uniform groups, but it seems like we don't support a "combination" of the two.
Tagging @Pennycook here for awareness, this feels like a gap: why can't we call group_ballot in a non-uniform control flow? There is OpGroupNonUniformBallot for us to implement that.

Notes about the extensions I was talking about:

sycl_ext_oneapi_sub_group_mask introduces the mask type and group_ballot function, which allows to get a mask which corresponds to work-items for which a predicate evaluates to true. But once again, group_ballot can only be called from uniform control-flow, meaning that you can't use it to query list of currently active work-items.

We have a couple of extensions for non-uniform groups:
sycl_ext_oneapi_non_uniform_groups and sycl_ext_oneapi_tangle. However, they won't let you know which exact work-items in a parent group are active, they will just give you an ability to perform a group operation on a subset of work-items from a parent group instead of doing that on the whole parent group.
Please also note that the extension specs have just been updated and our implementation for them doesn't yet match the latest revision.

[Pennycook]

Tagging @Pennycook here for awareness, this feels like a gap: why can't we call group_ballot in a non-uniform control flow? There is OpGroupNonUniformBallot for us to implement that.

The extensions we have are at two different levels of abstraction, and we need feedback about realistic use-cases.

We originally introduced sycl_ext_oneapi_sub_group_mask because it exposes OpGroupNonUniformBallot fairly directly. But we had trouble getting the concepts of masks and groups to play nicely together: you'd basically need to pass a sycl::sub_group and a mask to every group function, and the user would be responsible for any bitwise manipulation of the mask.

sycl_ext_oneapi_non_uniform_groups is a much nicer interface, largely because of its alignment with the rest of SYCL. When running on the CUDA backend, calling sycl::get_opportunistic_group() returns an opportunistic_group (fragment) that represents the set of work-items corresponding to the CUDA threads with bits set in the current __activemask().

We think everything somebody might want to do with a mask can be expressed much more easily via the group interface. For example, fragment::get_local_id() returns the calling work-item's position in the mask, and fragment::get_local_range() returns the number of bits set in the mask -- but a user doesn't need to understand or write any bitwise operations to get those results. Where in CUDA you'd typically have to write a loop iterating over the mask and shuffling to perform a reduction, you can just call reduce_over_group using the fragment, and the implementation does the reduction for you.

@jinz2014: My recommendation would be to try and rewrite the CUDA code in terms of an opportunistic_group instead of __activemask. If you find something that you can't express easily, or something where the mask representation is absolutely required, it would be very helpful if you could share that case with us to make sure we address it.

[jinz2014]

Thanks.

https://github.com/leannmlindsey/gpu_local_ht/blob/cff69f4209c0778c9f29004a822e95674ea8fc61/src/kernel.cpp#L617

        unsigned mask = __activemask();
        unsigned active = mask & 1; // zero if lane 0 has returned
        if(active == 0)
            break; // return if lane 0 has returned

@leannmlindsey