#include <cstdint>
#include <math.h>
#include "util/half.h"
#include "util/types.h"

Macros
#define	__KERNEL_GPU__

#define	__KERNEL_ONEAPI__

#define	__KERNEL_64_BIT__

#define	CCL_NAMESPACE_BEGIN

#define	CCL_NAMESPACE_END

#define	__NODES_MAX_GROUP__ NODE_GROUP_LEVEL_MAX

#define	__NODES_FEATURES__ NODE_FEATURE_ALL

#define	__device__

#define	ccl_device

#define	ccl_device_extern extern "C"

#define	ccl_global

#define	ccl_always_inline __attribute__((always_inline))

#define	ccl_device_inline inline

#define	ccl_noinline

#define	ccl_inline_constant const constexpr

#define	ccl_device_constant static constexpr

#define	ccl_static_constexpr static constexpr

#define	ccl_device_forceinline __attribute__((always_inline))

#define	ccl_device_noinline ccl_device ccl_noinline

#define	ccl_device_noinline_cpu ccl_device

#define	ccl_device_inline_method ccl_device

#define	ccl_restrict __restrict__

#define	ccl_loop_no_unroll

#define	ccl_optional_struct_init

#define	ccl_private

#define	ccl_ray_data ccl_private

#define	ccl_gpu_shared

#define	ATTR_FALLTHROUGH __attribute__((fallthrough))

#define	ccl_constant const

#define	ccl_try_align(...) __attribute__((aligned(__VA_ARGS__)))

#define	ccl_align(n) __attribute__((aligned(n)))

#define	kernel_assert(cond)

#define	ccl_may_alias

#define	ccl_gpu_kernel(block_num_threads, thread_num_registers)

#define	ccl_gpu_kernel_threads(block_num_threads)

#define	__ccl_gpu_kernel_signature(name, ...)

#define	ccl_gpu_kernel_signature __ccl_gpu_kernel_signature

#define	ccl_gpu_kernel_postfix

#define	ccl_gpu_kernel_call(x) ((ONEAPIKernelContext*)kg)->x

#define	ccl_gpu_kernel_within_bounds(i, n) ((i) < (n))

#define	ccl_gpu_kernel_lambda(func, ...)

#define	ccl_gpu_thread_idx_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_local_id(0))

#define	ccl_gpu_block_dim_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_local_range(0))

#define	ccl_gpu_block_idx_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_group(0))

#define	ccl_gpu_grid_dim_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_group_range(0))

#define	ccl_gpu_warp_size (sycl::ext::oneapi::experimental::this_sub_group().get_local_range()[0])

#define	ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))

#define	ccl_gpu_global_id_x() (sycl::ext::oneapi::experimental::this_nd_item<1>().get_global_id(0))

#define	ccl_gpu_global_size_x() (sycl::ext::oneapi::experimental::this_nd_item<1>().get_global_range(0))

#define	ccl_gpu_syncthreads() sycl::ext::oneapi::experimental::this_nd_item<1>().barrier()

#define	ccl_gpu_local_syncthreads() sycl::ext::oneapi::experimental::this_nd_item<1>().barrier(sycl::access::fence_space::local_space)

#define	ccl_gpu_ballot(predicate) (predicate ? 1 : 0)

#define	CCL_ONEAPI_CONSTANT

#define	sycl_printf(format, ...)

#define	sycl_printf_(format)

#define	fabsf(x) sycl::fabs((x))

#define	copysignf(x, y) sycl::copysign((x), (y))

#define	asinf(x) sycl::asin((x))

#define	acosf(x) sycl::acos((x))

#define	atanf(x) sycl::atan((x))

#define	floorf(x) sycl::floor((x))

#define	ceilf(x) sycl::ceil((x))

#define	sinhf(x) sycl::sinh((x))

#define	coshf(x) sycl::cosh((x))

#define	tanhf(x) sycl::tanh((x))

#define	hypotf(x, y) sycl::hypot((x), (y))

#define	atan2f(x, y) sycl::atan2((x), (y))

#define	fmaxf(x, y) sycl::fmax((x), (y))

#define	fminf(x, y) sycl::fmin((x), (y))

#define	fmodf(x, y) sycl::fmod((x), (y))

