Differential D8223 Diff 26558 intern/cycles/kernel/kernels/optix/kernel_optix.cu

Changeset View

Standalone View

intern/cycles/kernel/kernels/optix/kernel_optix.cu

Show All 9 Lines
*		*
* Unless required by applicable law or agreed to in writing, software		* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,		* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.		* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and		* See the License for the specific language governing permissions and
* limitations under the License.		* limitations under the License.
*/		*/

		// clang-format off
#include "kernel/kernel_compat_optix.h"		#include "kernel/kernel_compat_optix.h"
#include "util/util_atomic.h"		#include "util/util_atomic.h"
#include "kernel/kernel_types.h"		#include "kernel/kernel_types.h"
#include "kernel/kernel_globals.h"		#include "kernel/kernel_globals.h"
#include "../cuda/kernel_cuda_image.h" // Texture lookup uses normal CUDA intrinsics		#include "../cuda/kernel_cuda_image.h" // Texture lookup uses normal CUDA intrinsics

#include "kernel/kernel_path.h"		#include "kernel/kernel_path.h"
#include "kernel/kernel_bake.h"		#include "kernel/kernel_bake.h"
		// clang-format on

template<typename T> ccl_device_forceinline T *get_payload_ptr_0()		template<typename T> ccl_device_forceinline T *get_payload_ptr_0()
{		{
return (T *)(((uint64_t)optixGetPayload_1() << 32) \| optixGetPayload_0());		return (T *)(((uint64_t)optixGetPayload_1() << 32) \| optixGetPayload_0());
}		}
template<typename T> ccl_device_forceinline T *get_payload_ptr_2()		template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
{		{
return (T *)(((uint64_t)optixGetPayload_3() << 32) \| optixGetPayload_2());		return (T *)(((uint64_t)optixGetPayload_3() << 32) \| optixGetPayload_2());
▲ Show 20 Lines • Show All 100 Lines • ▼ Show 20 Lines	if (local_isect->num_hits > max_hits) {
hit = lcg_step_uint(lcg_state) % local_isect->num_hits;		hit = lcg_step_uint(lcg_state) % local_isect->num_hits;
if (hit >= max_hits) {		if (hit >= max_hits) {
return optixIgnoreIntersection();		return optixIgnoreIntersection();
}		}
}		}
}		}
else {		else {
if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {		if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {
// Record closest intersection only (do not terminate ray here, since there is no guarantee		// Record closest intersection only
// about distance ordering in anyhit)		// Do not terminate ray here, since there is no guarantee about distance ordering in any-hit
return optixIgnoreIntersection();		return optixIgnoreIntersection();
}		}

local_isect->num_hits = 1;		local_isect->num_hits = 1;
}		}

Intersection *isect = &local_isect->hits[hit];		Intersection *isect = &local_isect->hits[hit];
isect->t = optixGetRayTmax();		isect->t = optixGetRayTmax();
isect->prim = optixGetPrimitiveIndex();		isect->prim = optixGetPrimitiveIndex();
isect->object = get_object_id();		isect->object = get_object_id();
isect->type = kernel_tex_fetch(__prim_type, isect->prim);		isect->type = kernel_tex_fetch(__prim_type, isect->prim);

if (optixIsTriangleHit()) {
const float2 barycentrics = optixGetTriangleBarycentrics();		const float2 barycentrics = optixGetTriangleBarycentrics();
isect->u = 1.0f - barycentrics.y - barycentrics.x;		isect->u = 1.0f - barycentrics.y - barycentrics.x;
isect->v = barycentrics.x;		isect->v = barycentrics.x;
}
else {
isect->u = __uint_as_float(optixGetAttribute_0());
isect->v = __uint_as_float(optixGetAttribute_1());
}

// Record geometric normal		// Record geometric normal
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);		const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0));		const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0));
const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1));		const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1));
const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));		const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));		local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));

Show All 20 Lines	# endif
isect->object = get_object_id();		isect->object = get_object_id();
isect->type = kernel_tex_fetch(__prim_type, prim);		isect->type = kernel_tex_fetch(__prim_type, prim);

if (optixIsTriangleHit()) {		if (optixIsTriangleHit()) {
const float2 barycentrics = optixGetTriangleBarycentrics();		const float2 barycentrics = optixGetTriangleBarycentrics();
isect->u = 1.0f - barycentrics.y - barycentrics.x;		isect->u = 1.0f - barycentrics.y - barycentrics.x;
isect->v = barycentrics.x;		isect->v = barycentrics.x;
}		}
		# ifdef __HAIR__
else {		else {
isect->u = __uint_as_float(optixGetAttribute_0());		const float u = __uint_as_float(optixGetAttribute_0());
		isect->u = u;
isect->v = __uint_as_float(optixGetAttribute_1());		isect->v = __uint_as_float(optixGetAttribute_1());

		// Filter out curve endcaps
		if (u == 0.0f \|\| u == 1.0f) {
		return optixIgnoreIntersection();
}		}
		}
		# endif

