cuda::experimental::stf::exec_place_scope

class exec_place_scope

RAII guard that activates an execution place and restores the previous one on destruction.

This class provides a scoped mechanism for temporarily switching the active execution place. When constructed, it activates the given execution place (e.g., sets the current CUDA device). When destroyed, it restores the previous execution place that was active before construction.

For grids, the index specifies which sub-place to activate. For scalar places, the index should be 0 (the default).

The guard is non-copyable but movable (like std::unique_lock).

Example usage:

// Scalar place activation
{
  auto active = exec_place::device(1).activate();
  // Device 1 is now active
  // ... perform operations on device 1 ...
}
// Previous device is restored

// Grid iteration
exec_place grid = make_grid(...);
for (size_t i = 0; i < grid.size(); i++) {
  auto active = grid.activate(i);
  // grid[i] is now active
  kernel<<<..., active.place().getStream()>>>(...);
}

Public Functions

exec_place_scope() = default

Default constructor creates an inactive scope.

inline exec_place_scope(exec_place place, size_t idx = 0)

Constructs the guard and activates the sub-place at the given index.

Parameters:

• place – The execution place (or grid) containing the sub-place to activate

• idx – The index of the sub-place to activate (default 0 for scalar places)

template<typename T = void>
inline exec_place_scope(

const data_place&

)

Deleted constructor for data_place to prevent accidental misuse.

Use data_place::affine_exec_place() to get the exec_place first.

inline ~exec_place_scope()

Destructor that restores the previous execution place (if not moved-from).

exec_place_scope(const exec_place_scope&) = delete

exec_place_scope &operator=(const exec_place_scope&) = delete

inline exec_place_scope(exec_place_scope &&other) noexcept

inline exec_place_scope &operator=(

exec_place_scope &&other

) noexcept

inline const exec_place &place() const

Get the currently active sub-place.

inline size_t index() const

Get the index within the grid (0 for scalar places)

inline bool is_active() const

Check if this scope is active (not moved-from)

inline void reset()

Early deactivation - restores previous state and marks scope as inactive.

After calling reset(), the destructor becomes a no-op. Calling reset() on an inactive scope is safe (no-op).