Cycles: Add support for building with OptiX 7.4 SDK and use built-in catmull-rom curve type
Some enum names were changed/removed in OptiX 7.4, so some changes are necessary to make things compile still. In addition, OptiX 7.4 also adds built-in support for catmull-rom curves, so it is no longer necessary to convert the catmull-rom data to cubic bsplines first, and has endcaps disabled by default now, so can remove the special handling via any-hit programs that filtered them out before. Differential Revision: https://developer.blender.org/D13351
This commit is contained in:
parent
72acce43bc
commit
7a97e925fd
|
@ -208,11 +208,15 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
|
|||
}
|
||||
else {
|
||||
module_options.optLevel = OPTIX_COMPILE_OPTIMIZATION_LEVEL_3;
|
||||
module_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_LINEINFO;
|
||||
module_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_NONE;
|
||||
}
|
||||
|
||||
module_options.boundValues = nullptr;
|
||||
module_options.numBoundValues = 0;
|
||||
# if OPTIX_ABI_VERSION >= 55
|
||||
module_options.payloadTypes = nullptr;
|
||||
module_options.numPayloadTypes = 0;
|
||||
# endif
|
||||
|
||||
OptixPipelineCompileOptions pipeline_options = {};
|
||||
/* Default to no motion blur and two-level graph, since it is the fastest option. */
|
||||
|
@ -227,7 +231,11 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
|
|||
pipeline_options.usesPrimitiveTypeFlags = OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE;
|
||||
if (kernel_features & KERNEL_FEATURE_HAIR) {
|
||||
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
|
||||
# if OPTIX_ABI_VERSION >= 55
|
||||
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CATMULLROM;
|
||||
# else
|
||||
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CUBIC_BSPLINE;
|
||||
# endif
|
||||
}
|
||||
else
|
||||
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM;
|
||||
|
@ -324,7 +332,13 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
|
|||
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
|
||||
/* Built-in thick curve intersection. */
|
||||
OptixBuiltinISOptions builtin_options = {};
|
||||
# if OPTIX_ABI_VERSION >= 55
|
||||
builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM;
|
||||
builtin_options.buildFlags = OPTIX_BUILD_FLAG_PREFER_FAST_TRACE;
|
||||
builtin_options.curveEndcapFlags = OPTIX_CURVE_ENDCAP_DEFAULT; /* Disable endcaps. */
|
||||
# else
|
||||
builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE;
|
||||
# endif
|
||||
builtin_options.usesMotionBlur = false;
|
||||
|
||||
optix_assert(optixBuiltinISModuleGet(
|
||||
|
@ -411,7 +425,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
|
|||
link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_FULL;
|
||||
}
|
||||
else {
|
||||
link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_LINEINFO;
|
||||
link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_NONE;
|
||||
}
|
||||
|
||||
if (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) {
|
||||
|
@ -1178,6 +1192,15 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
int ka = max(k0 - 1, curve.first_key);
|
||||
int kb = min(k1 + 1, curve.first_key + curve.num_keys - 1);
|
||||
|
||||
index_data[i] = i * 4;
|
||||
float4 *const v = vertex_data.data() + step * num_vertices + index_data[i];
|
||||
|
||||
# if OPTIX_ABI_VERSION >= 55
|
||||
v[0] = make_float4(keys[ka].x, keys[ka].y, keys[ka].z, curve_radius[ka]);
|
||||
v[1] = make_float4(keys[k0].x, keys[k0].y, keys[k0].z, curve_radius[k0]);
|
||||
v[2] = make_float4(keys[k1].x, keys[k1].y, keys[k1].z, curve_radius[k1]);
|
||||
v[3] = make_float4(keys[kb].x, keys[kb].y, keys[kb].z, curve_radius[kb]);
|
||||
# else
|
||||
const float4 px = make_float4(keys[ka].x, keys[k0].x, keys[k1].x, keys[kb].x);
|
||||
const float4 py = make_float4(keys[ka].y, keys[k0].y, keys[k1].y, keys[kb].y);
|
||||
const float4 pz = make_float4(keys[ka].z, keys[k0].z, keys[k1].z, keys[kb].z);
|
||||
|
@ -1190,8 +1213,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
static const float4 cr2bsp2 = make_float4(+1, -4, 11, -2) / 6.f;
|
||||
static const float4 cr2bsp3 = make_float4(-2, +5, -4, +7) / 6.f;
|
||||
|
||||
index_data[i] = i * 4;
|
||||
float4 *const v = vertex_data.data() + step * num_vertices + index_data[i];
|
||||
v[0] = make_float4(
|
||||
dot(cr2bsp0, px), dot(cr2bsp0, py), dot(cr2bsp0, pz), dot(cr2bsp0, pw));
|
||||
v[1] = make_float4(
|
||||
|
@ -1200,6 +1221,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
dot(cr2bsp2, px), dot(cr2bsp2, py), dot(cr2bsp2, pz), dot(cr2bsp2, pw));
|
||||
v[3] = make_float4(
|
||||
dot(cr2bsp3, px), dot(cr2bsp3, py), dot(cr2bsp3, pz), dot(cr2bsp3, pw));
|
||||
# endif
|
||||
}
|
||||
else {
|
||||
BoundBox bounds = BoundBox::empty;
|
||||
|
@ -1241,7 +1263,11 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
OptixBuildInput build_input = {};
|
||||
if (hair->curve_shape == CURVE_THICK) {
|
||||
build_input.type = OPTIX_BUILD_INPUT_TYPE_CURVES;
|
||||
# if OPTIX_ABI_VERSION >= 55
|
||||
build_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM;
|
||||
# else
|
||||
build_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE;
|
||||
# endif
|
||||
build_input.curveArray.numPrimitives = num_segments;
|
||||
build_input.curveArray.vertexBuffers = (CUdeviceptr *)vertex_ptrs.data();
|
||||
build_input.curveArray.numVertices = num_vertices;
|
||||
|
@ -1422,9 +1448,12 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
instance.sbtOffset = PG_HITD_MOTION - PG_HITD;
|
||||
}
|
||||
}
|
||||
else {
|
||||
/* Can disable __anyhit__kernel_optix_visibility_test by default (except for thick curves,
|
||||
* since it needs to filter out end-caps there).
|
||||
# if OPTIX_ABI_VERSION < 55
|
||||
/* Cannot disable any-hit program for thick curves, since it needs to filter out endcaps. */
|
||||
else
|
||||
# endif
|
||||
{
|
||||
/* Can disable __anyhit__kernel_optix_visibility_test by default.
|
||||
* It is enabled where necessary (visibility mask exceeds 8 bits or the other any-hit
|
||||
* programs like __anyhit__kernel_optix_shadow_all_hit) via OPTIX_RAY_FLAG_ENFORCE_ANYHIT.
|
||||
*/
|
||||
|
@ -1494,9 +1523,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
cuMemcpyHtoD(motion_transform_gpu, &motion_transform, motion_transform_size);
|
||||
delete[] reinterpret_cast<uint8_t *>(&motion_transform);
|
||||
|
||||
/* Disable instance transform if object uses motion transform already. */
|
||||
instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
|
||||
|
||||
/* Get traversable handle to motion transform. */
|
||||
optixConvertPointerToTraversableHandle(context,
|
||||
motion_transform_gpu,
|
||||
|
@ -1510,10 +1536,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
|
|||
/* Set transform matrix. */
|
||||
memcpy(instance.transform, &ob->get_tfm(), sizeof(instance.transform));
|
||||
}
|
||||
else {
|
||||
/* Disable instance transform if geometry already has it applied to vertex data. */
|
||||
instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -31,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());
|
||||
|
@ -200,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
|
||||
|
||||
|
@ -310,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());
|
||||
|
@ -317,6 +322,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
|
|||
return optixIgnoreIntersection();
|
||||
}
|
||||
}
|
||||
# endif
|
||||
#endif
|
||||
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
|
|
Loading…
Reference in New Issue