cuda::experimental::hierarchy::index

Defined in /home/runner/work/cccl/cccl/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh

template<typename Unit, typename Level>
auto cuda::experimental::hierarchy::index(const Unit&, const Level&)

Returns a 3-dimensional index of an entity the calling thread belongs to in a CUDA hierarchy level.

Returned index is in line with intrinsic CUDA indexing like threadIdx and blockIdx, extentded to more unit/level combinations. Returns a hierarchy_query_result object, which can be used like cuda::std::extents or dim3. Unit and Level need to be a core CUDA hierarchy levels, for example grid_level or block_level. This function is also available as a level type member function, in that case it only takes a unit argument.

Snippet

#include <cudax/hierarchy_dimensions.cuh>
#include <cassert>

using namespace cuda::experimental;

__global__ void kernel()
{
    // Can be called with the instances of level types
    auto thread_index_in_block = hierarchy::index(thread, block);
    assert(thread_index_in_block == threadIdx);
    auto block_index_in_grid = grid.index(block);
    assert(block_index_in_grid == blockIdx);

    // Or using the level types as template arguments
    auto thread_index_in_grid = hierarchy::index<thread_level, grid_level>();
}

Template Parameters

• Unit – Specifies the entity that the index is requested for

• Level – Specifies at what hierarchy level the index is requested