common.hpp

Go to the documentation of this file.

/*******************************************************************************
* Copyright (c) 2022-2023 Intel Corporation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
*     http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*******************************************************************************/
 
 
#pragma once
 
#include <CL/sycl.hpp>
#include <ext/intel/esimd.hpp>
 
template <class T>
using remove_const_t = typename std::remove_const<T>::type;
 
 
#define KERNEL_MAIN SYCL_ESIMD_KERNEL
 
#define KERNEL_FUNC SYCL_ESIMD_FUNCTION
 
 
#define __XETLA_API inline
 
#ifndef __ESIMD_ENS
#define __ESIMD_ENS sycl::ext::intel::experimental::esimd
#endif
 
#ifndef __ESIMD_NS
#define __ESIMD_NS sycl::ext::intel::esimd
#endif
 
#define XETLA_MARKER(message) [[deprecated(message)]]
#define XETLA_WARNING(msg) __SYCL_WARNING(msg)
 
template <auto val>
XETLA_MARKER("Help function to print value")
inline constexpr void XETLA_PRINT() {}
template <typename type>
XETLA_MARKER("Help function to print type")
inline constexpr void XETLA_PRINT() {}
 
__XETLA_API int32_t xetla_get_hw_thread_id() {
    return __ESIMD_ENS::get_hw_thread_id();
}
 
__XETLA_API int32_t xetla_get_subdevice_id() {
    return __ESIMD_ENS::get_subdevice_id();
}
 
namespace gpu::xetla {
 
enum class gpu_arch : uint8_t { Xe = 0 };
enum class grf_mode : uint8_t { normal = 0, double_grf = 1 };
 
enum class mem_layout : uint8_t { row_major = 0, col_major = 1 };
enum class mem_space : uint8_t { global = 0, local = 1 };
enum class msg_type : uint8_t {
    block_2d = 0,
    block_1d = 1,
    scatter = 2,
    atomic_add = 3,
    unaligned_2d = 4
    // prefetch_2d = 4,
    // prefetch_1d = 5
};
 
enum class cache_hint : uint8_t {
    none = 0,
    uncached = 1,
    cached = 2,
    write_back = 3,
    write_through = 4,
    streaming = 5,
    read_invalidate = 6
};
 
enum class data_size : uint8_t {
    default_size = 0,
    u8 = 1,
    u16 = 2,
    u32 = 3,
    u64 = 4,
    u8u32 = 5, 
    u16u32 = 6, 
    u16u32h = 7, 
};
 
enum class memory_kind : uint8_t {
    untyped_global = 0, 
    untyped_global_low_pri = 1, 
    typed_global = 2, 
    shared_local = 3, 
};
 
enum class fence_op : uint8_t {
    none = 0, 
    evict = 1, 
    invalidate = 2, 
    discard = 3, 
    clean = 4, 
    flushl2 = 5, 
};
enum class fence_scope : uint8_t {
    group = 0, 
    local = 1, 
    tile = 2, 
    gpu = 3, 
    gpus = 4, 
    system = 5, 
    sysacq = 6, 
};
 
enum class atomic_op : uint8_t {
    iinc = 0x0,
    idec = 0x1,
    iadd = 0x2,
    isub = 0x3,
    smin = 0x4,
    smax = 0x5,
    cmpxchg = 0x6,
    fadd = 0x7,
    fsub = 0x8,
    fmin = 0x9,
    fmax = 0xa,
    fcmpxchg = 0xb,
    umin = 0xc,
    umax = 0xd,
    bit_and = 0xe,
    bit_or = 0xf,
    bit_xor = 0x10,
    load = 0x11,
    store = 0x12
};
 
enum class argument_type : uint8_t {
    U1 = 0, // unsigned 1 bit
    S1 = 1, // signed 1 bit
    U2 = 2, // unsigned 2 bits
    S2 = 3, // signed 2 bits
    U4 = 4, // unsigned 4 bits
    S4 = 5, // signed 4 bits
    U8 = 6, // unsigned 8 bits
    S8 = 7, // signed 8 bits
    BF16 = 8, // bfloat 16
    FP16 = 9, // half float
    TF32 = 12, // tensorfloat 32
    DF = 13, // double (64bits)
    NUM_ARG_TYPES = 14
};
 
// Saturation tag
class xetla_saturation_on_tag {
public:
    using sat_tag = typename __ESIMD_NS::saturation_on_tag;
    static constexpr sat_tag value = {};
};
 
class xetla_saturation_off_tag {
public:
    using sat_tag = typename __ESIMD_NS::saturation_off_tag;
    static constexpr sat_tag value = {};
};
 
template <typename T>
using is_xetla_scalar = typename __ESIMD_DNS::is_esimd_scalar<T>;
 
enum class reduce_op : uint8_t {
    sum = 0, // performance reduce_sum
    prod = 1, // performance reduce_prod
    min = 2, // performance reduce_min
    max = 3, // performance reduce_max
};
 
 
#define SW_BARRIER() __ESIMD_NS::fence<__ESIMD_NS::fence_mask::sw_barrier>()
 
__XETLA_API void xetla_wait(uint16_t val) {
    __ESIMD_ENS::wait(__ESIMD_NS::simd<uint16_t, 1>(val));
}
 
} // namespace gpu::xetla