#define	lgammaf(x) sycl::lgamma((x))

#define	cosf(x) sycl::native::cos(((float)(x)))

#define	sinf(x) sycl::native::sin(((float)(x)))

#define	powf(x, y) sycl::native::powr(((float)(x)), ((float)(y)))

#define	tanf(x) sycl::native::tan(((float)(x)))

#define	logf(x) sycl::native::log(((float)(x)))

#define	expf(x) sycl::native::exp(((float)(x)))

#define	__forceinline __attribute__((always_inline))

Typedefs
using	uchar = unsigned char

Functions
ccl_device_forceinline float	__uint_as_float (unsigned int x)

ccl_device_forceinline unsigned int	__float_as_uint (float x)

ccl_device_forceinline float	__int_as_float (int x)

ccl_device_forceinline int	__float_as_int (float x)

Macro Definition Documentation

◆ __ccl_gpu_kernel_signature

#define __ccl_gpu_kernel_signature	(		name,
			... )

Value:

void oneapi_kernel_##name(KernelGlobalsGPU *ccl_restrict kg, \
                          size_t kernel_global_size, \
                          size_t kernel_local_size, \
                          sycl::handler &cgh, \
                          __VA_ARGS__) { \
      (kg); \
      cgh.parallel_for( \
          sycl::nd_range<1>(kernel_global_size, kernel_local_size), \
          [=](sycl::nd_item<1> item) {

Definition at line 69 of file device/oneapi/compat.h.

◆ device

#define __device__

Definition at line 32 of file device/oneapi/compat.h.

◆ __forceinline

#define __forceinline __attribute__((always_inline))

Definition at line 234 of file device/oneapi/compat.h.

Referenced by ccl_try_align(), ccl_try_align(), and ccl_try_align().

◆ __KERNEL_64_BIT__

#define __KERNEL_64_BIT__

Definition at line 9 of file device/oneapi/compat.h.

◆ __KERNEL_GPU__

#define __KERNEL_GPU__

Definition at line 7 of file device/oneapi/compat.h.

◆ __KERNEL_ONEAPI__

#define __KERNEL_ONEAPI__

Definition at line 8 of file device/oneapi/compat.h.

◆ __NODES_FEATURES__

#define __NODES_FEATURES__ NODE_FEATURE_ALL

Definition at line 25 of file device/oneapi/compat.h.

◆ __NODES_MAX_GROUP__

#define __NODES_MAX_GROUP__ NODE_GROUP_LEVEL_MAX

Definition at line 22 of file device/oneapi/compat.h.

◆ acosf

#define acosf ( x ) sycl::acos((x))

Definition at line 213 of file device/oneapi/compat.h.

◆ asinf

#define asinf ( x ) sycl::asin((x))

Definition at line 212 of file device/oneapi/compat.h.

◆ atan2f

#define atan2f	(		x,
			y ) sycl::atan2((x), (y))

Definition at line 221 of file device/oneapi/compat.h.

◆ atanf

#define atanf ( x ) sycl::atan((x))

Definition at line 214 of file device/oneapi/compat.h.

◆ ATTR_FALLTHROUGH

#define ATTR_FALLTHROUGH __attribute__((fallthrough))

Definition at line 55 of file device/oneapi/compat.h.

◆ ccl_align

#define ccl_align ( n ) __attribute__((aligned(n)))

Definition at line 58 of file device/oneapi/compat.h.

◆ ccl_always_inline

#define ccl_always_inline __attribute__((always_inline))

Definition at line 39 of file device/oneapi/compat.h.

◆ ccl_constant

#define ccl_constant const

Definition at line 56 of file device/oneapi/compat.h.

◆ ccl_device

#define ccl_device

Definition at line 36 of file device/oneapi/compat.h.

◆ ccl_device_constant

#define ccl_device_constant static constexpr

Definition at line 43 of file device/oneapi/compat.h.

◆ ccl_device_extern

#define ccl_device_extern extern "C"

Definition at line 37 of file device/oneapi/compat.h.

◆ ccl_device_forceinline

#define ccl_device_forceinline __attribute__((always_inline))

Definition at line 45 of file device/oneapi/compat.h.

◆ ccl_device_inline

#define ccl_device_inline inline

Definition at line 40 of file device/oneapi/compat.h.

◆ ccl_device_inline_method

#define ccl_device_inline_method ccl_device

Definition at line 48 of file device/oneapi/compat.h.

◆ ccl_device_noinline

#define ccl_device_noinline ccl_device ccl_noinline

Definition at line 46 of file device/oneapi/compat.h.

◆ ccl_device_noinline_cpu

#define ccl_device_noinline_cpu ccl_device

Definition at line 47 of file device/oneapi/compat.h.

◆ ccl_global

#define ccl_global

Definition at line 38 of file device/oneapi/compat.h.

◆ ccl_gpu_ballot

#define ccl_gpu_ballot ( predicate ) (predicate ? 1 : 0)

Definition at line 145 of file device/oneapi/compat.h.

◆ ccl_gpu_block_dim_x

#define ccl_gpu_block_dim_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_local_range(0))

Definition at line 130 of file device/oneapi/compat.h.

◆ ccl_gpu_block_idx_x

#define ccl_gpu_block_idx_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_group(0))

Definition at line 131 of file device/oneapi/compat.h.

◆ ccl_gpu_global_id_x

#define ccl_gpu_global_id_x ( ) (sycl::ext::oneapi::experimental::this_nd_item<1>().get_global_id(0))

Definition at line 136 of file device/oneapi/compat.h.

◆ ccl_gpu_global_size_x

#define ccl_gpu_global_size_x ( ) (sycl::ext::oneapi::experimental::this_nd_item<1>().get_global_range(0))

Definition at line 137 of file device/oneapi/compat.h.

◆ ccl_gpu_grid_dim_x

#define ccl_gpu_grid_dim_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_group_range(0))

Definition at line 132 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel

#define ccl_gpu_kernel	(		block_num_threads,
			thread_num_registers )

Definition at line 65 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel_call

#define ccl_gpu_kernel_call ( x ) ((ONEAPIKernelContext*)kg)->x

Definition at line 114 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel_lambda

#define ccl_gpu_kernel_lambda	(		func,
			... )

