shamalgs::atomic::DynamicIdGenerator< int_t, group_size > Class Template Reference

Sycl utility to dynamically generate group ids. More...

Public Member Functions
	DynamicIdGenerator (sycl::queue &q)
	Construct DynamicIdGenerator
AccessedDynamicIdGenerator< int_t, group_size >	get_access (sycl::handler &cgh)
	Get the access to DynamicIdGenerator returning the accessed variants AccessedDynamicIdGenerator

Public Attributes
sycl::buffer< int_t >	group_id
	the buffer used for group_id synchronization

Detailed Description

template<class int_t, u32 group_size>
class shamalgs::atomic::DynamicIdGenerator< int_t, group_size >

Sycl utility to dynamically generate group ids.

The goal is to affect each worker with a unique id in growing order, e.g. worker group 2 can not start if worker group 1 is not started The performance overhead is minimal (10^-11 s/element on A100)

Todo:

add figure for overhead measurement

Example :

DynamicIdGenerator<i32, group_size> id_gen(q);
q.submit([&](sycl::handler &cgh) {
 
   auto dyn_id = id_gen.get_access(cgh);
 
   cgh.parallel_for(sycl::nd_range<1>{len, group_size},
       [=](sycl::nd_item<1> id) {
 
           atomic::DynamicId<i32> group_id = dyn_id.compute_id(id);
 
           u32 group_tile_id = group_id.dyn_group_id;
           u32 global_id = group_id.dyn_global_id;
 
           // kernel execution
 
       });
 
});

Template Parameters

int_t	the int type used by the counter (preferentially u32 or u64)
group_size	the group size used in SYCL

Definition at line 128 of file DynamicIdGenerator.hpp.

Constructor & Destructor Documentation

DynamicIdGenerator()

template<class int_t , u32 group_size>

shamalgs::atomic::DynamicIdGenerator< int_t, group_size >::DynamicIdGenerator ( sycl::queue &  q )

inlineexplicit

Construct DynamicIdGenerator

Parameters

q	the SYCL queue

Definition at line 141 of file DynamicIdGenerator.hpp.

Here is the call graph for this function:

Member Function Documentation

get_access()

template<class int_t , u32 group_size>

AccessedDynamicIdGenerator< int_t, group_size > shamalgs::atomic::DynamicIdGenerator< int_t, group_size >::get_access ( sycl::handler &  cgh )

inline

Get the access to DynamicIdGenerator returning the accessed variants AccessedDynamicIdGenerator

Parameters

cgh	the SYCL command group handler

Returns

AccessedDynamicIdGenerator<int_t, group_size> the accessed DynamicIdGenerator

Definition at line 152 of file DynamicIdGenerator.hpp.

Member Data Documentation

group_id

template<class int_t , u32 group_size>

sycl::buffer<int_t> shamalgs::atomic::DynamicIdGenerator< int_t, group_size >::group_id

the buffer used for group_id synchronization

Definition at line 134 of file DynamicIdGenerator.hpp.

---

The documentation for this class was generated from the following file:

• shamalgs/include/shamalgs/atomic/DynamicIdGenerator.hpp