[Bf-blender-cvs] [063ad8635ec] master: Cycles: reduce triangle memory usage with packed_float3
Brecht Van Lommel
noreply at git.blender.org
Wed Nov 17 18:25:38 CET 2021
Commit: 063ad8635ec87a490d6fc02c937387a3c6673b08
Author: Brecht Van Lommel
Date: Tue Nov 16 14:03:59 2021 +0100
Branches: master
https://developer.blender.org/rB063ad8635ec87a490d6fc02c937387a3c6673b08
Cycles: reduce triangle memory usage with packed_float3
Depends on D13243
Differential Revision: https://developer.blender.org/D13244
===================================================================
M intern/cycles/kernel/device/optix/kernel.cu
M intern/cycles/kernel/geom/attribute.h
M intern/cycles/kernel/geom/curve.h
M intern/cycles/kernel/geom/motion_curve.h
M intern/cycles/kernel/geom/motion_triangle.h
M intern/cycles/kernel/geom/patch.h
M intern/cycles/kernel/geom/primitive.h
M intern/cycles/kernel/geom/subd_triangle.h
M intern/cycles/kernel/geom/triangle.h
M intern/cycles/kernel/geom/triangle_intersect.h
M intern/cycles/kernel/geom/volume.h
M intern/cycles/kernel/textures.h
M intern/cycles/scene/attribute.cpp
M intern/cycles/scene/attribute.h
M intern/cycles/scene/geometry.cpp
M intern/cycles/scene/geometry.h
M intern/cycles/scene/mesh.cpp
M intern/cycles/scene/mesh.h
M intern/cycles/scene/scene.cpp
M intern/cycles/scene/scene.h
===================================================================
diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu
index b987aa7a817..849710ffe61 100644
--- a/intern/cycles/kernel/device/optix/kernel.cu
+++ b/intern/cycles/kernel/device/optix/kernel.cu
@@ -159,9 +159,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). */
diff --git a/intern/cycles/kernel/geom/attribute.h b/intern/cycles/kernel/geom/attribute.h
index 848e0430caa..ae96e7b76ef 100644
--- a/intern/cycles/kernel/geom/attribute.h
+++ b/intern/cycles/kernel/geom/attribute.h
@@ -106,9 +106,9 @@ ccl_device Transform primitive_attribute_matrix(KernelGlobals kg,
{
Transform tfm;
- tfm.x = kernel_tex_fetch(__attributes_float3, desc.offset + 0);
- tfm.y = kernel_tex_fetch(__attributes_float3, desc.offset + 1);
- tfm.z = kernel_tex_fetch(__attributes_float3, desc.offset + 2);
+ tfm.x = kernel_tex_fetch(__attributes_float4, desc.offset + 0);
+ tfm.y = kernel_tex_fetch(__attributes_float4, desc.offset + 1);
+ tfm.z = kernel_tex_fetch(__attributes_float4, desc.offset + 2);
return tfm;
}
diff --git a/intern/cycles/kernel/geom/curve.h b/intern/cycles/kernel/geom/curve.h
index 7271193eef8..4b6eecf9640 100644
--- a/intern/cycles/kernel/geom/curve.h
+++ b/intern/cycles/kernel/geom/curve.h
@@ -126,8 +126,8 @@ ccl_device float3 curve_attribute_float3(KernelGlobals kg,
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
- float3 f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + k0));
- float3 f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + k1));
+ float3 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
+ float3 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + k1);
# ifdef __RAY_DIFFERENTIALS__
if (dx)
@@ -149,7 +149,7 @@ ccl_device float3 curve_attribute_float3(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
desc.offset;
- return float4_to_float3(kernel_tex_fetch(__attributes_float3, offset));
+ return kernel_tex_fetch(__attributes_float3, offset);
}
else {
return make_float3(0.0f, 0.0f, 0.0f);
@@ -168,8 +168,8 @@ ccl_device float4 curve_attribute_float4(KernelGlobals kg,
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
- float4 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
- float4 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + k1);
+ float4 f0 = kernel_tex_fetch(__attributes_float4, desc.offset + k0);
+ float4 f1 = kernel_tex_fetch(__attributes_float4, desc.offset + k1);
# ifdef __RAY_DIFFERENTIALS__
if (dx)
@@ -191,7 +191,7 @@ ccl_device float4 curve_attribute_float4(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
desc.offset;
- return kernel_tex_fetch(__attributes_float3, offset);
+ return kernel_tex_fetch(__attributes_float4, offset);
}
else {
return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
diff --git a/intern/cycles/kernel/geom/motion_curve.h b/intern/cycles/kernel/geom/motion_curve.h
index 2dd213d43f6..8358c94360f 100644
--- a/intern/cycles/kernel/geom/motion_curve.h
+++ b/intern/cycles/kernel/geom/motion_curve.h
@@ -48,8 +48,8 @@ ccl_device_inline void motion_curve_keys_for_step_linear(KernelGlobals kg,
offset += step * numkeys;
- keys[0] = kernel_tex_fetch(__attributes_float3, offset + k0);
- keys[1] = kernel_tex_fetch(__attributes_float3, offset + k1);
+ keys[0] = kernel_tex_fetch(__attributes_float4, offset + k0);
+ keys[1] = kernel_tex_fetch(__attributes_float4, offset + k1);
}
}
@@ -106,10 +106,10 @@ ccl_device_inline void motion_curve_keys_for_step(KernelGlobals kg,
offset += step * numkeys;
- keys[0] = kernel_tex_fetch(__attributes_float3, offset + k0);
- keys[1] = kernel_tex_fetch(__attributes_float3, offset + k1);
- keys[2] = kernel_tex_fetch(__attributes_float3, offset + k2);
- keys[3] = kernel_tex_fetch(__attributes_float3, offset + k3);
+ keys[0] = kernel_tex_fetch(__attributes_float4, offset + k0);
+ keys[1] = kernel_tex_fetch(__attributes_float4, offset + k1);
+ keys[2] = kernel_tex_fetch(__attributes_float4, offset + k2);
+ keys[3] = kernel_tex_fetch(__attributes_float4, offset + k3);
}
}
diff --git a/intern/cycles/kernel/geom/motion_triangle.h b/intern/cycles/kernel/geom/motion_triangle.h
index 43f894938e0..62b7b630c89 100644
--- a/intern/cycles/kernel/geom/motion_triangle.h
+++ b/intern/cycles/kernel/geom/motion_triangle.h
@@ -43,9 +43,9 @@ ccl_device_inline void motion_triangle_verts_for_step(KernelGlobals kg,
{
if (step == numsteps) {
/* center step: regular vertex location */
- verts[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
- verts[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
- verts[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
+ verts[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
+ verts[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
+ verts[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
}
else {
/* center step not store in this array */
@@ -54,9 +54,9 @@ ccl_device_inline void motion_triangle_verts_for_step(KernelGlobals kg,
offset += step * numverts;
- verts[0] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x));
- verts[1] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y));
- verts[2] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z));
+ verts[0] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x);
+ verts[1] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y);
+ verts[2] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z);
}
}
@@ -70,9 +70,9 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals kg,
{
if (step == numsteps) {
/* center step: regular vertex location */
- normals[0] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
- normals[1] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
- normals[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
+ normals[0] = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
+ normals[1] = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
+ normals[2] = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
}
else {
/* center step is not stored in this array */
@@ -81,9 +81,9 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals kg,
offset += step * numverts;
- normals[0] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x));
- normals[1] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y));
- normals[2] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z));
+ normals[0] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x);
+ normals[1] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y);
+ normals[2] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z);
}
}
diff --git a/intern/cycles/kernel/geom/patch.h b/intern/cycles/kernel/geom/patch.h
index 7d24937a41e..432618aa243 100644
--- a/intern/cycles/kernel/geom/patch.h
+++ b/intern/cycles/kernel/geom/patch.h
@@ -380,7 +380,7 @@ ccl_device float3 patch_eval_float3(KernelGlobals kg,
*dv = make_float3(0.0f, 0.0f, 0.0f);
for (int i = 0; i < num_control; i++) {
- float3 v = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + indices[i]));
+ float3 v = kernel_tex_fetch(__attributes_float3, offset + indices[i]);
val += v * weights[i];
if (du)
@@ -417,7 +417,7 @@ ccl_device float4 patch_eval_float4(KernelGlobals kg,
*dv = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
for (int i = 0; i < num_control; i++) {
- float4 v = kernel_tex_fetch(__attributes_float3, offset + indices[i]);
+ float4 v = kernel_tex_fetch(__attributes_float4, offset + indices[i]);
val += v * weights[i];
if (du)
diff --git a/intern/cycles/kernel/geom/primitive.h b/intern/cycles/kernel/geom/primitive.h
index 7a8921b6d6e..6d7b550d82f 100644
--- a/intern/cycles/kernel/geom/primitive.h
+++ b/intern/cycles/kernel/geom/primitive.h
@@ -284,18 +284,33 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
int numverts, numkeys;
object_motion_info(kg, sd->object, NULL, &numverts, &numkeys);
- /* lookup attributes */
- motion_pre = primitive_surface_attribute_float3(kg, sd, desc, NULL, NULL);
-
- desc.offset += (sd->type & PRIMITIVE_ALL_TRIANGLE) ? numverts : numkeys;
- motion_post = primitive_surface_attribute_float3(kg, sd, desc, NULL, NULL);
-
#ifdef __HAIR__
- if (is_curve_primitive && (sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
- object_position_transform(kg, sd, &motion_pre);
- object_position_transform(kg, sd, &motion_post);
+ if (is_curve_primitiv
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list