1 /* 2 * Copyright 2019, NVIDIA Corporation. 3 * Copyright 2019, Blender Foundation. 4 * 5 * Licensed under the Apache License, Version 2.0 (the "License"); 6 * you may not use this file except in compliance with the License. 7 * You may obtain a copy of the License at 8 * 9 * http://www.apache.org/licenses/LICENSE-2.0 10 * 11 * Unless required by applicable law or agreed to in writing, software 12 * distributed under the License is distributed on an "AS IS" BASIS, 13 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 14 * See the License for the specific language governing permissions and 15 * limitations under the License. 16 */ 17 18 #ifndef __KERNEL_COMPAT_OPTIX_H__ 19 #define __KERNEL_COMPAT_OPTIX_H__ 20 21 #define OPTIX_DONT_INCLUDE_CUDA 22 #include <optix.h> 23 24 #define __KERNEL_GPU__ 25 #define __KERNEL_CUDA__ // OptiX kernels are implicitly CUDA kernels too 26 #define __KERNEL_OPTIX__ 27 #define CCL_NAMESPACE_BEGIN 28 #define CCL_NAMESPACE_END 29 30 #ifndef ATTR_FALLTHROUGH 31 # define ATTR_FALLTHROUGH 32 #endif 33 34 typedef unsigned int uint32_t; 35 typedef unsigned long long uint64_t; 36 typedef unsigned short half; 37 typedef unsigned long long CUtexObject; 38 39 #ifdef CYCLES_CUBIN_CC 40 # define FLT_MIN 1.175494350822287507969e-38f 41 # define FLT_MAX 340282346638528859811704183484516925440.0f 42 # define FLT_EPSILON 1.192092896e-07F 43 #endif 44 __float2half(const float f)45__device__ half __float2half(const float f) 46 { 47 half val; 48 asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f)); 49 return val; 50 } 51 52 /* Selective nodes compilation. */ 53 #ifndef __NODES_MAX_GROUP__ 54 # define __NODES_MAX_GROUP__ NODE_GROUP_LEVEL_MAX 55 #endif 56 #ifndef __NODES_FEATURES__ 57 # define __NODES_FEATURES__ NODE_FEATURE_ALL 58 #endif 59 60 #define ccl_device \ 61 __device__ __forceinline__ // Function calls are bad for OptiX performance, so inline everything 62 #define ccl_device_inline ccl_device 63 #define ccl_device_forceinline ccl_device 64 #define ccl_device_noinline __device__ __noinline__ 65 #define ccl_device_noinline_cpu ccl_device 66 #define ccl_global 67 #define ccl_static_constant __constant__ 68 #define ccl_constant const 69 #define ccl_local 70 #define ccl_local_param 71 #define ccl_private 72 #define ccl_may_alias 73 #define ccl_addr_space 74 #define ccl_loop_no_unroll 75 #define ccl_restrict __restrict__ 76 #define ccl_ref 77 #define ccl_align(n) __align__(n) 78 79 // Zero initialize structs to help the compiler figure out scoping 80 #define ccl_optional_struct_init = {} 81 82 #define kernel_data __params.data // See kernel_globals.h 83 #define kernel_tex_array(t) __params.t 84 #define kernel_tex_fetch(t, index) __params.t[(index)] 85 86 #define kernel_assert(cond) 87 88 /* Types */ 89 90 #include "util/util_half.h" 91 #include "util/util_types.h" 92 93 #endif /* __KERNEL_COMPAT_OPTIX_H__ */ 94