Value:

  struct KernelLambda \
  { \
    KernelLambda(const ONEAPIKernelContext *_kg) : kg(_kg) {} \
    ccl_private const ONEAPIKernelContext *kg; \
    __VA_ARGS__; \
    int operator()(const int state) const { return (func); } \
  } ccl_gpu_kernel_lambda_pass((ONEAPIKernelContext *)kg)

Definition at line 117 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel_postfix

#define ccl_gpu_kernel_postfix

Value:

}); \

}

Definition at line 82 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel_signature

#define ccl_gpu_kernel_signature __ccl_gpu_kernel_signature

Definition at line 80 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel_threads

#define ccl_gpu_kernel_threads ( block_num_threads )

Definition at line 66 of file device/oneapi/compat.h.

◆ ccl_gpu_kernel_within_bounds

#define ccl_gpu_kernel_within_bounds	(		i,
			n ) ((i) < (n))

Definition at line 115 of file device/oneapi/compat.h.

◆ ccl_gpu_local_syncthreads

#define ccl_gpu_local_syncthreads ( ) sycl::ext::oneapi::experimental::this_nd_item<1>().barrier(sycl::access::fence_space::local_space)

Definition at line 141 of file device/oneapi/compat.h.

Referenced by gpu_parallel_active_index_array_impl().

◆ ccl_gpu_shared

#define ccl_gpu_shared

Definition at line 54 of file device/oneapi/compat.h.

◆ ccl_gpu_syncthreads

#define ccl_gpu_syncthreads ( ) sycl::ext::oneapi::experimental::this_nd_item<1>().barrier()

Definition at line 140 of file device/oneapi/compat.h.

◆ ccl_gpu_thread_idx_x

#define ccl_gpu_thread_idx_x (sycl::ext::oneapi::experimental::this_nd_item<1>().get_local_id(0))

Definition at line 129 of file device/oneapi/compat.h.

◆ ccl_gpu_thread_mask

#define ccl_gpu_thread_mask ( thread_warp ) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))

Definition at line 134 of file device/oneapi/compat.h.

◆ ccl_gpu_warp_size

#define ccl_gpu_warp_size (sycl::ext::oneapi::experimental::this_sub_group().get_local_range()[0])

