common.hpp

Go to the documentation of this file.

//==---------------- common.hpp - DPC++ Explicit SIMD API   ----------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// Common definitions used in experimental Explicit SIMD APIs.
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/ext/intel/esimd/common.hpp>
#include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>
#include <sycl/ext/intel/esimd/native/common.hpp>
#include <sycl/ext/intel/esimd/xmx/common.hpp>
 
#include <cstdint>
#include <type_traits>
 
namespace sycl {
inline namespace _V1 {
namespace ext::intel::experimental::esimd {
 
 
enum class __SYCL_DEPRECATED(
    "use sycl::ext::intel::esimd::fence_scope") lsc_scope : uint8_t {
  group = 0,  
  local = 1,  
  tile = 2,   
  gpu = 3,    
  gpus = 4,   
  system = 5, 
  sysacq = 6, 
};
 
enum class __SYCL_DEPRECATED("use sycl::ext::intel::esimd::fence_flush_op")
    lsc_fence_op : uint8_t {
      none = 0,       
      evict = 1,      
      invalidate = 2, 
      discard = 3,    
      clean = 4,   
      flushl3 = 5, 
    };
 
enum class __SYCL_DEPRECATED("use sycl::ext::intel::esimd::memory_kind")
    lsc_memory_kind : uint8_t {
      untyped_global = 0,         
      untyped_global_low_pri = 1, 
      typed_global = 2,           
      shared_local = 3,           
    };
 
using lsc_data_size = __ESIMD_DNS::lsc_data_size;
 
namespace detail {
 
using lsc_vector_size = __ESIMD_DNS::lsc_vector_size;
 
using lsc_data_order = __ESIMD_DNS::lsc_data_order;
 
template <lsc_vector_size VS> constexpr void check_lsc_vector_size() {
  __ESIMD_DNS::check_lsc_vector_size<VS>();
}
 
template <int VS> constexpr void check_lsc_vector_size() {
  __ESIMD_DNS::check_lsc_vector_size<VS>();
}
 
template <typename T, lsc_data_size DS> constexpr void check_lsc_data_size() {
  __ESIMD_DNS::check_lsc_data_size<T, DS>();
}
 
template <lsc_vector_size VS> constexpr uint8_t to_int() {
  return __ESIMD_DNS::to_int<VS>();
}
 
template <int VS> constexpr lsc_vector_size to_lsc_vector_size() {
  return __ESIMD_DNS::to_lsc_vector_size<VS>();
}
 
template <typename T, lsc_data_size DS>
constexpr lsc_data_size finalize_data_size() {
  return __ESIMD_DNS::finalize_data_size<T, DS>();
}
 
constexpr lsc_data_size expand_data_size(lsc_data_size DS) {
  return __ESIMD_DNS::expand_data_size(DS);
}
 
template <typename T> struct lsc_expand_type {
  using type = __ESIMD_DNS::lsc_expand_type<T>::type;
};
 
} // namespace detail
 
using cache_hint = __ESIMD_NS::cache_hint;
 
enum class split_barrier_action : uint8_t {
  wait = 0,   // split barrier wait
  signal = 1, // split barrier signal
};
 
 
} // namespace ext::intel::experimental::esimd
} // namespace _V1
} // namespace sycl