# ifdef __TRANSPARENT_SHADOWS__		# ifdef __TRANSPARENT_SHADOWS__
// Detect if this surface has a shader with transparent shadows		// Detect if this surface has a shader with transparent shadows
if (!shader_transparent_shadow(NULL, isect) \|\| optixGetPayload_2() >= optixGetPayload_3()) {		if (!shader_transparent_shadow(NULL, isect) \|\| optixGetPayload_2() >= optixGetPayload_3()) {
# endif		# endif
// This is an opaque hit or the hit limit has been reached, abort traversal		// This is an opaque hit or the hit limit has been reached, abort traversal
optixSetPayload_5(true);		optixSetPayload_5(true);
return optixTerminateRay();		return optixTerminateRay();
# ifdef __TRANSPARENT_SHADOWS__		# ifdef __TRANSPARENT_SHADOWS__
}		}

// TODO(pmours): Do we need REQUIRE_UNIQUE_ANYHIT for this to work?
optixSetPayload_2(optixGetPayload_2() + 1); // num_hits++		optixSetPayload_2(optixGetPayload_2() + 1); // num_hits++

// Continue tracing		// Continue tracing
optixIgnoreIntersection();		optixIgnoreIntersection();
# endif		# endif
#endif		#endif
}		}

extern "C" __global__ void __anyhit__kernel_optix_visibility_test()		extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
{		{
uint visibility = optixGetPayload_4();		uint visibility = optixGetPayload_4();
#ifdef __VISIBILITY_FLAG__		#ifdef __VISIBILITY_FLAG__
const uint prim = optixGetPrimitiveIndex();		const uint prim = optixGetPrimitiveIndex();
if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0)		if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
		return optixIgnoreIntersection();
		}
		#endif

		#ifdef __HAIR__
		if (!optixIsTriangleHit()) {
		// Filter out curve endcaps
		const float u = __uint_as_float(optixGetAttribute_0());
		if (u == 0.0f \|\| u == 1.0f) {
return optixIgnoreIntersection();		return optixIgnoreIntersection();
		}
		}
#endif		#endif

// Shadow ray early termination		// Shadow ray early termination
if (visibility & PATH_RAY_SHADOW_OPAQUE)		if (visibility & PATH_RAY_SHADOW_OPAQUE) {
return optixTerminateRay();		return optixTerminateRay();
}		}
		}

extern "C" __global__ void __closesthit__kernel_optix_hit()		extern "C" __global__ void __closesthit__kernel_optix_hit()
{		{
optixSetPayload_0(__float_as_uint(optixGetRayTmax())); // Intersection distance		optixSetPayload_0(__float_as_uint(optixGetRayTmax())); // Intersection distance
optixSetPayload_3(optixGetPrimitiveIndex());		optixSetPayload_3(optixGetPrimitiveIndex());
optixSetPayload_4(get_object_id());		optixSetPayload_4(get_object_id());
// Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index		// Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index
optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex()));		optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex()));

if (optixIsTriangleHit()) {		if (optixIsTriangleHit()) {
const float2 barycentrics = optixGetTriangleBarycentrics();		const float2 barycentrics = optixGetTriangleBarycentrics();
optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));		optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
optixSetPayload_2(__float_as_uint(barycentrics.x));		optixSetPayload_2(__float_as_uint(barycentrics.x));
}		}
else {		else {
optixSetPayload_1(optixGetAttribute_0());		optixSetPayload_1(optixGetAttribute_0()); // Same as 'optixGetCurveParameter()'
optixSetPayload_2(optixGetAttribute_1());		optixSetPayload_2(optixGetAttribute_1());
}		}
}		}

#ifdef __HAIR__		#ifdef __HAIR__
ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)		ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)
{		{
const uint object = get_object_id<true>();		const uint object = get_object_id<true>();
Show All 19 Lines	if (isect.t != FLT_MAX)
isect.t *= len;		isect.t *= len;

if (curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) {		if (curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) {
optixReportIntersection(isect.t / len,		optixReportIntersection(isect.t / len,
type & PRIMITIVE_ALL,		type & PRIMITIVE_ALL,
__float_as_int(isect.u), // Attribute_0		__float_as_int(isect.u), // Attribute_0
__float_as_int(isect.v)); // Attribute_1		__float_as_int(isect.v)); // Attribute_1
}		}

}		}

extern "C" __global__ void __intersection__curve_ribbon()		extern "C" __global__ void __intersection__curve_ribbon()
{		{
const uint prim = optixGetPrimitiveIndex();		const uint prim = optixGetPrimitiveIndex();
const uint type = kernel_tex_fetch(__prim_type, prim);		const uint type = kernel_tex_fetch(__prim_type, prim);

if (type & (PRIMITIVE_CURVE_RIBBON \| PRIMITIVE_MOTION_CURVE_RIBBON)) {		if (type & (PRIMITIVE_CURVE_RIBBON \| PRIMITIVE_MOTION_CURVE_RIBBON)) {
Show All 18 Lines