gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ > Struct Template Reference

GROUP_SWIZZLE implementation of snake curve. More...

#include <dispatch_policy.hpp>

Public Member Functions
	group_swizzle_snake ()=default

Static Public Member Functions
template<int idx>
static __XETLA_API std::enable_if_t< idx==0, int >	get_tile_idx (sycl::nd_item< 3 > &item)
template<int idx>
static __XETLA_API std::enable_if_t< idx==1, int >	get_tile_idx (sycl::nd_item< 3 > &item)
template<int idx>
static __XETLA_API std::enable_if_t< idx==2, int >	get_tile_idx (sycl::nd_item< 3 > &item)
static __XETLA_API void	update_group_range (uint32_t &group_range_m, uint32_t &group_range_n)

Static Public Attributes
static constexpr gpu_arch	arch_tag = arch_tag_

Detailed Description

template<int wg_num_n_, gpu_arch arch_tag_>
struct gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >

GROUP_SWIZZLE implementation of snake curve.

A GROUP_SWIZZLE implementation to remap linear workgroup id to a 2d coordination in snake order.

Template Parameters

wg_num_n_	Is the number of workgroup in horizontal direction, given by users.
arch_tag_	Is the HW architecture.

Constructor & Destructor Documentation

group_swizzle_snake()

template<int wg_num_n_, gpu_arch arch_tag_>

gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >::group_swizzle_snake ( )

inlinedefault

Member Function Documentation

get_tile_idx() [1/3]

template<int wg_num_n_, gpu_arch arch_tag_>

template<int idx>

static __XETLA_API std::enable_if_t< idx==0, int > gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >::get_tile_idx ( sycl::nd_item< 3 > &  item )

inlinestatic

get_tile_idx() [2/3]

template<int wg_num_n_, gpu_arch arch_tag_>

template<int idx>

static __XETLA_API std::enable_if_t< idx==1, int > gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >::get_tile_idx ( sycl::nd_item< 3 > &  item )

inlinestatic

get_tile_idx() [3/3]

template<int wg_num_n_, gpu_arch arch_tag_>

template<int idx>

static __XETLA_API std::enable_if_t< idx==2, int > gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >::get_tile_idx ( sycl::nd_item< 3 > &  item )

inlinestatic

update_group_range()

template<int wg_num_n_, gpu_arch arch_tag_>

static __XETLA_API void gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >::update_group_range ( uint32_t &  group_range_m, uint32_t &  group_range_n  )

inlinestatic

Member Data Documentation

arch_tag

template<int wg_num_n_, gpu_arch arch_tag_>

constexpr gpu_arch gpu::xetla::kernel::group_swizzle_snake< wg_num_n_, arch_tag_ >::arch_tag = arch_tag_

staticconstexpr