diff options
Diffstat (limited to 'intern/cycles/kernel/device/optix')
| -rw-r--r-- | intern/cycles/kernel/device/optix/compat.h | 5 | ||||
| -rw-r--r-- | intern/cycles/kernel/device/optix/kernel.cu | 18 |
2 files changed, 16 insertions, 7 deletions
diff --git a/intern/cycles/kernel/device/optix/compat.h b/intern/cycles/kernel/device/optix/compat.h index 835e4621d47..0619c135c39 100644 --- a/intern/cycles/kernel/device/optix/compat.h +++ b/intern/cycles/kernel/device/optix/compat.h @@ -49,10 +49,11 @@ typedef unsigned long long uint64_t; __device__ __forceinline__ // Function calls are bad for OptiX performance, so inline everything #define ccl_device_inline ccl_device #define ccl_device_forceinline ccl_device +#define ccl_device_inline_method ccl_device #define ccl_device_noinline __device__ __noinline__ #define ccl_device_noinline_cpu ccl_device #define ccl_global -#define ccl_static_constant __constant__ +#define ccl_inline_constant __constant__ #define ccl_device_constant __constant__ __device__ #define ccl_constant const #define ccl_gpu_shared __shared__ @@ -76,6 +77,7 @@ typedef unsigned long long uint64_t; #define ccl_gpu_block_idx_x (blockIdx.x) #define ccl_gpu_grid_dim_x (gridDim.x) #define ccl_gpu_warp_size (warpSize) +#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp)) #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x) #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x) @@ -85,7 +87,6 @@ typedef unsigned long long uint64_t; #define ccl_gpu_syncthreads() __syncthreads() #define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate) #define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down_sync(mask, var, detla) -#define ccl_gpu_popc(x) __popc(x) /* GPU texture objects */ diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index b987aa7a817..70b977b3d84 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -21,6 +21,8 @@ #include "kernel/device/gpu/image.h" /* Texture lookup uses normal CUDA intrinsics. */ +#include "kernel/tables.h" + #include "kernel/integrator/state.h" #include "kernel/integrator/state_flow.h" #include "kernel/integrator/state_util.h" @@ -29,9 +31,11 @@ #include "kernel/integrator/intersect_shadow.h" #include "kernel/integrator/intersect_subsurface.h" #include "kernel/integrator/intersect_volume_stack.h" - // clang-format on +#define OPTIX_DEFINE_ABI_VERSION_ONLY +#include <optix_function_table.h> + template<typename T> ccl_device_forceinline T *get_payload_ptr_0() { return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1()); @@ -44,7 +48,7 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2() ccl_device_forceinline int get_object_id() { #ifdef __OBJECT_MOTION__ - /* Always get the the instance ID from the TLAS + /* Always get the instance ID from the TLAS * There might be a motion transform node between TLAS and BLAS which does not have one. */ return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0)); #else @@ -159,9 +163,9 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit() /* Record geometric normal. */ const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w; - const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)); - const float3 tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)); - const float3 tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2)); + const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0); + const float3 tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1); + const float3 tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2); local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); /* Continue tracing (without this the trace call would return after the first hit). */ @@ -198,10 +202,12 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() type = segment.type; prim = segment.prim; +# if OPTIX_ABI_VERSION < 55 /* Filter out curve endcaps. */ if (u == 0.0f || u == 1.0f) { return optixIgnoreIntersection(); } +# endif } # endif @@ -308,6 +314,7 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test() extern "C" __global__ void __anyhit__kernel_optix_visibility_test() { #ifdef __HAIR__ +# if OPTIX_ABI_VERSION < 55 if (!optixIsTriangleHit()) { /* Filter out curve endcaps. */ const float u = __uint_as_float(optixGetAttribute_0()); @@ -315,6 +322,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test() return optixIgnoreIntersection(); } } +# endif #endif #ifdef __VISIBILITY_FLAG__ |