Definition at line 133 of file device/oneapi/compat.h.

◆ ccl_inline_constant

#define ccl_inline_constant const constexpr

Definition at line 42 of file device/oneapi/compat.h.

◆ ccl_loop_no_unroll

#define ccl_loop_no_unroll

Definition at line 50 of file device/oneapi/compat.h.

◆ ccl_may_alias

#define ccl_may_alias

Definition at line 60 of file device/oneapi/compat.h.

◆ CCL_NAMESPACE_BEGIN

#define CCL_NAMESPACE_BEGIN

Definition at line 15 of file device/oneapi/compat.h.

◆ CCL_NAMESPACE_END

#define CCL_NAMESPACE_END

Definition at line 16 of file device/oneapi/compat.h.

◆ ccl_noinline

#define ccl_noinline

Definition at line 41 of file device/oneapi/compat.h.

◆ CCL_ONEAPI_CONSTANT

#define CCL_ONEAPI_CONSTANT

Definition at line 167 of file device/oneapi/compat.h.

◆ ccl_optional_struct_init

#define ccl_optional_struct_init

Definition at line 51 of file device/oneapi/compat.h.

◆ ccl_private

#define ccl_private

Definition at line 52 of file device/oneapi/compat.h.

◆ ccl_ray_data

#define ccl_ray_data ccl_private

Definition at line 53 of file device/oneapi/compat.h.

◆ ccl_restrict

#define ccl_restrict __restrict__

Definition at line 49 of file device/oneapi/compat.h.

◆ ccl_static_constexpr

#define ccl_static_constexpr static constexpr

Definition at line 44 of file device/oneapi/compat.h.

◆ ccl_try_align

#define ccl_try_align ( ... ) __attribute__((aligned(__VA_ARGS__)))

Definition at line 57 of file device/oneapi/compat.h.

◆ ceilf

#define ceilf ( x ) sycl::ceil((x))

Definition at line 216 of file device/oneapi/compat.h.

◆ copysignf

#define copysignf	(		x,
			y ) sycl::copysign((x), (y))

Definition at line 211 of file device/oneapi/compat.h.

◆ cosf

#define cosf ( x ) sycl::native::cos(((float)(x)))

Definition at line 227 of file device/oneapi/compat.h.

◆ coshf

#define coshf ( x ) sycl::cosh((x))

Definition at line 218 of file device/oneapi/compat.h.

◆ expf

#define expf ( x ) sycl::native::exp(((float)(x)))

Definition at line 232 of file device/oneapi/compat.h.

◆ fabsf

#define fabsf ( x ) sycl::fabs((x))

Definition at line 210 of file device/oneapi/compat.h.

◆ floorf

#define floorf ( x ) sycl::floor((x))

Definition at line 215 of file device/oneapi/compat.h.

◆ fmaxf

#define fmaxf	(		x,
			y ) sycl::fmax((x), (y))

Definition at line 222 of file device/oneapi/compat.h.

◆ fminf

#define fminf	(		x,
			y ) sycl::fmin((x), (y))

Definition at line 223 of file device/oneapi/compat.h.

◆ fmodf

#define fmodf	(		x,
			y ) sycl::fmod((x), (y))

Definition at line 224 of file device/oneapi/compat.h.

◆ hypotf

#define hypotf	(		x,
			y ) sycl::hypot((x), (y))

Definition at line 220 of file device/oneapi/compat.h.

◆ kernel_assert

#define kernel_assert ( cond )

Definition at line 59 of file device/oneapi/compat.h.

◆ lgammaf

#define lgammaf ( x ) sycl::lgamma((x))

Definition at line 225 of file device/oneapi/compat.h.

Referenced by beta().

◆ logf

#define logf ( x ) sycl::native::log(((float)(x)))

Definition at line 231 of file device/oneapi/compat.h.

◆ powf

#define powf	(		x,
			y ) sycl::native::powr(((float)(x)), ((float)(y)))

Definition at line 229 of file device/oneapi/compat.h.

◆ sinf

#define sinf ( x ) sycl::native::sin(((float)(x)))

Definition at line 228 of file device/oneapi/compat.h.

◆ sinhf

#define sinhf ( x ) sycl::sinh((x))

