diff options
| author | Siddhartha Jejurkar <f20180617@goa.bits-pilani.ac.in> | 2021-11-19 21:34:44 +0300 |
|---|---|---|
| committer | Siddhartha Jejurkar <f20180617@goa.bits-pilani.ac.in> | 2021-11-19 21:34:44 +0300 |
| commit | 99a2af76d10e05a18987be5d554ada197b1ca086 (patch) | |
| tree | 9f9ceccccbd117c05a701b6769eb69055be6a28e /intern/cycles/kernel/device/optix/kernel.cu | |
| parent | 139606bd370f96e0a8685547d515a2335591d5de (diff) | |
| parent | 50ad0e15fe0adde7335c89c6fcfa0948c341a08a (diff) | |
Merge branch 'master' into soc-2021-uv-edge-select-support
Diffstat (limited to 'intern/cycles/kernel/device/optix/kernel.cu')
| -rw-r--r-- | intern/cycles/kernel/device/optix/kernel.cu | 12 |
1 files changed, 7 insertions, 5 deletions
diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index 6989219cd9f..4feed59d018 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" @@ -44,7 +46,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 @@ -57,7 +59,7 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest() const int global_index = optixGetLaunchIndex().x; const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] : global_index; - integrator_intersect_closest(nullptr, path_index); + integrator_intersect_closest(nullptr, path_index, __params.render_buffer); } extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow() @@ -159,9 +161,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). */ |