#ifndef GRAPH_BACKEND_DNNL_EXECUTABLES_GEN_INDEX_HPP
#define GRAPH_BACKEND_DNNL_EXECUTABLES_GEN_INDEX_HPP
#include "graph/backend/dnnl/executables/base.hpp"
namespace dnnl {
namespace impl {
namespace graph {
namespace dnnl_impl {
#if DNNL_GPU_RUNTIME != DNNL_RUNTIME_NONE \
&& DNNL_GPU_VENDOR == DNNL_VENDOR_INTEL
using namespace dnnl::impl::gpu::intel;
#define MAX_NDIMS 6
#endif
struct genindex_executable_t : public op_executable_t {
DECLARE_ARG_INDICES_GETTER;
genindex_executable_t(std::shared_ptr<op_t> &op,
const dnnl::engine &p_engine, pd_cache_t &pd_cache,
const fpmath_t &fpmath, bool use_block_layout);
void execute(const stream &stream,
const std::unordered_map<int, memory> &args) const override;
void execute_impl(const stream &stream,
const std::unordered_map<int, memory> &args) const;
#ifdef DNNL_WITH_SYCL
std::optional<::sycl::event> execute_sycl(const stream &stream,
const std::unordered_map<int, memory> &args,
const std::vector<::sycl::event> &deps) const override;
std::optional<::sycl::event> execute_sycl_impl(const stream &stream,
const std::unordered_map<int, memory> &args,
const std::vector<::sycl::event> &deps) const;
#endif
#if DNNL_GPU_RUNTIME == DNNL_RUNTIME_OCL
cl_event execute_ocl(const stream &stream,
const std::unordered_map<int, memory> &args,
const std::vector<cl_event> &deps) const override;
cl_event execute_ocl_impl(const stream &stream,
const std::unordered_map<int, memory> &args,
const std::vector<cl_event> &deps) const;
#endif
private:
int axis_, nelems_, ndims_;
dims_t output_dims_, output_strides_;
std::string info_;
#if (DNNL_GPU_RUNTIME != DNNL_RUNTIME_NONE) \
&& (DNNL_GPU_VENDOR == DNNL_VENDOR_INTEL)
compute::kernel_t kernel_;
#endif
};
} } } }
#endif