Definition at line 217 of file device/oneapi/compat.h.

◆ sycl_printf

#define sycl_printf	(		format,
			... )

Value:

    {               \
    static const CCL_ONEAPI_CONSTANT char fmt[] = format;          \
    sycl::ext::oneapi::experimental::printf(fmt, __VA_ARGS__ );    \
  }

Definition at line 170 of file device/oneapi/compat.h.

◆ sycl_printf_

#define sycl_printf_ ( format )

Value:

    {               \
    static const CCL_ONEAPI_CONSTANT char fmt[] = format;          \
    sycl::ext::oneapi::experimental::printf(fmt);                  \
  }

Definition at line 175 of file device/oneapi/compat.h.

◆ tanf

#define tanf ( x ) sycl::native::tan(((float)(x)))

Definition at line 230 of file device/oneapi/compat.h.

◆ tanhf

#define tanhf ( x ) sycl::tanh((x))

Definition at line 219 of file device/oneapi/compat.h.

Typedef Documentation

◆ uchar

using uchar = unsigned char

Definition at line 189 of file device/oneapi/compat.h.

Function Documentation

◆ __float_as_int()

ccl_device_forceinline int __float_as_int ( float x )

Definition at line 205 of file device/oneapi/compat.h.

◆ __float_as_uint()

ccl_device_forceinline unsigned int __float_as_uint ( float x )

Definition at line 197 of file device/oneapi/compat.h.

◆ __int_as_float()

ccl_device_forceinline float __int_as_float ( int x )

Definition at line 201 of file device/oneapi/compat.h.

◆ __uint_as_float()

ccl_device_forceinline float __uint_as_float ( unsigned int x )

Definition at line 193 of file device/oneapi/compat.h.

Macros

Typedefs

Functions

Macro Definition Documentation

◆ __ccl_gpu_kernel_signature

◆ __device__

◆ __forceinline

◆ __KERNEL_64_BIT__

◆ __KERNEL_GPU__

◆ __KERNEL_ONEAPI__

◆ __NODES_FEATURES__

◆ __NODES_MAX_GROUP__

◆ acosf

◆ asinf

◆ atan2f

◆ atanf

◆ ATTR_FALLTHROUGH

◆ ccl_align

◆ ccl_always_inline

◆ ccl_constant

◆ ccl_device

◆ ccl_device_constant

◆ ccl_device_extern

◆ ccl_device_forceinline

◆ ccl_device_inline

◆ ccl_device_inline_method

◆ ccl_device_noinline

◆ ccl_device_noinline_cpu

◆ ccl_global

◆ ccl_gpu_ballot

◆ ccl_gpu_block_dim_x

◆ ccl_gpu_block_idx_x

◆ ccl_gpu_global_id_x

◆ ccl_gpu_global_size_x

◆ ccl_gpu_grid_dim_x

◆ ccl_gpu_kernel

◆ ccl_gpu_kernel_call

◆ ccl_gpu_kernel_lambda

◆ ccl_gpu_kernel_postfix

◆ ccl_gpu_kernel_signature

◆ ccl_gpu_kernel_threads

◆ ccl_gpu_kernel_within_bounds

◆ ccl_gpu_local_syncthreads

◆ ccl_gpu_shared

◆ ccl_gpu_syncthreads

◆ ccl_gpu_thread_idx_x

◆ ccl_gpu_thread_mask

◆ ccl_gpu_warp_size

◆ ccl_inline_constant

◆ ccl_loop_no_unroll

◆ ccl_may_alias

◆ CCL_NAMESPACE_BEGIN

◆ CCL_NAMESPACE_END

◆ ccl_noinline

◆ CCL_ONEAPI_CONSTANT

◆ ccl_optional_struct_init

◆ ccl_private

◆ ccl_ray_data

◆ ccl_restrict

◆ ccl_static_constexpr

◆ ccl_try_align

◆ ceilf

◆ copysignf

◆ cosf

◆ coshf

◆ expf

◆ fabsf

◆ floorf

◆ fmaxf

◆ fminf

◆ fmodf

◆ hypotf

◆ kernel_assert

◆ lgammaf

◆ logf

◆ powf

◆ sinf

◆ sinhf

◆ sycl_printf

◆ sycl_printf_

◆ device