// Copyright 2009-2021 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "platform.h"
using sycl::float16;
using sycl::float8;
using sycl::float4;
using sycl::float3;
using sycl::float2;
using sycl::int16;
using sycl::int8;
using sycl::int4;
using sycl::int3;
using sycl::int2;
using sycl::uint16;
using sycl::uint8;
using sycl::uint4;
using sycl::uint3;
using sycl::uint2;
using sycl::uchar16;
using sycl::uchar8;
using sycl::uchar4;
using sycl::uchar3;
using sycl::uchar2;
using sycl::ushort16;
using sycl::ushort8;
using sycl::ushort4;
using sycl::ushort3;
using sycl::ushort2;
#ifdef __SYCL_DEVICE_ONLY__
#define GLOBAL __attribute__((opencl_global))
#define LOCAL __attribute__((opencl_local))
SYCL_EXTERNAL extern int work_group_reduce_add(int x);
SYCL_EXTERNAL extern float work_group_reduce_min(float x);
SYCL_EXTERNAL extern float work_group_reduce_max(float x);
SYCL_EXTERNAL extern float atomic_min(volatile GLOBAL float *p, float val);
SYCL_EXTERNAL extern float atomic_min(volatile LOCAL float *p, float val);
SYCL_EXTERNAL extern float atomic_max(volatile GLOBAL float *p, float val);
SYCL_EXTERNAL extern float atomic_max(volatile LOCAL float *p, float val);
SYCL_EXTERNAL extern "C" unsigned int intel_sub_group_ballot(bool valid);
SYCL_EXTERNAL extern "C" void __builtin_IB_assume_uniform(void *p);
// Load message caching control
enum LSC_LDCC {
LSC_LDCC_DEFAULT,
LSC_LDCC_L1UC_L3UC, // Override to L1 uncached and L3 uncached
LSC_LDCC_L1UC_L3C, // Override to L1 uncached and L3 cached
LSC_LDCC_L1C_L3UC, // Override to L1 cached and L3 uncached
LSC_LDCC_L1C_L3C, // Override to L1 cached and L3 cached
LSC_LDCC_L1S_L3UC, // Override to L1 streaming load and L3 uncached
LSC_LDCC_L1S_L3C, // Override to L1 streaming load and L3 cached
LSC_LDCC_L1IAR_L3C, // Override to L1 invalidate-after-read, and L3 cached
};
// Store message caching control (also used for atomics)
enum LSC_STCC {
LSC_STCC_DEFAULT,
LSC_STCC_L1UC_L3UC, // Override to L1 uncached and L3 uncached
LSC_STCC_L1UC_L3WB, // Override to L1 uncached and L3 written back
LSC_STCC_L1WT_L3UC, // Override to L1 written through and L3 uncached
LSC_STCC_L1WT_L3WB, // Override to L1 written through and L3 written back
LSC_STCC_L1S_L3UC, // Override to L1 streaming and L3 uncached
LSC_STCC_L1S_L3WB, // Override to L1 streaming and L3 written back
LSC_STCC_L1WB_L3WB, // Override to L1 written through and L3 written back
};
///////////////////////////////////////////////////////////////////////
// LSC Loads
///////////////////////////////////////////////////////////////////////
SYCL_EXTERNAL /* extern "C" */ uint32_t __builtin_IB_lsc_load_global_uchar_to_uint (const GLOBAL uint8_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D8U32
SYCL_EXTERNAL /* extern "C" */ uint32_t __builtin_IB_lsc_load_global_ushort_to_uint(const GLOBAL uint16_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D16U32
SYCL_EXTERNAL /* extern "C" */ uint32_t __builtin_IB_lsc_load_global_uint (const GLOBAL uint32_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V1
SYCL_EXTERNAL /* extern "C" */ sycl::uint2 __builtin_IB_lsc_load_global_uint2 (const GLOBAL sycl::uint2 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V2
SYCL_EXTERNAL /* extern "C" */ sycl::uint3 __builtin_IB_lsc_load_global_uint3 (const GLOBAL sycl::uint3 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V3
SYCL_EXTERNAL /* extern "C" */ sycl::uint4 __builtin_IB_lsc_load_global_uint4 (const GLOBAL sycl::uint4 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V4
SYCL_EXTERNAL /* extern "C" */ sycl::uint8 __builtin_IB_lsc_load_global_uint8 (const GLOBAL sycl::uint8 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V8
SYCL_EXTERNAL /* extern "C" */ uint64_t __builtin_IB_lsc_load_global_ulong (const GLOBAL uint64_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V1
SYCL_EXTERNAL /* extern "C" */ sycl::ulong2 __builtin_IB_lsc_load_global_ulong2 (const GLOBAL sycl::ulong2 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V2
SYCL_EXTERNAL /* extern "C" */ sycl::ulong3 __builtin_IB_lsc_load_global_ulong3 (const GLOBAL sycl::ulong3 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V3
SYCL_EXTERNAL /* extern "C" */ sycl::ulong4 __builtin_IB_lsc_load_global_ulong4 (const GLOBAL sycl::ulong4 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V4
SYCL_EXTERNAL /* extern "C" */ sycl::ulong8 __builtin_IB_lsc_load_global_ulong8 (const GLOBAL sycl::ulong8 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V8
// global address space
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uchar_from_uint (GLOBAL uint8_t *base, int immElemOff, uint32_t val, enum LSC_STCC cacheOpt); //D8U32
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ushort_from_uint(GLOBAL uint16_t *base, int immElemOff, uint32_t val, enum LSC_STCC cacheOpt); //D16U32
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint (GLOBAL uint32_t *base, int immElemOff, uint32_t val, enum LSC_STCC cacheOpt); //D32V1
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint2 (GLOBAL sycl::uint2 *base, int immElemOff, sycl::uint2 val, enum LSC_STCC cacheOpt); //D32V2
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint3 (GLOBAL sycl::uint3 *base, int immElemOff, sycl::uint3 val, enum LSC_STCC cacheOpt); //D32V3
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint4 (GLOBAL sycl::uint4 *base, int immElemOff, sycl::uint4 val, enum LSC_STCC cacheOpt); //D32V4
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint8 (GLOBAL sycl::uint8 *base, int immElemOff, sycl::uint8 val, enum LSC_STCC cacheOpt); //D32V8
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong (GLOBAL uint64_t *base, int immElemOff, uint64_t val, enum LSC_STCC cacheOpt); //D64V1
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong2 (GLOBAL sycl::ulong2 *base, int immElemOff, sycl::ulong2 val, enum LSC_STCC cacheOpt); //D64V2
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong3 (GLOBAL sycl::ulong3 *base, int immElemOff, sycl::ulong3 val, enum LSC_STCC cacheOpt); //D64V3
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong4 (GLOBAL sycl::ulong4 *base, int immElemOff, sycl::ulong4 val, enum LSC_STCC cacheOpt); //D64V4
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong8 (GLOBAL sycl::ulong8 *base, int immElemOff, sycl::ulong8 val, enum LSC_STCC cacheOpt); //D64V8
///////////////////////////////////////////////////////////////////////
// prefetching
///////////////////////////////////////////////////////////////////////
//
// LSC Pre-Fetch Load functions with CacheControls
// global address space
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uchar (const GLOBAL uint8_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D8U32
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ushort(const GLOBAL uint16_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D16U32
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint (const GLOBAL uint32_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V1
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint2 (const GLOBAL sycl::uint2 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V2
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint3 (const GLOBAL sycl::uint3 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V3
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint4 (const GLOBAL sycl::uint4 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V4
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint8 (const GLOBAL sycl::uint8 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V8
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong (const GLOBAL uint64_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V1
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong2(const GLOBAL sycl::ulong2 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V2
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong3(const GLOBAL sycl::ulong3 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V3
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong4(const GLOBAL sycl::ulong4 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V4
SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong8(const GLOBAL sycl::ulong8 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V8
#else
#define GLOBAL
#define LOCAL
/* dummy functions for host */
inline int work_group_reduce_add(int x) { return x; }
inline float work_group_reduce_min(float x) { return x; }
inline float work_group_reduce_max(float x) { return x; }
inline float atomic_min(volatile float *p, float val) { return val; };
inline float atomic_max(volatile float *p, float val) { return val; };
inline uint32_t intel_sub_group_ballot(bool valid) { return 0; }
#endif
/* creates a temporary that is enforced to be uniform */
#define SYCL_UNIFORM_VAR(Ty,tmp,k) \
Ty tmp##_data; \
Ty* p##tmp##_data = (Ty*) sub_group_broadcast((uint64_t)&tmp##_data,k); \
Ty& tmp = *p##tmp##_data;
#if !defined(__forceinline)
#define __forceinline inline __attribute__((always_inline))
#endif
#if __SYCL_COMPILER_VERSION < 20210801
#define all_of_group all_of
#define any_of_group any_of
#define none_of_group none_of
#define group_broadcast broadcast
#define reduce_over_group reduce
#define exclusive_scan_over_group exclusive_scan
#define inclusive_scan_over_group inclusive_scan
#endif
namespace embree
{
template<typename T>
__forceinline T cselect(const bool mask, const T &a, const T &b)
{
return sycl::select(b,a,(int)mask);
}
template<typename T, typename M>
__forceinline T cselect(const M &mask, const T &a, const T &b)
{
return sycl::select(b,a,mask);
}
__forceinline const sycl::sub_group this_sub_group() {
return sycl::ext::oneapi::experimental::this_sub_group();
}
__forceinline const uint32_t get_sub_group_local_id() {
return this_sub_group().get_local_id()[0];
}
__forceinline const uint32_t get_sub_group_size() {
return this_sub_group().get_max_local_range().size();
}
__forceinline const uint32_t get_sub_group_id() {
return this_sub_group().get_group_id()[0];
}
__forceinline const uint32_t get_num_sub_groups() {
return this_sub_group().get_group_range().size();
}
__forceinline uint32_t sub_group_ballot(bool pred) {
return intel_sub_group_ballot(pred);
}
__forceinline bool sub_group_all_of(bool pred) {
return sycl::all_of_group(this_sub_group(),pred);
}
__forceinline bool sub_group_any_of(bool pred) {
return sycl::any_of_group(this_sub_group(),pred);
}
__forceinline bool sub_group_none_of(bool pred) {
return sycl::none_of_group(this_sub_group(),pred);
}
template <typename T> __forceinline T sub_group_broadcast(T x, sycl::id<1> local_id) {
return sycl::group_broadcast<sycl::sub_group>(this_sub_group(),x,local_id);
}
template <typename T> __forceinline T sub_group_make_uniform(T x) {
return sub_group_broadcast(x,sycl::ctz(intel_sub_group_ballot(true)));
}
__forceinline void assume_uniform_array(void* ptr) {
#ifdef __SYCL_DEVICE_ONLY__
__builtin_IB_assume_uniform(ptr);
#endif
}
template <typename T, class BinaryOperation> __forceinline T sub_group_reduce(T x, BinaryOperation binary_op) {
return sycl::reduce_over_group<sycl::sub_group>(this_sub_group(),x,binary_op);
}
template <typename T, class BinaryOperation> __forceinline T sub_group_reduce(T x, T init, BinaryOperation binary_op) {
return sycl::reduce_over_group<sycl::sub_group>(this_sub_group(),x,init,binary_op);
}
template <typename T> __forceinline T sub_group_reduce_min(T x, T init) {
return sub_group_reduce(x, init, sycl::ext::oneapi::minimum<T>());
}
template <typename T> __forceinline T sub_group_reduce_min(T x) {
return sub_group_reduce(x, sycl::ext::oneapi::minimum<T>());
}
template <typename T> __forceinline T sub_group_reduce_max(T x) {
return sub_group_reduce(x, sycl::ext::oneapi::maximum<T>());
}
template <typename T> __forceinline T sub_group_reduce_add(T x) {
return sub_group_reduce(x, sycl::ext::oneapi::plus<T>());
}
template <typename T, class BinaryOperation> __forceinline T sub_group_exclusive_scan(T x, BinaryOperation binary_op) {
return sycl::exclusive_scan_over_group(this_sub_group(),x,binary_op);
}
template <typename T, class BinaryOperation> __forceinline T sub_group_exclusive_scan_min(T x) {
return sub_group_exclusive_scan(x,sycl::ext::oneapi::minimum<T>());
}
template <typename T, class BinaryOperation> __forceinline T sub_group_exclusive_scan(T x, T init, BinaryOperation binary_op) {
return sycl::exclusive_scan_over_group(this_sub_group(),x,init,binary_op);
}
template <typename T, class BinaryOperation> __forceinline T sub_group_inclusive_scan(T x, BinaryOperation binary_op) {
return sycl::inclusive_scan_over_group(this_sub_group(),x,binary_op);
}
template <typename T, class BinaryOperation> __forceinline T sub_group_inclusive_scan(T x, BinaryOperation binary_op, T init) {
return sycl::inclusive_scan_over_group(this_sub_group(),x,binary_op,init);
}
template <typename T> __forceinline T sub_group_shuffle(T x, sycl::id<1> local_id) {
return this_sub_group().shuffle(x, local_id);
}
template <typename T> __forceinline T sub_group_shuffle_down(T x, uint32_t delta) {
return this_sub_group().shuffle_down(x, delta);
}
template <typename T> __forceinline T sub_group_shuffle_up(T x, uint32_t delta) {
return this_sub_group().shuffle_up(x, delta);
}
template <typename T> __forceinline T sub_group_load(const void* src) {
return this_sub_group().load(sycl::multi_ptr<T,sycl::access::address_space::global_space>((T*)src));
}
template <typename T> __forceinline void sub_group_store(void* dst, const T& x) {
this_sub_group().store(sycl::multi_ptr<T,sycl::access::address_space::global_space>((T*)dst),x);
}
}
#if __SYCL_COMPILER_VERSION < 20210801
#undef all_of_group
#undef any_of_group
#undef none_of_group
#undef group_broadcast
#undef reduce_over_group
#undef exclusive_scan_over_group
#undef inclusive_scan_over_group
#endif
Codebase Browser
godot