blender
/
blender
128
395
Fork 567
Code Issues 6.2k Pull Requests 750 Projects 19 Wiki Activity

Cycles: Enable MetalRT pointclouds & other fixes 

Differential Revision: https://developer.blender.org/D16499
This commit is contained in:
Michael Jones (Apple) 2022-11-14 16:51:48 +00:00 committed by Brecht Van Lommel
parent 021c8c7cd0
commit 3e247f0f76
Notes: blender-bot 2023-02-14 05:37:19 +01:00 
Referenced by issue #100749, Blender LTS: Maintenance Task 3.3
8 changed files with 103 additions and 44 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
intern/cycles/device/metal/bvh.mm
View File

@ -496,7 +496,7 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
num_motion_steps = pointcloud->get_motion_steps();
}
const size_t num_aabbs = num_motion_steps;
const size_t num_aabbs = num_motion_steps * num_points;
MTLResourceOptions storage_mode;
if (device.hasUnifiedMemory) {
@ -757,6 +757,10 @@ bool BVHMetal::build_TLAS(Progress &progress,
}
}
if (num_instances == 0) {
return false;
}
/*------------------------------------------------*/
BVH_status("Building TLAS | %7d instances", (int)num_instances);
/*------------------------------------------------*/

3
intern/cycles/device/metal/device_impl.mm
View File

@ -301,6 +301,9 @@ void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_feat
MD5Hash md5;
md5.append(baked_constants);
md5.append(source);
if (use_metalrt) {
md5.append(std::to_string(kernel_features & METALRT_FEATURE_MASK));
}
source_md5[pso_type] = md5.get_hex();
}

8
intern/cycles/device/metal/kernel.h
View File

