[Bf-blender-cvs] [3e247f0f76e] blender-v3.3-release: Cycles: Enable MetalRT pointclouds & other fixes
Michael Jones
noreply at git.blender.org
Mon Nov 28 19:27:06 CET 2022
Commit: 3e247f0f76ec98a09ce0f206a7e6878cb1521807
Author: Michael Jones
Date: Mon Nov 14 16:51:48 2022 +0000
Branches: blender-v3.3-release
https://developer.blender.org/rB3e247f0f76ec98a09ce0f206a7e6878cb1521807
Cycles: Enable MetalRT pointclouds & other fixes
Differential Revision: https://developer.blender.org/D16499
===================================================================
M intern/cycles/device/metal/bvh.mm
M intern/cycles/device/metal/device_impl.mm
M intern/cycles/device/metal/kernel.h
M intern/cycles/device/metal/kernel.mm
M intern/cycles/kernel/data_template.h
M intern/cycles/kernel/device/metal/bvh.h
M intern/cycles/kernel/device/metal/kernel.metal
M intern/cycles/scene/object.cpp
===================================================================
diff --git a/intern/cycles/device/metal/bvh.mm b/intern/cycles/device/metal/bvh.mm
index 09c4ace081e..a7fd64d3c98 100644
--- a/intern/cycles/device/metal/bvh.mm
+++ b/intern/cycles/device/metal/bvh.mm
@@ -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);
/*------------------------------------------------*/
diff --git a/intern/cycles/device/metal/device_impl.mm b/intern/cycles/device/metal/device_impl.mm
index d1250b83d22..6feeaa0707c 100644
--- a/intern/cycles/device/metal/device_impl.mm
+++ b/intern/cycles/device/metal/device_impl.mm
@@ -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();
}
diff --git a/intern/cycles/device/metal/kernel.h b/intern/cycles/device/metal/kernel.h
index 11393f8b7e1..3e88d2daea7 100644
--- a/intern/cycles/device/metal/kernel.h
+++ b/intern/cycles/device/metal/kernel.h
@@ -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;
diff --git a/intern/cycles/device/metal/kernel.mm b/intern/cycles/device/metal/kernel.mm
index f3a2fc9ec6c..e22b0159108 100644
--- a/intern/cycles/device/metal/kernel.mm
+++ b/intern/cycles/device/metal/kernel.mm
@@ -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];
diff --git a/intern/cycles/kernel/data_template.h b/intern/cycles/kernel/data_template.h
index 807d0650fc3..9f194e1ff57 100644
--- a/intern/cycles/kernel/data_template.h
+++ b/intern/cycles/kernel/data_template.h
@@ -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. */
diff --git a/intern/cycles/kernel/device/metal/bvh.h b/intern/cycles/kernel/device/metal/bvh.h
index 03faa3f020f..2ea2d9c2601 100644
--- a/intern/cycles/kernel/device/metal/bvh.h
+++ b/intern/cycles/kernel/device/metal/bvh.h
@@ -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_in
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list