@ -54,6 +54,10 @@ enum MetalPipelineType {
PSO_NUM
};
# define METALRT_FEATURE_MASK \
(KERNEL_FEATURE_HAIR | KERNEL_FEATURE_HAIR_THICK | KERNEL_FEATURE_POINTCLOUD | \
KERNEL_FEATURE_OBJECT_MOTION)
const char *kernel_type_as_string(MetalPipelineType pso_type);
struct MetalKernelPipeline {
@ -67,9 +71,7 @@ struct MetalKernelPipeline {
KernelData kernel_data_;
bool use_metalrt;
bool metalrt_hair;
bool metalrt_hair_thick;
bool metalrt_pointcloud;
uint32_t metalrt_features = 0;
int threads_per_threadgroup;

53
intern/cycles/device/metal/kernel.mm
View File

@ -225,12 +225,9 @@ void ShaderCache::load_kernel(DeviceKernel device_kernel,
/* metalrt options */
request.pipeline->use_metalrt = device->use_metalrt;
request.pipeline->metalrt_hair = device->use_metalrt &&
(device->kernel_features & KERNEL_FEATURE_HAIR);
request.pipeline->metalrt_hair_thick = device->use_metalrt &&
(device->kernel_features & KERNEL_FEATURE_HAIR_THICK);
request.pipeline->metalrt_pointcloud = device->use_metalrt &&
(device->kernel_features & KERNEL_FEATURE_POINTCLOUD);
request.pipeline->metalrt_features = device->use_metalrt ?
(device->kernel_features & METALRT_FEATURE_MASK) :
0;
{
thread_scoped_lock lock(cache_mutex);
@ -267,9 +264,13 @@ MetalKernelPipeline *ShaderCache::get_best_pipeline(DeviceKernel kernel, const M
/* metalrt options */
bool use_metalrt = device->use_metalrt;
bool metalrt_hair = use_metalrt && (device->kernel_features & KERNEL_FEATURE_HAIR);
bool metalrt_hair_thick = use_metalrt && (device->kernel_features & KERNEL_FEATURE_HAIR_THICK);
bool metalrt_pointcloud = use_metalrt && (device->kernel_features & KERNEL_FEATURE_POINTCLOUD);
bool device_metalrt_hair = use_metalrt && device->kernel_features & KERNEL_FEATURE_HAIR;
bool device_metalrt_hair_thick = use_metalrt &&
device->kernel_features & KERNEL_FEATURE_HAIR_THICK;
bool device_metalrt_pointcloud = use_metalrt &&
device->kernel_features & KERNEL_FEATURE_POINTCLOUD;
bool device_metalrt_motion = use_metalrt &&
device->kernel_features & KERNEL_FEATURE_OBJECT_MOTION;
MetalKernelPipeline *best_pipeline = nullptr;
for (auto &pipeline : collection) {
@ -278,9 +279,16 @@ MetalKernelPipeline *ShaderCache::get_best_pipeline(DeviceKernel kernel, const M
continue;
}
if (pipeline->use_metalrt != use_metalrt || pipeline->metalrt_hair != metalrt_hair ||
pipeline->metalrt_hair_thick != metalrt_hair_thick ||
pipeline->metalrt_pointcloud != metalrt_pointcloud) {
bool pipeline_metalrt_hair = pipeline->metalrt_features & KERNEL_FEATURE_HAIR;
bool pipeline_metalrt_hair_thick = pipeline->metalrt_features & KERNEL_FEATURE_HAIR_THICK;
bool pipeline_metalrt_pointcloud = pipeline->metalrt_features & KERNEL_FEATURE_POINTCLOUD;
bool pipeline_metalrt_motion = use_metalrt &&
pipeline->metalrt_features & KERNEL_FEATURE_OBJECT_MOTION;
if (pipeline->use_metalrt != use_metalrt || pipeline_metalrt_hair != device_metalrt_hair ||
pipeline_metalrt_hair_thick != device_metalrt_hair_thick ||
pipeline_metalrt_pointcloud != device_metalrt_pointcloud ||
pipeline_metalrt_motion != device_metalrt_motion) {
/* wrong combination of metalrt options */
continue;
}
@ -345,6 +353,8 @@ static MTLFunctionConstantValues *GetConstantValues(KernelData const *data = nul
if (!data) {
data = &zero_data;
}
int zero_int = 0;
[constant_values setConstantValue:&zero_int type:MTLDataType_int atIndex:Kernel_DummyConstant];
# define KERNEL_STRUCT_MEMBER(parent, _type, name) \
[constant_values setConstantValue:&data->parent.name \
@ -375,10 +385,7 @@ void MetalKernelPipeline::compile()
MTLFunctionDescriptor *func_desc = [MTLIntersectionFunctionDescriptor functionDescriptor];
func_desc.name = entryPoint;
if (pso_type == PSO_SPECIALIZED_SHADE) {
func_desc.constantValues = GetConstantValues(&kernel_data_);
}
else if (pso_type == PSO_SPECIALIZED_INTERSECT) {
if (pso_type != PSO_GENERIC) {
func_desc.constantValues = GetConstantValues(&kernel_data_);
}
else {
@ -423,6 +430,13 @@ void MetalKernelPipeline::compile()
const char *function_name = function_names[i];
desc.name = [@(function_name) copy];
if (pso_type != PSO_GENERIC) {
desc.constantValues = GetConstantValues(&kernel_data_);
}
else {
desc.constantValues = GetConstantValues();
}
NSError *error = NULL;
rt_intersection_function[i] = [mtlLibrary newFunctionWithDescriptor:desc error:&error];
@ -443,6 +457,10 @@ void MetalKernelPipeline::compile()
NSArray *table_functions[METALRT_TABLE_NUM] = {nil};
NSArray *linked_functions = nil;
bool metalrt_hair = use_metalrt && (metalrt_features & KERNEL_FEATURE_HAIR);
bool metalrt_hair_thick = use_metalrt && (metalrt_features & KERNEL_FEATURE_HAIR_THICK);
bool metalrt_pointcloud = use_metalrt && (metalrt_features & KERNEL_FEATURE_POINTCLOUD);
if (use_metalrt) {
id<MTLFunction> curve_intersect_default = nil;
id<MTLFunction> curve_intersect_shadow = nil;
@ -680,7 +698,8 @@ void MetalKernelPipeline::compile()
newIntersectionFunctionTableWithDescriptor:ift_desc];
/* Finally write the function handles into this pipeline's table */
for (int i = 0; i < 2; i++) {
int size = (int)[table_functions[table] count];
for (int i = 0; i < size; i++) {
id<MTLFunctionHandle> handle = [pipeline
functionHandleWithFunction:table_functions[table][i]];
[intersection_func_table[table] setFunction:handle atIndex:i];

4
intern/cycles/kernel/data_template.h
View File

@ -49,11 +49,11 @@ KERNEL_STRUCT_BEGIN(KernelBVH, bvh)
KERNEL_STRUCT_MEMBER(bvh, int, root)
KERNEL_STRUCT_MEMBER(bvh, int, have_motion)
KERNEL_STRUCT_MEMBER(bvh, int, have_curves)
KERNEL_STRUCT_MEMBER(bvh, int, have_points)
KERNEL_STRUCT_MEMBER(bvh, int, have_volumes)
KERNEL_STRUCT_MEMBER(bvh, int, bvh_layout)
KERNEL_STRUCT_MEMBER(bvh, int, use_bvh_steps)
KERNEL_STRUCT_MEMBER(bvh, int, curve_subdivisions)
KERNEL_STRUCT_MEMBER(bvh, int, pad1)
KERNEL_STRUCT_MEMBER(bvh, int, pad2)
KERNEL_STRUCT_END(KernelBVH)
/* Film. */

19
intern/cycles/kernel/device/metal/bvh.h
View File

@ -79,7 +79,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
if (!kernel_data.bvh.have_curves) {
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
@ -177,7 +178,9 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
if (!kernel_data.bvh.have_curves) {
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
@ -205,7 +208,9 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
if (lcg_state) {
*lcg_state = payload.lcg_state;
}
*local_isect = payload.local_isect;
if (local_isect) {
*local_isect = payload.local_isect;
}
return payload.result;
}
@ -240,7 +245,9 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
if (!kernel_data.bvh.have_curves) {
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
@ -307,7 +314,9 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
if (!kernel_data.bvh.have_curves) {
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}

41
intern/cycles/kernel/device/metal/kernel.metal
View File

@ -182,20 +182,20 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
const float u = barycentrics.x;
const float v = barycentrics.y;
int type = 0;
if (intersection_type == METALRT_HIT_TRIANGLE) {
type = kernel_data_fetch(objects, object).primitive_type;
}
const int prim_type = kernel_data_fetch(objects, object).primitive_type;
int type = prim_type;
# ifdef __HAIR__
else {
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
type = segment.type;
prim = segment.prim;
/* Filter out curve endcaps */
if (u == 0.0f || u == 1.0f) {
/* continue search */
return true;
if (intersection_type != METALRT_HIT_TRIANGLE) {
if ( (prim_type == PRIMITIVE_CURVE_THICK || prim_type == PRIMITIVE_CURVE_RIBBON)) {
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
type = segment.type;
prim = segment.prim;
/* Filter out curve endcaps */
if (u == 0.0f || u == 1.0f) {
/* continue search */
return true;
}
}
}
# endif
@ -279,7 +279,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type;
/* Continue tracing. */
# endif /* __TRANSPARENT_SHADOWS__ */
#endif /* __SHADOW_RECORD_ALL__ */
@ -327,7 +327,8 @@ inline TReturnType metalrt_visibility_test(
TReturnType result;
#ifdef __HAIR__
if (intersection_type == METALRT_HIT_BOUNDING_BOX) {
const int type = kernel_data_fetch(objects, object).primitive_type;
if (intersection_type == METALRT_HIT_BOUNDING_BOX && (type == PRIMITIVE_CURVE_THICK || type == PRIMITIVE_CURVE_RIBBON)) {
/* Filter out curve endcaps. */
if (u == 0.0f || u == 1.0f) {
result.accept = false;
@ -463,7 +464,12 @@ ccl_device_inline void metalrt_intersection_curve_shadow(
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
Intersection isect;
isect.t = ray_tmax;
@ -685,7 +691,12 @@ ccl_device_inline void metalrt_intersection_point_shadow(
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
Intersection isect;
isect.t = ray_tmax;

13
intern/cycles/scene/object.cpp
View File

@ -57,7 +57,8 @@ struct UpdateObjectTransformState {
/* Flags which will be synchronized to Integrator. */
bool have_motion;
bool have_curves;
// bool have_points;
bool have_points;
bool have_volumes;
/* ** Scheduling queue. ** */
Scene *scene;
@ -545,6 +546,12 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
if (geom->geometry_type == Geometry::HAIR) {
state->have_curves = true;
}
if (geom->geometry_type == Geometry::POINTCLOUD) {
state->have_points = true;
}
if (geom->geometry_type == Geometry::VOLUME) {
state->have_volumes = true;
}
/* Light group. */
auto it = scene->lightgroups.find(ob->lightgroup);
@ -591,6 +598,8 @@ void ObjectManager::device_update_transforms(DeviceScene *dscene, Scene *scene,
state.need_motion = scene->need_motion();
state.have_motion = false;
state.have_curves = false;
state.have_points = false;
state.have_volumes = false;
state.scene = scene;
state.queue_start_object = 0;
@ -658,6 +667,8 @@ void ObjectManager::device_update_transforms(DeviceScene *dscene, Scene *scene,
dscene->data.bvh.have_motion = state.have_motion;
dscene->data.bvh.have_curves = state.have_curves;
dscene->data.bvh.have_points = state.have_points;
dscene->data.bvh.have_volumes = state.have_volumes;
dscene->objects.clear_modified();
dscene->object_motion_pass.clear_modified();