blender
/
blender
136
406
Fork 588
Code Issues 6.2k Pull Requests 755 Projects 19 Wiki Activity

Merge branch 'master' into sculpt-dev

This commit is contained in:
Joseph Eagar 2021-10-06 18:33:08 -07:00
parent b970edb97e 70cc80ea1c
commit c2f2a8260c
403 changed files with 9542 additions and 6398 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
.clang-format
View File

@ -180,6 +180,7 @@ ForEachMacros:
- CTX_DATA_BEGIN_WITH_ID
- DEG_OBJECT_ITER_BEGIN
- DEG_OBJECT_ITER_FOR_RENDER_ENGINE_BEGIN
- DRW_ENABLED_ENGINE_ITER
- DRIVER_TARGETS_LOOPER_BEGIN
- DRIVER_TARGETS_USED_LOOPER_BEGIN
- FOREACH_BASE_IN_EDIT_MODE_BEGIN

1
CMakeLists.txt
View File

@ -406,6 +406,7 @@ mark_as_advanced(WITH_CYCLES_CUDA_BUILD_SERIAL)
set(CYCLES_TEST_DEVICES CPU CACHE STRING "Run regression tests on the specified device types (CPU CUDA OPTIX)" )
set(CYCLES_CUDA_BINARIES_ARCH sm_30 sm_35 sm_37 sm_50 sm_52 sm_60 sm_61 sm_70 sm_75 sm_86 compute_75 CACHE STRING "CUDA architectures to build binaries for")
mark_as_advanced(CYCLES_CUDA_BINARIES_ARCH)
option(WITH_CYCLES_HIP_BINARIES "Build Cycles HIP binaries" OFF)
unset(PLATFORM_DEFAULT)
option(WITH_CYCLES_LOGGING "Build Cycles with logging support" ON)
option(WITH_CYCLES_DEBUG_NAN "Build Cycles with additional asserts for detecting NaNs and invalid values" OFF)

81
build_files/cmake/Modules/FindHIP.cmake Normal file
View File

@ -0,0 +1,81 @@
# - Find HIP compiler
#
# This module defines
# HIP_HIPCC_EXECUTABLE, the full path to the hipcc executable
# HIP_VERSION, the HIP compiler version
#
# HIP_FOUND, if the HIP toolkit is found.
#=============================================================================
# Copyright 2021 Blender Foundation.
#
# Distributed under the OSI-approved BSD 3-Clause License,
# see accompanying file BSD-3-Clause-license.txt for details.
#=============================================================================
# If HIP_ROOT_DIR was defined in the environment, use it.
if(NOT HIP_ROOT_DIR AND NOT $ENV{HIP_ROOT_DIR} STREQUAL "")
set(HIP_ROOT_DIR $ENV{HIP_ROOT_DIR})
endif()
set(_hip_SEARCH_DIRS
${HIP_ROOT_DIR}
)
find_program(HIP_HIPCC_EXECUTABLE
NAMES
hipcc
HINTS
${_hip_SEARCH_DIRS}
PATH_SUFFIXES
bin
)
if(HIP_HIPCC_EXECUTABLE AND NOT EXISTS ${HIP_HIPCC_EXECUTABLE})
message(WARNING "Cached or directly specified hipcc executable does not exist.")
set(HIP_FOUND FALSE)
elseif(HIP_HIPCC_EXECUTABLE)
set(HIP_FOUND TRUE)
set(HIP_VERSION_MAJOR 0)
set(HIP_VERSION_MINOR 0)
set(HIP_VERSION_PATCH 0)
# Get version from the output.
execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} --version
OUTPUT_VARIABLE HIP_VERSION_RAW
ERROR_QUIET
OUTPUT_STRIP_TRAILING_WHITESPACE)
# Parse parts.
if(HIP_VERSION_RAW MATCHES "HIP version: .*")
# Strip the HIP prefix and get list of individual version components.
string(REGEX REPLACE
".*HIP version: ([.0-9]+).*" "\\1"
HIP_SEMANTIC_VERSION "${HIP_VERSION_RAW}")
string(REPLACE "." ";" HIP_VERSION_PARTS "${HIP_SEMANTIC_VERSION}")
list(LENGTH HIP_VERSION_PARTS NUM_HIP_VERSION_PARTS)
# Extract components into corresponding variables.
if(NUM_HIP_VERSION_PARTS GREATER 0)
list(GET HIP_VERSION_PARTS 0 HIP_VERSION_MAJOR)
endif()
if(NUM_HIP_VERSION_PARTS GREATER 1)
list(GET HIP_VERSION_PARTS 1 HIP_VERSION_MINOR)
endif()
if(NUM_HIP_VERSION_PARTS GREATER 2)
list(GET HIP_VERSION_PARTS 2 HIP_VERSION_PATCH)
endif()
# Unset temp variables.
unset(NUM_HIP_VERSION_PARTS)
unset(HIP_SEMANTIC_VERSION)
unset(HIP_VERSION_PARTS)
endif()
# Construct full semantic version.
set(HIP_VERSION "${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}")
unset(HIP_VERSION_RAW)
else()
set(HIP_FOUND FALSE)
endif()

21
build_files/cmake/cmake_static_check_cppcheck.py
View File

@ -24,6 +24,7 @@ import project_source_info
import subprocess
import sys
import os
import tempfile
from typing import (
Any,
@ -35,7 +36,6 @@ USE_QUIET = (os.environ.get("QUIET", None) is not None)
CHECKER_IGNORE_PREFIX = [
"extern",
"intern/moto",
]
CHECKER_BIN = "cppcheck"
@ -47,13 +47,19 @@ CHECKER_ARGS = [
"--max-configs=1", # speeds up execution
# "--check-config", # when includes are missing
"--enable=all", # if you want sixty hundred pedantic suggestions
# Quiet output, otherwise all defines/includes are printed (overly verbose).
# Only enable this for troubleshooting (if defines are not set as expected for example).
"--quiet",
# NOTE: `--cppcheck-build-dir=<dir>` is added later as a temporary directory.
]
if USE_QUIET:
CHECKER_ARGS.append("--quiet")
def main() -> None:
def cppcheck() -> None:
source_info = project_source_info.build_info(ignore_prefix_list=CHECKER_IGNORE_PREFIX)
source_defines = project_source_info.build_defines_as_args()
@ -78,7 +84,10 @@ def main() -> None:
percent_str = "[" + ("%.2f]" % percent).rjust(7) + " %:"
sys.stdout.flush()
sys.stdout.write("%s " % percent_str)
sys.stdout.write("%s %s\n" % (
percent_str,
os.path.relpath(c, project_source_info.SOURCE_DIR)
))
return subprocess.Popen(cmd)
@ -90,5 +99,11 @@ def main() -> None:
print("Finished!")
def main() -> None:
with tempfile.TemporaryDirectory() as temp_dir:
CHECKER_ARGS.append("--cppcheck-build-dir=" + temp_dir)
cppcheck()
if __name__ == "__main__":
main()

12
build_files/cmake/project_source_info.py
View File

@ -243,7 +243,9 @@ def build_defines_as_args() -> List[str]:
# use this module.
def queue_processes(
process_funcs: Sequence[Tuple[Callable[..., subprocess.Popen[Any]], Tuple[Any, ...]]],
*,
job_total: int =-1,
sleep: float = 0.1,
) -> None:
""" Takes a list of function arg pairs, each function must return a process
"""
@ -271,14 +273,20 @@ def queue_processes(
if len(processes) <= job_total:
break
else:
time.sleep(0.1)
time.sleep(sleep)
sys.stdout.flush()
sys.stderr.flush()
processes.append(func(*args))
# Don't return until all jobs have finished.
while 1:
processes[:] = [p for p in processes if p.poll() is None]
if not processes:
break
time.sleep(sleep)
def main() -> None:
if not os.path.exists(join(CMAKE_DIR, "CMakeCache.txt")):

1
doc/python_api/sphinx_doc_gen.py
View File

@ -1101,6 +1101,7 @@ context_type_map = {
"scene": ("Scene", False),
"sculpt_object": ("Object", False),
"selectable_objects": ("Object", True),
"selected_asset_files": ("FileSelectEntry", True),
"selected_bones": ("EditBone", True),
"selected_editable_bones": ("EditBone", True),
"selected_editable_fcurves": ("FCurve", True),

1
extern/cuew/include/cuew.h vendored
View File

@ -609,6 +609,7 @@ typedef enum cudaError_enum {
CUDA_ERROR_INVALID_GRAPHICS_CONTEXT = 219,
CUDA_ERROR_NVLINK_UNCORRECTABLE = 220,
CUDA_ERROR_JIT_COMPILER_NOT_FOUND = 221,
CUDA_ERROR_UNSUPPORTED_PTX_VERSION = 222,
CUDA_ERROR_INVALID_SOURCE = 300,
CUDA_ERROR_FILE_NOT_FOUND = 301,
CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302,

1
extern/cuew/src/cuew.c vendored
View File

@ -736,6 +736,7 @@ const char *cuewErrorString(CUresult result) {
case CUDA_ERROR_INVALID_GRAPHICS_CONTEXT: return "Invalid graphics context";
case CUDA_ERROR_NVLINK_UNCORRECTABLE: return "Nvlink uncorrectable";
case CUDA_ERROR_JIT_COMPILER_NOT_FOUND: return "Jit compiler not found";
case CUDA_ERROR_UNSUPPORTED_PTX_VERSION: return "Unsupported PTX version";
case CUDA_ERROR_INVALID_SOURCE: return "Invalid source";
case CUDA_ERROR_FILE_NOT_FOUND: return "File not found";
case CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND: return "Link to a shared object failed to resolve";

2
extern/smaa_areatex/smaa_areatex.cpp vendored
View File

@ -574,6 +574,8 @@ Dbl2 AreaDiag::area(Dbl2 p1, Dbl2 p2, int left)
Dbl2 d = p2 - p1;
if (d.x == 0.0)
return Dbl2(0.0, 1.0);
if (d.y == 0.0)
return Dbl2(1.0, 0.0);
double x1 = (double)(1 + left);
double x2 = x1 + 1.0;

1
intern/cycles/blender/addon/engine.py
View File

@ -211,7 +211,6 @@ def list_render_passes(scene, srl):
if crl.use_pass_shadow_catcher: yield ("Shadow Catcher", "RGB", 'COLOR')
# Debug passes.
if crl.pass_debug_render_time: yield ("Debug Render Time", "X", 'VALUE')
if crl.pass_debug_sample_count: yield ("Debug Sample Count", "X", 'VALUE')
# Cryptomatte passes.

6
intern/cycles/blender/addon/properties.py
View File

@ -1197,12 +1197,6 @@ class CyclesCurveRenderSettings(bpy.types.PropertyGroup):
class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
pass_debug_render_time: BoolProperty(
name="Debug Render Time",
description="Render time in milliseconds per sample and pixel",
default=False,
update=update_render_passes,
)
pass_debug_sample_count: BoolProperty(
name="Debug Sample Count",
description="Number of samples/camera rays per pixel",

1
intern/cycles/blender/addon/ui.py
View File

@ -792,7 +792,6 @@ class CYCLES_RENDER_PT_passes_data(CyclesButtonsPanel, Panel):
col.prop(view_layer, "use_pass_material_index")
col = layout.column(heading="Debug", align=True)
col.prop(cycles_view_layer, "pass_debug_render_time", text="Render Time")
col.prop(cycles_view_layer, "pass_debug_sample_count", text="Sample Count")
layout.prop(view_layer, "pass_alpha_threshold")

3
intern/cycles/blender/blender_camera.cpp
View File

@ -927,6 +927,9 @@ BufferParams BlenderSync::get_buffer_params(
params.height = height;
}
params.window_width = params.width;
params.window_height = params.height;
return params;
}

2
intern/cycles/blender/blender_display_driver.cpp
View File

@ -460,7 +460,7 @@ void BlenderDisplayDriver::draw(const Params &params)
/* Texture is requested to be cleared and was not yet cleared.
*
* Do early return which should be equivalent of drawing all-zero texture.
* Watchout for the lock though so that the clear happening during update is properly
* Watch out for the lock though so that the clear happening during update is properly
* synchronized here. */
gl_context_mutex_.unlock();
return;

4
intern/cycles/blender/blender_session.cpp
View File

@ -504,6 +504,10 @@ void BlenderSession::render_frame_finish()
/* Clear driver. */
session->set_output_driver(nullptr);
session->full_buffer_written_cb = function_null;
/* All the files are handled.
* Clear the list so that this session can be re-used by Persistent Data. */
full_buffer_files_.clear();
}
static PassType bake_type_to_pass(const string &bake_type_str, const int bake_filter)

6
intern/cycles/blender/blender_sync.cpp
View File

@ -545,8 +545,6 @@ static PassType get_blender_pass_type(BL::RenderPass &b_pass)
MAP_PASS("Shadow Catcher", PASS_SHADOW_CATCHER);
MAP_PASS("Noisy Shadow Catcher", PASS_SHADOW_CATCHER);
MAP_PASS("Debug Render Time", PASS_RENDER_TIME);
MAP_PASS("AdaptiveAuxBuffer", PASS_ADAPTIVE_AUX_BUFFER);
MAP_PASS("Debug Sample Count", PASS_SAMPLE_COUNT);
@ -604,10 +602,6 @@ void BlenderSync::sync_render_passes(BL::RenderLayer &b_rlay, BL::ViewLayer &b_v
PointerRNA crl = RNA_pointer_get(&b_view_layer.ptr, "cycles");
/* Debug passes. */
if (get_boolean(crl, "pass_debug_render_time")) {
b_engine.add_pass("Debug Render Time", 1, "X", b_view_layer.name().c_str());
pass_add(scene, PASS_RENDER_TIME, "Debug Render Time");
}
if (get_boolean(crl, "pass_debug_sample_count")) {
b_engine.add_pass("Debug Sample Count", 1, "X", b_view_layer.name().c_str());
pass_add(scene, PASS_SAMPLE_COUNT, "Debug Sample Count");

4
intern/cycles/bvh/bvh.h
View File

@ -50,10 +50,6 @@ struct PackedBVH {
array<int4> leaf_nodes;
/* object index to BVH node index mapping for instances */
array<int> object_node;
/* Mapping from primitive index to index in triangle array. */
array<uint> prim_tri_index;
/* Continuous storage of triangle vertices. */
array<float4> prim_tri_verts;
/* primitive type - triangle or strand */
array<int> prim_type;
/* visibility visibilitys for primitives */

61
intern/cycles/bvh/bvh2.cpp
View File

@ -439,61 +439,20 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility
/* Triangles */
void BVH2::pack_triangle(int idx, float4 tri_verts[3])
{
int tob = pack.prim_object[idx];
assert(tob >= 0 && tob < objects.size());
const Mesh *mesh = static_cast<const Mesh *>(objects[tob]->get_geometry());
int tidx = pack.prim_index[idx];
Mesh::Triangle t = mesh->get_triangle(tidx);
const float3 *vpos = &mesh->verts[0];
float3 v0 = vpos[t.v[0]];
float3 v1 = vpos[t.v[1]];
float3 v2 = vpos[t.v[2]];
tri_verts[0] = float3_to_float4(v0);
tri_verts[1] = float3_to_float4(v1);
tri_verts[2] = float3_to_float4(v2);
}
void BVH2::pack_primitives()
{
const size_t tidx_size = pack.prim_index.size();
size_t num_prim_triangles = 0;
/* Count number of triangles primitives in BVH. */
for (unsigned int i = 0; i < tidx_size; i++) {
if ((pack.prim_index[i] != -1)) {
if ((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
++num_prim_triangles;
}
}
}
/* Reserve size for arrays. */
pack.prim_tri_index.clear();
pack.prim_tri_index.resize(tidx_size);
pack.prim_tri_verts.clear();
pack.prim_tri_verts.resize(num_prim_triangles * 3);
pack.prim_visibility.clear();
pack.prim_visibility.resize(tidx_size);
/* Fill in all the arrays. */
size_t prim_triangle_index = 0;
for (unsigned int i = 0; i < tidx_size; i++) {
if (pack.prim_index[i] != -1) {
int tob = pack.prim_object[i];
Object *ob = objects[tob];
if ((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
pack_triangle(i, (float4 *)&pack.prim_tri_verts[3 * prim_triangle_index]);
pack.prim_tri_index[i] = 3 * prim_triangle_index;
++prim_triangle_index;
}
else {
pack.prim_tri_index[i] = -1;
}
pack.prim_visibility[i] = ob->visibility_for_tracing();
}
else {
pack.prim_tri_index[i] = -1;
pack.prim_visibility[i] = 0;
}
}
@ -522,10 +481,8 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
/* reserve */
size_t prim_index_size = pack.prim_index.size();
size_t prim_tri_verts_size = pack.prim_tri_verts.size();
size_t pack_prim_index_offset = prim_index_size;
size_t pack_prim_tri_verts_offset = prim_tri_verts_size;
size_t pack_nodes_offset = nodes_size;
size_t pack_leaf_nodes_offset = leaf_nodes_size;
size_t object_offset = 0;
@ -535,7 +492,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
if (geom->need_build_bvh(params.bvh_layout)) {
prim_index_size += bvh->pack.prim_index.size();
prim_tri_verts_size += bvh->pack.prim_tri_verts.size();
nodes_size += bvh->pack.nodes.size();
leaf_nodes_size += bvh->pack.leaf_nodes.size();
}
@ -545,8 +501,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
pack.prim_type.resize(prim_index_size);
pack.prim_object.resize(prim_index_size);
pack.prim_visibility.resize(prim_index_size);
pack.prim_tri_verts.resize(prim_tri_verts_size);
pack.prim_tri_index.resize(prim_index_size);
pack.nodes.resize(nodes_size);
pack.leaf_nodes.resize(leaf_nodes_size);
pack.object_node.resize(objects.size());
@ -559,8 +513,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
int *pack_prim_type = (pack.prim_type.size()) ? &pack.prim_type[0] : NULL;
int *pack_prim_object = (pack.prim_object.size()) ? &pack.prim_object[0] : NULL;
uint *pack_prim_visibility = (pack.prim_visibility.size()) ? &pack.prim_visibility[0] : NULL;
float4 *pack_prim_tri_verts = (pack.prim_tri_verts.size()) ? &pack.prim_tri_verts[0] : NULL;
uint *pack_prim_tri_index = (pack.prim_tri_index.size()) ? &pack.prim_tri_index[0] : NULL;
int4 *pack_nodes = (pack.nodes.size()) ? &pack.nodes[0] : NULL;
int4 *pack_leaf_nodes = (pack.leaf_nodes.size()) ? &pack.leaf_nodes[0] : NULL;
float2 *pack_prim_time = (pack.prim_time.size()) ? &pack.prim_time[0] : NULL;
@ -609,18 +561,14 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
int *bvh_prim_index = &bvh->pack.prim_index[0];
int *bvh_prim_type = &bvh->pack.prim_type[0];
uint *bvh_prim_visibility = &bvh->pack.prim_visibility[0];
uint *bvh_prim_tri_index = &bvh->pack.prim_tri_index[0];
float2 *bvh_prim_time = bvh->pack.prim_time.size() ? &bvh->pack.prim_time[0] : NULL;
for (size_t i = 0; i < bvh_prim_index_size; i++) {
if (bvh->pack.prim_type[i] & PRIMITIVE_ALL_CURVE) {
pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
pack_prim_tri_index[pack_prim_index_offset] = -1;
}
else {
pack_prim_index[pack_prim_index_offset] = bvh_prim_index[i] + geom_prim_offset;
pack_prim_tri_index[pack_prim_index_offset] = bvh_prim_tri_index[i] +
pack_prim_tri_verts_offset;
}
pack_prim_type[pack_prim_index_offset] = bvh_prim_type[i];
@ -633,15 +581,6 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
}
}
/* Merge triangle vertices data. */
if (bvh->pack.prim_tri_verts.size()) {
const size_t prim_tri_size = bvh->pack.prim_tri_verts.size();
memcpy(pack_prim_tri_verts + pack_prim_tri_verts_offset,
&bvh->pack.prim_tri_verts[0],
prim_tri_size * sizeof(float4));
pack_prim_tri_verts_offset += prim_tri_size;
}
/* merge nodes */
if (bvh->pack.leaf_nodes.size()) {
int4 *leaf_nodes_offset = &bvh->pack.leaf_nodes[0];

36
intern/cycles/bvh/bvh_build.cpp
View File

@ -67,8 +67,12 @@ BVHBuild::~BVHBuild()
/* Adding References */
void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *mesh, int i)
void BVHBuild::add_reference_triangles(BoundBox &root,
BoundBox &center,
Mesh *mesh,
int object_index)
{
const PrimitiveType primitive_type = mesh->primitive_type();
const Attribute *attr_mP = NULL;
if (mesh->has_motion_blur()) {
attr_mP = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@ -81,7 +85,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
BoundBox bounds = BoundBox::empty;
t.bounds_grow(verts, bounds);
if (bounds.valid() && t.valid(verts)) {
references.push_back(BVHReference(bounds, j, i, PRIMITIVE_TRIANGLE));
references.push_back(BVHReference(bounds, j, object_index, primitive_type));
root.grow(bounds);
center.grow(bounds.center2());
}
@ -101,7 +105,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
t.bounds_grow(vert_steps + step * num_verts, bounds);
}
if (bounds.valid()) {
references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE));
references.push_back(BVHReference(bounds, j, object_index, primitive_type));
root.grow(bounds);
center.grow(bounds.center2());
}
@ -140,7 +144,7 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
if (bounds.valid()) {
const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
references.push_back(
BVHReference(bounds, j, i, PRIMITIVE_MOTION_TRIANGLE, prev_time, curr_time));
BVHReference(bounds, j, object_index, primitive_type, prev_time, curr_time));
root.grow(bounds);
center.grow(bounds.center2());
}
@ -153,18 +157,14 @@ void BVHBuild::add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *m
}
}
void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int i)
void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int object_index)
{
const Attribute *curve_attr_mP = NULL;
if (hair->has_motion_blur()) {
curve_attr_mP = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
}
const PrimitiveType primitive_type =
(curve_attr_mP != NULL) ?
((hair->curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
PRIMITIVE_MOTION_CURVE_THICK) :
((hair->curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK);
const PrimitiveType primitive_type = hair->primitive_type();
const size_t num_curves = hair->num_curves();
for (uint j = 0; j < num_curves; j++) {
@ -177,7 +177,7 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
curve.bounds_grow(k, &hair->get_curve_keys()[0], curve_radius, bounds);
if (bounds.valid()) {
int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
references.push_back(BVHReference(bounds, j, i, packed_type));
references.push_back(BVHReference(bounds, j, object_index, packed_type));
root.grow(bounds);
center.grow(bounds.center2());
}
@ -198,7 +198,7 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
}
if (bounds.valid()) {
int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
references.push_back(BVHReference(bounds, j, i, packed_type));
references.push_back(BVHReference(bounds, j, object_index, packed_type));
root.grow(bounds);
center.grow(bounds.center2());
}
@ -254,7 +254,8 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
if (bounds.valid()) {
const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
int packed_type = PRIMITIVE_PACK_SEGMENT(primitive_type, k);
references.push_back(BVHReference(bounds, j, i, packed_type, prev_time, curr_time));
references.push_back(
BVHReference(bounds, j, object_index, packed_type, prev_time, curr_time));
root.grow(bounds);
center.grow(bounds.center2());
}
@ -268,15 +269,18 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
}
}
void BVHBuild::add_reference_geometry(BoundBox &root, BoundBox &center, Geometry *geom, int i)
void BVHBuild::add_reference_geometry(BoundBox &root,
BoundBox &center,
Geometry *geom,
int object_index)
{
if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
Mesh *mesh = static_cast<Mesh *>(geom);
add_reference_triangles(root, center, mesh, i);
add_reference_triangles(root, center, mesh, object_index);
}
else if (geom->geometry_type == Geometry::HAIR) {
Hair *hair = static_cast<Hair *>(geom);
add_reference_curves(root, center, hair, i);
add_reference_curves(root, center, hair, object_index);
}
}

42
intern/cycles/bvh/bvh_embree.cpp
View File

@ -89,20 +89,9 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
/* Test if we need to record this transparent intersection. */
if (ctx->num_hits < ctx->max_hits || ray->tfar < ctx->max_t) {
/* Skip already recorded intersections. */
int num_recorded_hits = min(ctx->num_hits, ctx->max_hits);
for (int i = 0; i < num_recorded_hits; ++i) {
if (current_isect.object == ctx->isect_s[i].object &&
current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
/* This intersection was already recorded, skip it. */
*args->valid = 0;
return;
}
}
/* If maximum number of hits was reached, replace the intersection with the
* highest distance. We want to find the N closest intersections. */
const int num_recorded_hits = min(ctx->num_hits, ctx->max_hits);
int isect_index = num_recorded_hits;
if (num_recorded_hits + 1 >= ctx->max_hits) {
float max_t = ctx->isect_s[0].t;
@ -147,10 +136,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
}
else {
kernel_embree_convert_hit(kg, ray, hit, &current_isect);
int object = (current_isect.object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, current_isect.prim) :
current_isect.object;
if (ctx->local_object_id != object) {
if (ctx->local_object_id != current_isect.object) {
/* This tells Embree to continue tracing. */
*args->valid = 0;
break;
@ -213,21 +199,11 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
if (ctx->num_hits < ctx->max_hits) {
Intersection current_isect;
kernel_embree_convert_hit(kg, ray, hit, &current_isect);
for (size_t i = 0; i < ctx->num_hits; ++i) {
if (current_isect.object == ctx->isect_s[i].object &&
current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
/* This intersection was already recorded, skip it. */
*args->valid = 0;
break;
}
}
Intersection *isect = &ctx->isect_s[ctx->num_hits];
++ctx->num_hits;
*isect = current_isect;
/* Only primitives from volume object. */
uint tri_object = (isect->object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, isect->prim) :
isect->object;
uint tri_object = isect->object;
int object_flag = kernel_tex_fetch(__object_flag, tri_object);
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
--ctx->num_hits;
@ -249,7 +225,7 @@ static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args)
const RTCRay *ray = (RTCRay *)args->ray;
RTCHit *hit = (RTCHit *)args->hit;
/* Always ignore backfacing intersections. */
/* Always ignore back-facing intersections. */
if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
*args->valid = 0;
@ -262,7 +238,7 @@ static void rtc_filter_occluded_func_thick_curve(const RTCFilterFunctionNArgumen
const RTCRay *ray = (RTCRay *)args->ray;
RTCHit *hit = (RTCHit *)args->hit;
/* Always ignore backfacing intersections. */
/* Always ignore back-facing intersections. */
if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z),
make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) {
*args->valid = 0;
@ -456,7 +432,7 @@ void BVHEmbree::add_instance(Object *ob, int i)
void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
{
size_t prim_offset = mesh->optix_prim_offset;
size_t prim_offset = mesh->prim_offset;
const Attribute *attr_mP = NULL;
size_t num_motion_steps = 1;
@ -625,7 +601,7 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c
void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
{
size_t prim_offset = hair->optix_prim_offset;
size_t prim_offset = hair->curve_segment_offset;
const Attribute *attr_mP = NULL;
size_t num_motion_steps = 1;
@ -702,7 +678,7 @@ void BVHEmbree::refit(Progress &progress)
if (mesh->num_triangles() > 0) {
RTCGeometry geom = rtcGetGeometry(scene, geom_id);
set_tri_vertex_buffer(geom, mesh, true);
rtcSetGeometryUserData(geom, (void *)mesh->optix_prim_offset);
rtcSetGeometryUserData(geom, (void *)mesh->prim_offset);
rtcCommitGeometry(geom);
}
}
@ -711,7 +687,7 @@ void BVHEmbree::refit(Progress &progress)
if (hair->num_curves() > 0) {
RTCGeometry geom = rtcGetGeometry(scene, geom_id + 1);
set_curve_vertex_buffer(geom, hair, true);
rtcSetGeometryUserData(geom, (void *)hair->optix_prim_offset);
rtcSetGeometryUserData(geom, (void *)hair->curve_segment_offset);
rtcCommitGeometry(geom);
}
}

12
intern/cycles/cmake/external_libs.cmake
View File

@ -521,7 +521,7 @@ endif()
if(WITH_CYCLES_CUDA_BINARIES OR NOT WITH_CUDA_DYNLOAD)
find_package(CUDA) # Try to auto locate CUDA toolkit
if(CUDA_FOUND)
message(STATUS "CUDA nvcc = ${CUDA_NVCC_EXECUTABLE}")
message(STATUS "Found CUDA ${CUDA_NVCC_EXECUTABLE} (${CUDA_VERSION})")
else()
message(STATUS "CUDA compiler not found, disabling WITH_CYCLES_CUDA_BINARIES")
set(WITH_CYCLES_CUDA_BINARIES OFF)
@ -537,6 +537,16 @@ endif()
# HIP
###########################################################################
if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
find_package(HIP)
if(HIP_FOUND)
message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
else()
message(STATUS "HIP compiler not found, disabling WITH_CYCLES_HIP_BINARIES")
set(WITH_CYCLES_HIP_BINARIES OFF)
endif()
endif()
if(NOT WITH_HIP_DYNLOAD)
set(WITH_HIP_DYNLOAD ON)
endif()

2
intern/cycles/device/hip/kernel.cpp
View File

@ -28,7 +28,7 @@ void HIPDeviceKernels::load(HIPDevice *device)
for (int i = 0; i < (int)DEVICE_KERNEL_NUM; i++) {
HIPDeviceKernel &kernel = kernels_[i];
/* No megakernel used for GPU. */
/* No mega-kernel used for GPU. */
if (i == DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL) {
continue;
}

23
intern/cycles/device/optix/device_impl.cpp
View File

@ -768,7 +768,13 @@ void OptiXDevice::denoise_color_read(DenoiseContext &context, const DenoisePass
destination.num_components = 3;
destination.pixel_stride = context.buffer_params.pass_stride;
pass_accessor.get_render_tile_pixels(context.render_buffers, context.buffer_params, destination);
BufferParams buffer_params = context.buffer_params;
buffer_params.window_x = 0;
buffer_params.window_y = 0;
buffer_params.window_width = buffer_params.width;
buffer_params.window_height = buffer_params.height;
pass_accessor.get_render_tile_pixels(context.render_buffers, buffer_params, destination);
}
bool OptiXDevice::denoise_filter_color_preprocess(DenoiseContext &context, const DenoisePass &pass)
@ -1246,7 +1252,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
build_input.curveArray.indexBuffer = (CUdeviceptr)index_data.device_pointer;
build_input.curveArray.indexStrideInBytes = sizeof(int);
build_input.curveArray.flag = build_flags;
build_input.curveArray.primitiveIndexOffset = hair->optix_prim_offset;
build_input.curveArray.primitiveIndexOffset = hair->curve_segment_offset;
}
else {
/* Disable visibility test any-hit program, since it is already checked during
@ -1259,7 +1265,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
build_input.customPrimitiveArray.strideInBytes = sizeof(OptixAabb);
build_input.customPrimitiveArray.flags = &build_flags;
build_input.customPrimitiveArray.numSbtRecords = 1;
build_input.customPrimitiveArray.primitiveIndexOffset = hair->optix_prim_offset;
build_input.customPrimitiveArray.primitiveIndexOffset = hair->curve_segment_offset;
}
if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
@ -1328,7 +1334,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
* buffers for that purpose. OptiX does not allow this to be zero though, so just pass in
* one and rely on that having the same meaning in this case. */
build_input.triangleArray.numSbtRecords = 1;
build_input.triangleArray.primitiveIndexOffset = mesh->optix_prim_offset;
build_input.triangleArray.primitiveIndexOffset = mesh->prim_offset;
if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
progress.set_error("Failed to build OptiX acceleration structure");
@ -1395,8 +1401,8 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
instance.transform[5] = 1.0f;
instance.transform[10] = 1.0f;
/* Set user instance ID to object index (but leave low bit blank). */
instance.instanceId = ob->get_device_index() << 1;
/* Set user instance ID to object index. */
instance.instanceId = ob->get_device_index();
/* Add some of the object visibility bits to the mask.
* __prim_visibility contains the combined visibility bits of all instances, so is not
@ -1419,7 +1425,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
}
else {
/* Can disable __anyhit__kernel_optix_visibility_test by default (except for thick curves,
* since it needs to filter out endcaps there).
* since it needs to filter out end-caps there).
* 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.
*/
@ -1508,9 +1514,6 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
else {
/* Disable instance transform if geometry already has it applied to vertex data. */
instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
/* Non-instanced objects read ID from 'prim_object', so distinguish
* them from instanced objects with the low bit set. */
instance.instanceId |= 1;
}
}
}

8
intern/cycles/integrator/denoiser_oidn.cpp
View File

@ -289,7 +289,13 @@ class OIDNDenoiseContext {
* pixels. */
const PassAccessorCPU pass_accessor(pass_access_info, 1.0f, num_samples_);
pass_accessor.get_render_tile_pixels(render_buffers_, buffer_params_, destination);
BufferParams buffer_params = buffer_params_;
buffer_params.window_x = 0;
buffer_params.window_y = 0;
buffer_params.window_width = buffer_params.width;
buffer_params.window_height = buffer_params.height;
pass_accessor.get_render_tile_pixels(render_buffers_, buffer_params, destination);
}
/* Read pass pixels using PassAccessor into a temporary buffer which is owned by the pass.. */

3
intern/cycles/integrator/pass_accessor.cpp
View File

@ -149,9 +149,6 @@ bool PassAccessor::get_render_tile_pixels(const RenderBuffers *render_buffers,
/* Denoised passes store their final pixels, no need in special calculation. */
get_pass_float(render_buffers, buffer_params, destination);
}
else if (type == PASS_RENDER_TIME) {
/* TODO(sergey): Needs implementation. */
}
else if (type == PASS_DEPTH) {
get_pass_depth(render_buffers, buffer_params, destination);
}

45
intern/cycles/integrator/pass_accessor_cpu.cpp
View File

@ -99,17 +99,22 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_float(
{
DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
const float *buffer_data = render_buffers->buffer.data();
const int64_t pass_stride = buffer_params.pass_stride;
const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
const float *window_data = render_buffers->buffer.data() + buffer_params.window_x * pass_stride +
buffer_params.window_y * buffer_row_stride;
const int pixel_stride = destination.pixel_stride ? destination.pixel_stride :
destination.num_components;
tbb::parallel_for(0, buffer_params.height, [&](int64_t y) {
int64_t pixel_index = y * buffer_params.width;
for (int64_t x = 0; x < buffer_params.width; ++x, ++pixel_index) {
const int64_t input_pixel_offset = pixel_index * buffer_params.pass_stride;
const float *buffer = buffer_data + input_pixel_offset;
float *pixel = destination.pixels + (pixel_index + destination.offset) * pixel_stride;
tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
const float *buffer = window_data + y * buffer_row_stride;
float *pixel = destination.pixels +
(y * buffer_params.width + destination.offset) * pixel_stride;
for (int64_t x = 0; x < buffer_params.window_width;
++x, buffer += pass_stride, pixel += pixel_stride) {
processor(kfilm_convert, buffer, pixel);
}
});
@ -123,26 +128,28 @@ inline void PassAccessorCPU::run_get_pass_kernel_processor_half_rgba(
const Destination &destination,
const Processor &processor) const
{
const float *buffer_data = render_buffers->buffer.data();
const int64_t pass_stride = buffer_params.pass_stride;
const int64_t buffer_row_stride = buffer_params.stride * buffer_params.pass_stride;
const float *window_data = render_buffers->buffer.data() + buffer_params.window_x * pass_stride +
buffer_params.window_y * buffer_row_stride;
half4 *dst_start = destination.pixels_half_rgba + destination.offset;
const int destination_stride = destination.stride != 0 ? destination.stride :
buffer_params.width;
tbb::parallel_for(0, buffer_params.height, [&](int64_t y) {
int64_t pixel_index = y * buffer_params.width;
half4 *dst_row_start = dst_start + y * destination_stride;
for (int64_t x = 0; x < buffer_params.width; ++x, ++pixel_index) {
const int64_t input_pixel_offset = pixel_index * buffer_params.pass_stride;
const float *buffer = buffer_data + input_pixel_offset;
tbb::parallel_for(0, buffer_params.window_height, [&](int64_t y) {
const float *buffer = window_data + y * buffer_row_stride;
half4 *pixel = dst_start + y * destination_stride;
for (int64_t x = 0; x < buffer_params.window_width; ++x, buffer += pass_stride, ++pixel) {
float pixel[4];
processor(kfilm_convert, buffer, pixel);
float pixel_rgba[4];
processor(kfilm_convert, buffer, pixel_rgba);
film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel_rgba);
half4 *pixel_half_rgba = dst_row_start + x;
float4_store_half(&pixel_half_rgba->x, make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
float4_store_half(&pixel->x,
make_float4(pixel_rgba[0], pixel_rgba[1], pixel_rgba[2], pixel_rgba[3]));
}
});
}

15
intern/cycles/integrator/pass_accessor_gpu.cpp
View File

@ -43,10 +43,13 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
KernelFilmConvert kfilm_convert;
init_kernel_film_convert(&kfilm_convert, buffer_params, destination);
const int work_size = buffer_params.width * buffer_params.height;
const int work_size = buffer_params.window_width * buffer_params.window_height;
const int destination_stride = destination.stride != 0 ? destination.stride :
buffer_params.width;
buffer_params.window_width;
const int offset = buffer_params.window_x * buffer_params.pass_stride +
buffer_params.window_y * buffer_params.stride * buffer_params.pass_stride;
if (destination.d_pixels) {
DCHECK_EQ(destination.stride, 0) << "Custom stride for float destination is not implemented.";
@ -55,8 +58,8 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
const_cast<device_ptr *>(&destination.d_pixels),
const_cast<device_ptr *>(&render_buffers->buffer.device_pointer),
const_cast<int *>(&work_size),
const_cast<int *>(&buffer_params.width),
const_cast<int *>(&buffer_params.offset),
const_cast<int *>(&buffer_params.window_width),
const_cast<int *>(&offset),
const_cast<int *>(&buffer_params.stride),
const_cast<int *>(&destination.offset),
const_cast<int *>(&destination_stride)};
@ -70,8 +73,8 @@ void PassAccessorGPU::run_film_convert_kernels(DeviceKernel kernel,
const_cast<device_ptr *>(&destination.d_pixels_half_rgba),
const_cast<device_ptr *>(&render_buffers->buffer.device_pointer),
const_cast<int *>(&work_size),
const_cast<int *>(&buffer_params.width),
const_cast<int *>(&buffer_params.offset),
const_cast<int *>(&buffer_params.window_width),
const_cast<int *>(&offset),
const_cast<int *>(&buffer_params.stride),
const_cast<int *>(&destination.offset),
const_cast<int *>(&destination_stride)};

14
intern/cycles/integrator/path_trace.cpp
View File

@ -282,6 +282,12 @@ static BufferParams scale_buffer_params(const BufferParams &params, int resoluti
scaled_params.width = max(1, params.width / resolution_divider);
scaled_params.height = max(1, params.height / resolution_divider);
scaled_params.window_x = params.window_x / resolution_divider;
scaled_params.window_y = params.window_y / resolution_divider;
scaled_params.window_width = params.window_width / resolution_divider;
scaled_params.window_height = params.window_height / resolution_divider;
scaled_params.full_x = params.full_x / resolution_divider;
scaled_params.full_y = params.full_y / resolution_divider;
scaled_params.full_width = params.full_width / resolution_divider;
@ -1005,12 +1011,12 @@ bool PathTrace::set_render_tile_pixels(PassAccessor &pass_accessor,
int2 PathTrace::get_render_tile_size() const
{
if (full_frame_state_.render_buffers) {
return make_int2(full_frame_state_.render_buffers->params.width,
full_frame_state_.render_buffers->params.height);
return make_int2(full_frame_state_.render_buffers->params.window_width,
full_frame_state_.render_buffers->params.window_height);
}
const Tile &tile = tile_manager_.get_current_tile();
return make_int2(tile.width, tile.height);
return make_int2(tile.window_width, tile.window_height);
}
int2 PathTrace::get_render_tile_offset() const
@ -1020,7 +1026,7 @@ int2 PathTrace::get_render_tile_offset() const
}
const Tile &tile = tile_manager_.get_current_tile();
return make_int2(tile.x, tile.y);
return make_int2(tile.x + tile.window_x, tile.y + tile.window_y);
}
int2 PathTrace::get_render_size() const

6
intern/cycles/integrator/path_trace_work.cpp
View File

@ -191,8 +191,10 @@ PassAccessor::Destination PathTraceWork::get_display_destination_template(
PassAccessor::Destination destination(film_->get_display_pass());
const int2 display_texture_size = display->get_texture_size();
const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x;
const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y;
const int texture_x = effective_buffer_params_.full_x - effective_full_params_.full_x +
effective_buffer_params_.window_x;
const int texture_y = effective_buffer_params_.full_y - effective_full_params_.full_y +
effective_buffer_params_.window_y;
destination.offset = texture_y * display_texture_size.x + texture_x;
destination.stride = display_texture_size.x;

53
intern/cycles/integrator/path_trace_work_gpu.cpp
View File

@ -23,6 +23,7 @@
#include "render/buffers.h"
#include "render/scene.h"
#include "util/util_logging.h"
#include "util/util_string.h"
#include "util/util_tbb.h"
#include "util/util_time.h"
@ -30,6 +31,33 @@
CCL_NAMESPACE_BEGIN
static size_t estimate_single_state_size(DeviceScene *device_scene)
{
size_t state_size = 0;
#define KERNEL_STRUCT_BEGIN(name) for (int array_index = 0;; array_index++) {
#define KERNEL_STRUCT_MEMBER(parent_struct, type, name, feature) state_size += sizeof(type);
#define KERNEL_STRUCT_ARRAY_MEMBER(parent_struct, type, name, feature) state_size += sizeof(type);
#define KERNEL_STRUCT_END(name) \
break; \
}
#define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
if (array_index == gpu_array_size - 1) { \
break; \
} \
}
#define KERNEL_STRUCT_VOLUME_STACK_SIZE (device_scene->data.volume_stack_size)
#include "kernel/integrator/integrator_state_template.h"
#undef KERNEL_STRUCT_BEGIN
#undef KERNEL_STRUCT_MEMBER
#undef KERNEL_STRUCT_ARRAY_MEMBER
#undef KERNEL_STRUCT_END
#undef KERNEL_STRUCT_END_ARRAY
#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
return state_size;
}
PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
Film *film,
DeviceScene *device_scene,
@ -47,7 +75,7 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
num_queued_paths_(device, "num_queued_paths", MEM_READ_WRITE),
work_tiles_(device, "work_tiles", MEM_READ_WRITE),
display_rgba_half_(device, "display buffer half", MEM_READ_WRITE),
max_num_paths_(queue_->num_concurrent_states(sizeof(IntegratorStateCPU))),
max_num_paths_(queue_->num_concurrent_states(estimate_single_state_size(device_scene))),
min_num_active_paths_(queue_->num_concurrent_busy_states()),
max_active_path_index_(0)
{
@ -100,12 +128,23 @@ void PathTraceWorkGPU::alloc_integrator_soa()
break; \
} \
}
#define KERNEL_STRUCT_VOLUME_STACK_SIZE (device_scene_->data.volume_stack_size)
#include "kernel/integrator/integrator_state_template.h"
#undef KERNEL_STRUCT_BEGIN
#undef KERNEL_STRUCT_MEMBER
#undef KERNEL_STRUCT_ARRAY_MEMBER
#undef KERNEL_STRUCT_END
#undef KERNEL_STRUCT_END_ARRAY
#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
if (VLOG_IS_ON(3)) {
size_t total_soa_size = 0;
for (auto &&soa_memory : integrator_state_soa_) {
total_soa_size += soa_memory->memory_size();
}
VLOG(3) << "GPU SoA state size: " << string_human_readable_size(total_soa_size);
}
}
void PathTraceWorkGPU::alloc_integrator_queue()
@ -712,13 +751,13 @@ void PathTraceWorkGPU::copy_to_display_naive(PathTraceDisplay *display,
{
const int full_x = effective_buffer_params_.full_x;
const int full_y = effective_buffer_params_.full_y;
const int width = effective_buffer_params_.width;
const int height = effective_buffer_params_.height;
const int final_width = buffers_->params.width;
const int final_height = buffers_->params.height;
const int width = effective_buffer_params_.window_width;
const int height = effective_buffer_params_.window_height;
const int final_width = buffers_->params.window_width;
const int final_height = buffers_->params.window_height;
const int texture_x = full_x - effective_full_params_.full_x;
const int texture_y = full_y - effective_full_params_.full_y;
const int texture_x = full_x - effective_full_params_.full_x + effective_buffer_params_.window_x;
const int texture_y = full_y - effective_full_params_.full_y + effective_buffer_params_.window_y;
/* Re-allocate display memory if needed, and make sure the device pointer is allocated.
*

60
intern/cycles/kernel/CMakeLists.txt
View File

@ -404,16 +404,27 @@ if(WITH_CYCLES_CUDA_BINARIES)
-cuda-toolkit-dir "${cuda_toolkit_root_dir}"
DEPENDS ${kernel_sources} cycles_cubin_cc)
else()
add_custom_command(
OUTPUT ${cuda_file}
COMMAND ${cuda_nvcc_executable}
set(_cuda_nvcc_args
-arch=${arch}
${CUDA_NVCC_FLAGS}
--${format}
${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
--ptxas-options="-v"
${cuda_flags}
DEPENDS ${kernel_sources})
${cuda_flags})
if(WITH_COMPILER_CCACHE AND CCACHE_PROGRAM)
add_custom_command(
OUTPUT ${cuda_file}
COMMAND ${CCACHE_PROGRAM} ${cuda_nvcc_executable} ${_cuda_nvcc_args}
DEPENDS ${kernel_sources})
else()
add_custom_command(
OUTPUT ${cuda_file}
COMMAND ${cuda_nvcc_executable} ${_cuda_nvcc_args}
DEPENDS ${kernel_sources})
endif()
unset(_cuda_nvcc_args)
endif()
delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_file}" ${CYCLES_INSTALL_PATH}/lib)
list(APPEND cuda_cubins ${cuda_file})
@ -472,20 +483,10 @@ endif()
# HIP module
if(WITH_CYCLES_HIP_BINARIES)
if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
# 64 bit only
set(HIP_BITS 64)
# HIP version
execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} "--version" OUTPUT_VARIABLE HIPCC_OUT)
string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\1" HIP_VERSION_MAJOR "${HIPCC_OUT}")
string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\2" HIP_VERSION_MINOR "${HIPCC_OUT}")
set(HIP_VERSION "${HIP_VERSION_MAJOR}${HIP_VERSION_MINOR}")
message(WARNING
"HIP version ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR} detected")
# build for each arch
set(hip_sources device/hip/kernel.cpp
${SRC_HEADERS}
@ -542,23 +543,24 @@ if(WITH_CYCLES_HIP_BINARIES)
-D WITH_NANOVDB
-I "${NANOVDB_INCLUDE_DIR}")
endif()
add_custom_command(
OUTPUT ${hip_file}
COMMAND ${HIP_HIPCC_EXECUTABLE}
-arch=${arch}
${HIP_HIPCC_FLAGS}
--${format}
${CMAKE_CURRENT_SOURCE_DIR}${hip_kernel_src}
${hip_flags}
DEPENDS ${kernel_sources})
delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${hip_file}" ${CYCLES_INSTALL_PATH}/lib)
list(APPEND hip_fatbins ${hip_file})
endmacro()
set(prev_arch "none")
foreach(arch ${CYCLES_HIP_BINARIES_ARCH})
set(hip_hipcc_executable ${HIP_HIPCC_EXECUTABLE})
set(hip_toolkit_root_dir ${HIP_TOOLKIT_ROOT_DIR})
if(DEFINED hip_hipcc_executable AND DEFINED hip_toolkit_root_dir)
# Compile regular kernel
CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
if(WITH_CYCLES_HIP_BUILD_SERIAL)
set(prev_arch ${arch})
endif()
unset(hip_hipcc_executable)
unset(hip_toolkit_root_dir)
endif()
# Compile regular kernel
CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
endforeach()
add_custom_target(cycles_kernel_hip ALL DEPENDS ${hip_fatbins})

33
intern/cycles/kernel/bvh/bvh_embree.h
View File

@ -106,9 +106,6 @@ ccl_device_inline void kernel_embree_convert_hit(const KernelGlobals *kg,
const RTCHit *hit,
Intersection *isect)
{
bool is_hair = hit->geomID & 1;
isect->u = is_hair ? hit->u : 1.0f - hit->v - hit->u;
isect->v = is_hair ? hit->v : hit->u;
isect->t = ray->tfar;
isect->Ng = make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z);
if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) {
@ -121,27 +118,37 @@ ccl_device_inline void kernel_embree_convert_hit(const KernelGlobals *kg,
else {
isect->prim = hit->primID + (intptr_t)rtcGetGeometryUserData(
rtcGetGeometry(kernel_data.bvh.scene, hit->geomID));
isect->object = OBJECT_NONE;
isect->object = hit->geomID / 2;
}
const bool is_hair = hit->geomID & 1;
if (is_hair) {
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, isect->prim);
isect->type = segment.type;
isect->prim = segment.prim;
isect->u = hit->u;
isect->v = hit->v;
}
else {
isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
isect->u = 1.0f - hit->v - hit->u;
isect->v = hit->u;
}
isect->type = kernel_tex_fetch(__prim_type, isect->prim);
}
ccl_device_inline void kernel_embree_convert_sss_hit(const KernelGlobals *kg,
const RTCRay *ray,
const RTCHit *hit,
Intersection *isect,
int local_object_id)
ccl_device_inline void kernel_embree_convert_sss_hit(
const KernelGlobals *kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object)
{
isect->u = 1.0f - hit->v - hit->u;
isect->v = hit->u;
isect->t = ray->tfar;
isect->Ng = make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z);
RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
rtcGetGeometry(kernel_data.bvh.scene, local_object_id * 2));
rtcGetGeometry(kernel_data.bvh.scene, object * 2));
isect->prim = hit->primID +
(intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
isect->object = local_object_id;
isect->type = kernel_tex_fetch(__prim_type, isect->prim);
isect->object = object;
isect->type = kernel_tex_fetch(__objects, object).primitive_type;
}
CCL_NAMESPACE_END

31
intern/cycles/kernel/bvh/bvh_shadow_all.h
View File

@ -130,7 +130,6 @@ ccl_device_inline
if (prim_addr >= 0) {
const int prim_addr2 = __float_as_int(leaf.y);
const uint type = __float_as_int(leaf.w);
const uint p_type = type & PRIMITIVE_ALL;
/* pop */
node_addr = traversal_stack[stack_ptr];
@ -138,14 +137,15 @@ ccl_device_inline
/* primitive intersection */
while (prim_addr < prim_addr2) {
kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) == p_type);
kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
(type & PRIMITIVE_ALL));
bool hit;
/* todo: specialized intersect functions which don't fill in
* isect unless needed and check SD_HAS_TRANSPARENT_SHADOW?
* might give a few % performance improvement */
switch (p_type) {
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {
hit = triangle_intersect(
kg, isect, P, dir, isect_t, visibility, object, prim_addr);
@ -163,17 +163,20 @@ ccl_device_inline
case PRIMITIVE_MOTION_CURVE_THICK:
case PRIMITIVE_CURVE_RIBBON:
case PRIMITIVE_MOTION_CURVE_RIBBON: {
const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
hit = curve_intersect(kg,
isect,
P,
dir,
isect_t,
visibility,
object,
prim_addr,
ray->time,
curve_type);
if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) {
hit = false;
break;
}
}
const int curve_object = kernel_tex_fetch(__prim_object, prim_addr);
const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
hit = curve_intersect(
kg, isect, P, dir, isect_t, curve_object, curve_prim, ray->time, curve_type);
break;
}
#endif

24
intern/cycles/kernel/bvh/bvh_traversal.h
View File

@ -165,18 +165,18 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(const KernelGlobals *kg,
case PRIMITIVE_CURVE_RIBBON:
case PRIMITIVE_MOTION_CURVE_RIBBON: {
for (; prim_addr < prim_addr2; prim_addr++) {
const uint curve_type = kernel_tex_fetch(__prim_type, prim_addr);
kernel_assert((curve_type & PRIMITIVE_ALL) == (type & PRIMITIVE_ALL));
const bool hit = curve_intersect(kg,
isect,
P,
dir,
isect->t,
visibility,
object,
prim_addr,
ray->time,
curve_type);
if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
if (ray->time < prim_time.x || ray->time > prim_time.y) {
continue;
}
}
const int curve_object = kernel_tex_fetch(__prim_object, prim_addr);
const int curve_prim = kernel_tex_fetch(__prim_index, prim_addr);
const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
const bool hit = curve_intersect(
kg, isect, P, dir, isect->t, curve_object, curve_prim, ray->time, curve_type);
if (hit) {
/* shadow ray early termination */
if (visibility & PATH_RAY_SHADOW_OPAQUE)

31
intern/cycles/kernel/bvh/bvh_util.h
View File

@ -118,19 +118,18 @@ ccl_device_inline void sort_intersections(Intersection *hits, uint num_hits)
ccl_device_forceinline int intersection_get_shader_flags(const KernelGlobals *ccl_restrict kg,
const Intersection *ccl_restrict isect)
{
const int prim = kernel_tex_fetch(__prim_index, isect->prim);
const int prim = isect->prim;
int shader = 0;
#ifdef __HAIR__
if (kernel_tex_fetch(__prim_type, isect->prim) & PRIMITIVE_ALL_TRIANGLE)
if (isect->type & PRIMITIVE_ALL_TRIANGLE)
#endif
{
shader = kernel_tex_fetch(__tri_shader, prim);
}
#ifdef __HAIR__
else {
float4 str = kernel_tex_fetch(__curves, prim);
shader = __float_as_int(str.z);
shader = kernel_tex_fetch(__curves, prim).shader_id;
}
#endif
@ -138,21 +137,19 @@ ccl_device_forceinline int intersection_get_shader_flags(const KernelGlobals *cc
}
ccl_device_forceinline int intersection_get_shader_from_isect_prim(
const KernelGlobals *ccl_restrict kg, const int isect_prim)
const KernelGlobals *ccl_restrict kg, const int prim, const int isect_type)
{
const int prim = kernel_tex_fetch(__prim_index, isect_prim);
int shader = 0;
#ifdef __HAIR__
if (kernel_tex_fetch(__prim_type, isect_prim) & PRIMITIVE_ALL_TRIANGLE)
if (isect_type & PRIMITIVE_ALL_TRIANGLE)
#endif
{
shader = kernel_tex_fetch(__tri_shader, prim);
}
#ifdef __HAIR__
else {
float4 str = kernel_tex_fetch(__curves, prim);
shader = __float_as_int(str.z);
shader = kernel_tex_fetch(__curves, prim).shader_id;
}
#endif
@ -162,25 +159,13 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(
ccl_device_forceinline int intersection_get_shader(const KernelGlobals *ccl_restrict kg,
const Intersection *ccl_restrict isect)
{
return intersection_get_shader_from_isect_prim(kg, isect->prim);
}
ccl_device_forceinline int intersection_get_object(const KernelGlobals *ccl_restrict kg,
const Intersection *ccl_restrict isect)
{
if (isect->object != OBJECT_NONE) {
return isect->object;
}
return kernel_tex_fetch(__prim_object, isect->prim);
return intersection_get_shader_from_isect_prim(kg, isect->prim, isect->type);
}
ccl_device_forceinline int intersection_get_object_flags(const KernelGlobals *ccl_restrict kg,
const Intersection *ccl_restrict isect)
{
const int object = intersection_get_object(kg, isect);
return kernel_tex_fetch(__object_flag, object);
return kernel_tex_fetch(__object_flag, isect->object);
}
CCL_NAMESPACE_END

25
intern/cycles/kernel/closure/bsdf_principled_diffuse.h
View File

@ -35,21 +35,25 @@ static_assert(sizeof(ShaderClosure) >= sizeof(PrincipledDiffuseBsdf),
"PrincipledDiffuseBsdf is too large!");
ccl_device float3 calculate_principled_diffuse_brdf(
const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float3 H, float *pdf)
const PrincipledDiffuseBsdf *bsdf, float3 N, float3 V, float3 L, float *pdf)
{
float NdotL = dot(N, L);
float NdotV = dot(N, V);
if (NdotL <= 0 || NdotV <= 0) {
*pdf = 0.0f;
if (NdotL <= 0) {
return make_float3(0.0f, 0.0f, 0.0f);
}
float LdotH = dot(L, H);
float NdotV = dot(N, V);
/* H = normalize(L + V); // Bissector of an angle between L and V
* LH2 = 2 * dot(L, H)^2 = 2cos(x)^2 = cos(2x) + 1 = dot(L, V) + 1,
* half-angle x between L and V is at most 90 deg
*/
float LH2 = dot(L, V) + 1;
float FL = schlick_fresnel(NdotL), FV = schlick_fresnel(NdotV);
const float Fd90 = 0.5f + 2.0f * LdotH * LdotH * bsdf->roughness;
float Fd = (1.0f * (1.0f - FL) + Fd90 * FL) * (1.0f * (1.0f - FV) + Fd90 * FV);
const float Fd90 = 0.5f + LH2 * bsdf->roughness;
float Fd = (1.0f - FL + Fd90 * FL) * (1.0f - FV + Fd90 * FV);
float value = M_1_PI_F * NdotL * Fd;
@ -72,11 +76,10 @@ ccl_device float3 bsdf_principled_diffuse_eval_reflect(const ShaderClosure *sc,
float3 N = bsdf->N;
float3 V = I; // outgoing
float3 L = omega_in; // incoming
float3 H = normalize(L + V);
if (dot(N, omega_in) > 0.0f) {
*pdf = fmaxf(dot(N, omega_in), 0.0f) * M_1_PI_F;
return calculate_principled_diffuse_brdf(bsdf, N, V, L, H, pdf);
return calculate_principled_diffuse_brdf(bsdf, N, V, L, pdf);
}
else {
*pdf = 0.0f;
@ -112,9 +115,7 @@ ccl_device int bsdf_principled_diffuse_sample(const ShaderClosure *sc,
sample_cos_hemisphere(N, randu, randv, omega_in, pdf);
if (dot(Ng, *omega_in) > 0) {
float3 H = normalize(I + *omega_in);
*eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, H, pdf);
*eval = calculate_principled_diffuse_brdf(bsdf, N, I, *omega_in, pdf);
#ifdef __RAY_DIFFERENTIALS__
// TODO: find a better approximation for the diffuse bounce

18
intern/cycles/kernel/device/gpu/kernel.h
View File

@ -424,8 +424,12 @@ ccl_device_inline void kernel_gpu_film_convert_common(const KernelFilmConvert *k
return;
}
const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kfilm_convert->pass_stride;
ccl_global const float *buffer = render_buffer + render_buffer_offset;
const int x = render_pixel_index % width;
const int y = render_pixel_index / width;
ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert->pass_stride +
y * stride * kfilm_convert->pass_stride;
ccl_global float *pixel = pixels +
(render_pixel_index + dst_offset) * kfilm_convert->pixel_stride;
@ -451,17 +455,17 @@ ccl_device_inline void kernel_gpu_film_convert_half_rgba_common_rgba(
return;
}
const uint64_t render_buffer_offset = (uint64_t)render_pixel_index * kfilm_convert->pass_stride;
ccl_global const float *buffer = render_buffer + render_buffer_offset;
const int x = render_pixel_index % width;
const int y = render_pixel_index / width;
ccl_global const float *buffer = render_buffer + offset + x * kfilm_convert->pass_stride +
y * stride * kfilm_convert->pass_stride;
float pixel[4];
processor(kfilm_convert, buffer, pixel);
film_apply_pass_pixel_overlays_rgba(kfilm_convert, buffer, pixel);
const int x = render_pixel_index % width;
const int y = render_pixel_index / width;
ccl_global half4 *out = ((ccl_global half4 *)rgba) + rgba_offset + y * rgba_stride + x;
float4_store_half((ccl_global half *)out, make_float4(pixel[0], pixel[1], pixel[2], pixel[3]));
}

86
intern/cycles/kernel/device/optix/kernel.cu
View File

@ -41,22 +41,15 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
return (T *)(((uint64_t)optixGetPayload_3() << 32) | optixGetPayload_2());
}
template<bool always = false> ccl_device_forceinline uint get_object_id()
ccl_device_forceinline int get_object_id()
{
#ifdef __OBJECT_MOTION__
/* Always get the the instance ID from the TLAS.
/* Always get the the instance ID from the TLAS
* There might be a motion transform node between TLAS and BLAS which does not have one. */
uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
#else
uint object = optixGetInstanceId();
return optixGetInstanceId();
#endif
/* Choose between always returning object ID or only for instances. */
if (always || (object & 1) == 0)
/* Can just remove the low bit since instance always contains object ID. */
return object >> 1;
else
/* Set to OBJECT_NONE if this is not an instanced object. */
return OBJECT_NONE;
}
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
@ -108,7 +101,7 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
#endif
#ifdef __BVH_LOCAL__
const uint object = get_object_id<true>();
const int object = get_object_id();
if (object != optixGetPayload_4() /* local_object */) {
/* Only intersect with matching object. */
return optixIgnoreIntersection();
@ -152,21 +145,23 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
local_isect->num_hits = 1;
}
const int prim = optixGetPrimitiveIndex();
Intersection *isect = &local_isect->hits[hit];
isect->t = optixGetRayTmax();
isect->prim = optixGetPrimitiveIndex();
isect->prim = prim;
isect->object = get_object_id();
isect->type = kernel_tex_fetch(__prim_type, isect->prim);
isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
const float2 barycentrics = optixGetTriangleBarycentrics();
isect->u = 1.0f - barycentrics.y - barycentrics.x;
isect->v = barycentrics.x;
/* Record geometric normal. */
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_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 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));
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). */
@ -179,25 +174,32 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
#ifdef __SHADOW_RECORD_ALL__
bool ignore_intersection = false;
const uint prim = optixGetPrimitiveIndex();
int prim = optixGetPrimitiveIndex();
const uint object = get_object_id();
# ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4();
if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
ignore_intersection = true;
}
# endif
float u = 0.0f, v = 0.0f;
int type = 0;
if (optixIsTriangleHit()) {
const float2 barycentrics = optixGetTriangleBarycentrics();
u = 1.0f - barycentrics.y - barycentrics.x;
v = barycentrics.x;
type = kernel_tex_fetch(__objects, object).primitive_type;
}
# ifdef __HAIR__
else {
u = __uint_as_float(optixGetAttribute_0());
v = __uint_as_float(optixGetAttribute_1());
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
type = segment.type;
prim = segment.prim;
/* Filter out curve endcaps. */
if (u == 0.0f || u == 1.0f) {
ignore_intersection = true;
@ -245,8 +247,8 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
isect->v = v;
isect->t = optixGetRayTmax();
isect->prim = prim;
isect->object = get_object_id();
isect->type = kernel_tex_fetch(__prim_type, prim);
isect->object = object;
isect->type = type;
# ifdef __TRANSPARENT_SHADOWS__
/* Detect if this surface has a shader with transparent shadows. */
@ -274,15 +276,14 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
}
#endif
const uint object = get_object_id();
#ifdef __VISIBILITY_FLAG__
const uint prim = optixGetPrimitiveIndex();
const uint visibility = optixGetPayload_4();
if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();
}
#endif
const uint object = get_object_id<true>();
if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
return optixIgnoreIntersection();
}
@ -301,9 +302,9 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
#endif
#ifdef __VISIBILITY_FLAG__
const uint prim = optixGetPrimitiveIndex();
const uint object = get_object_id();
const uint visibility = optixGetPayload_4();
if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
return optixIgnoreIntersection();
}
@ -316,28 +317,39 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
extern "C" __global__ void __closesthit__kernel_optix_hit()
{
const int object = get_object_id();
const int prim = optixGetPrimitiveIndex();
optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */
optixSetPayload_3(optixGetPrimitiveIndex());
optixSetPayload_4(get_object_id());
/* Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index. */
optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex()));
optixSetPayload_4(object);
if (optixIsTriangleHit()) {
const float2 barycentrics = optixGetTriangleBarycentrics();
optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
optixSetPayload_2(__float_as_uint(barycentrics.x));
optixSetPayload_3(prim);
optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
}
else {
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
optixSetPayload_2(optixGetAttribute_1());
optixSetPayload_3(segment.prim);
optixSetPayload_5(segment.type);
}
}
#ifdef __HAIR__
ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)
ccl_device_inline void optix_intersection_curve(const int prim, const int type)
{
const uint object = get_object_id<true>();
const int object = get_object_id();
# ifdef __VISIBILITY_FLAG__
const uint visibility = optixGetPayload_4();
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
return;
}
# endif
float3 P = optixGetObjectRayOrigin();
float3 dir = optixGetObjectRayDirection();
@ -358,7 +370,7 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
if (isect.t != FLT_MAX)
isect.t *= len;
if (curve_intersect(NULL, &isect, P, dir, isect.t, visibility, object, prim, time, type)) {
if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
optixReportIntersection(isect.t / len,
type & PRIMITIVE_ALL,
__float_as_int(isect.u), /* Attribute_0 */
@ -368,9 +380,9 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
extern "C" __global__ void __intersection__curve_ribbon()
{
const uint prim = optixGetPrimitiveIndex();
const uint type = kernel_tex_fetch(__prim_type, prim);
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
const int prim = segment.prim;
const int type = segment.type;
if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
optix_intersection_curve(prim, type);
}

24
intern/cycles/kernel/geom/geom_curve.h
View File

@ -34,8 +34,8 @@ ccl_device float curve_attribute_float(const KernelGlobals *kg,
float *dy)
{
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float f0 = kernel_tex_fetch(__attributes_float, desc.offset + k0);
@ -76,8 +76,8 @@ ccl_device float2 curve_attribute_float2(const KernelGlobals *kg,
float2 *dy)
{
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float2 f0 = kernel_tex_fetch(__attributes_float2, desc.offset + k0);
@ -122,8 +122,8 @@ ccl_device float3 curve_attribute_float3(const KernelGlobals *kg,
float3 *dy)
{
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
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));
@ -164,8 +164,8 @@ ccl_device float4 curve_attribute_float4(const KernelGlobals *kg,
float4 *dy)
{
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float4 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
@ -206,8 +206,8 @@ ccl_device float curve_thickness(const KernelGlobals *kg, const ShaderData *sd)
float r = 0.0f;
if (sd->type & PRIMITIVE_ALL_CURVE) {
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float4 P_curve[2];
@ -231,8 +231,8 @@ ccl_device float curve_thickness(const KernelGlobals *kg, const ShaderData *sd)
ccl_device float3 curve_motion_center_location(const KernelGlobals *kg, const ShaderData *sd)
{
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
int k0 = __float_as_int(curvedata.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float4 P_curve[2];

55
intern/cycles/kernel/geom/geom_curve_intersect.h
View File

@ -630,33 +630,19 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
const float3 P,
const float3 dir,
const float tmax,
uint visibility,
int object,
int curveAddr,
int prim,
float time,
int type)
{
const bool is_motion = (type & PRIMITIVE_ALL_MOTION);
# ifndef __KERNEL_OPTIX__ /* See OptiX motion flag OPTIX_MOTION_FLAG_[START|END]_VANISH */
if (is_motion && kernel_data.bvh.use_bvh_steps) {
const float2 prim_time = kernel_tex_fetch(__prim_time, curveAddr);
if (time < prim_time.x || time > prim_time.y) {
return false;
}
}
# endif
KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
int segment = PRIMITIVE_UNPACK_SEGMENT(type);
int prim = kernel_tex_fetch(__prim_index, curveAddr);
float4 v00 = kernel_tex_fetch(__curves, prim);
int k0 = __float_as_int(v00.x) + segment;
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
int k1 = k0 + 1;
int ka = max(k0 - 1, __float_as_int(v00.x));
int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
int ka = max(k0 - 1, kcurve.first_key);
int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
float4 curve[4];
if (!is_motion) {
@ -666,21 +652,14 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
curve[3] = kernel_tex_fetch(__curve_keys, kb);
}
else {
int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
motion_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, curve);
motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
}
# ifdef __VISIBILITY_FLAG__
if (!(kernel_tex_fetch(__prim_visibility, curveAddr) & visibility)) {
return false;
}
# endif
if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
/* todo: adaptive number of subdivisions could help performance here. */
const int subdivisions = kernel_data.bvh.curve_subdivisions;
if (ribbon_intersect(P, dir, tmax, subdivisions, curve, isect)) {
isect->prim = curveAddr;
isect->prim = prim;
isect->object = object;
isect->type = type;
return true;
@ -690,7 +669,7 @@ ccl_device_forceinline bool curve_intersect(const KernelGlobals *kg,
}
else {
if (curve_intersect_recursive(P, dir, tmax, curve, isect)) {
isect->prim = curveAddr;
isect->prim = prim;
isect->object = object;
isect->type = type;
return true;
@ -708,7 +687,7 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
const int isect_object,
const int isect_prim)
{
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
@ -716,14 +695,12 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
D = safe_normalize_len(D, &t);
}
int prim = kernel_tex_fetch(__prim_index, isect_prim);
float4 v00 = kernel_tex_fetch(__curves, prim);
KernelCurve kcurve = kernel_tex_fetch(__curves, isect_prim);
int k0 = __float_as_int(v00.x) + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
int ka = max(k0 - 1, __float_as_int(v00.x));
int kb = min(k1 + 1, __float_as_int(v00.x) + __float_as_int(v00.y) - 1);
int ka = max(k0 - 1, kcurve.first_key);
int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
float4 P_curve[4];
@ -780,15 +757,13 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
sd->dPdv = cross(dPdu, sd->Ng);
# endif
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}
sd->P = P;
float4 curvedata = kernel_tex_fetch(__curves, sd->prim);
sd->shader = __float_as_int(curvedata.z);
sd->shader = kernel_tex_fetch(__curves, sd->prim).shader_id;
}
#endif

6
intern/cycles/kernel/geom/geom_motion_triangle.h
View File

@ -72,9 +72,9 @@ ccl_device_inline void motion_triangle_verts_for_step(const KernelGlobals *kg,
{
if (step == numsteps) {
/* center step: regular vertex location */
verts[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
verts[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
verts[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
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));
}
else {
/* center step not store in this array */

17
intern/cycles/kernel/geom/geom_motion_triangle_intersect.h
View File

@ -44,7 +44,7 @@ ccl_device_inline float3 motion_triangle_refine(const KernelGlobals *kg,
float3 verts[3])
{
#ifdef __INTERSECTION_REFINE__
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
if (UNLIKELY(t == 0.0f)) {
return P;
}
@ -70,7 +70,7 @@ ccl_device_inline float3 motion_triangle_refine(const KernelGlobals *kg,
/* Compute refined position. */
P = P + D * rt;
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}
@ -106,7 +106,7 @@ ccl_device_inline
return motion_triangle_refine(kg, sd, P, D, t, isect_object, isect_prim, verts);
# else
# ifdef __INTERSECTION_REFINE__
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
@ -128,7 +128,7 @@ ccl_device_inline
P = P + D * rt;
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}
@ -186,8 +186,9 @@ ccl_device_inline bool motion_triangle_intersect(const KernelGlobals *kg,
isect->t = t;
isect->u = u;
isect->v = v;
isect->prim = prim_addr;
isect->object = object;
isect->prim = prim;
isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
object;
isect->type = PRIMITIVE_MOTION_TRIANGLE;
return true;
}
@ -288,8 +289,8 @@ ccl_device_inline bool motion_triangle_intersect_local(const KernelGlobals *kg,
isect->t = t;
isect->u = u;
isect->v = v;
isect->prim = prim_addr;
isect->object = object;
isect->prim = prim;
isect->object = local_object;
isect->type = PRIMITIVE_MOTION_TRIANGLE;
/* Record geometric normal. */

9
intern/cycles/kernel/geom/geom_shader_data.h
View File

@ -37,7 +37,7 @@ ccl_device void shader_setup_object_transforms(const KernelGlobals *ccl_restrict
#endif
/* TODO: break this up if it helps reduce register pressure to load data from
* global memory as we write it to shaderdata. */
* global memory as we write it to shader-data. */
ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict kg,
ShaderData *ccl_restrict sd,
const Ray *ccl_restrict ray,
@ -52,10 +52,9 @@ ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict k
sd->v = isect->v;
sd->ray_length = isect->t;
sd->type = isect->type;
sd->object = (isect->object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, isect->prim) :
isect->object;
sd->object = isect->object;
sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
sd->prim = isect->prim;
sd->lamp = LAMP_NONE;
sd->flag = 0;
@ -103,7 +102,7 @@ ccl_device_inline void shader_setup_from_ray(const KernelGlobals *ccl_restrict k
sd->flag |= kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
if (isect->object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
/* instance transform */
object_normal_transform_auto(kg, sd, &sd->N);
object_normal_transform_auto(kg, sd, &sd->Ng);

30
intern/cycles/kernel/geom/geom_triangle.h
View File

@ -29,9 +29,9 @@ ccl_device_inline float3 triangle_normal(const KernelGlobals *kg, ShaderData *sd
{
/* load triangle vertices */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
const float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
const float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
const float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
const float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
const float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
const float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
/* return normal */
if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
@ -54,9 +54,9 @@ ccl_device_inline void triangle_point_normal(const KernelGlobals *kg,
{
/* load triangle vertices */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
float3 v0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
float3 v1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
float3 v2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
/* compute point */
float t = 1.0f - u - v;
*P = (u * v0 + v * v1 + t * v2);
@ -78,9 +78,9 @@ ccl_device_inline void triangle_point_normal(const KernelGlobals *kg,
ccl_device_inline void triangle_vertices(const KernelGlobals *kg, int prim, float3 P[3])
{
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
P[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
P[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
P[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
}
/* Triangle vertex locations and vertex normals */
@ -91,9 +91,9 @@ ccl_device_inline void triangle_vertices_and_normals(const KernelGlobals *kg,
float3 N[3])
{
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
P[0] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
P[1] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
P[2] = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
N[0] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
N[1] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
N[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
@ -145,9 +145,9 @@ ccl_device_inline void triangle_dPdudv(const KernelGlobals *kg,
{
/* fetch triangle vertex coordinates */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
const float3 p0 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 0));
const float3 p1 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 1));
const float3 p2 = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex.w + 2));
const float3 p0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
const float3 p1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
const float3 p2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
/* compute derivatives of P w.r.t. uv */
*dPdu = (p0 - p2);

61
intern/cycles/kernel/geom/geom_triangle_intersect.h
View File

@ -35,13 +35,14 @@ ccl_device_inline bool triangle_intersect(const KernelGlobals *kg,
int object,
int prim_addr)
{
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts = (ssef *)&kg->__prim_tri_verts.data[tri_vindex];
const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
#else
const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
#endif
float t, u, v;
if (ray_triangle_intersect(P,
@ -64,8 +65,9 @@ ccl_device_inline bool triangle_intersect(const KernelGlobals *kg,
if (kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
#endif
{
isect->prim = prim_addr;
isect->object = object;
isect->object = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, prim_addr) :
object;
isect->prim = prim;
isect->type = PRIMITIVE_TRIANGLE;
isect->u = u;
isect->v = v;
@ -102,13 +104,14 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
}
}
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
# if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts = (ssef *)&kg->__prim_tri_verts.data[tri_vindex];
const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
# else
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)),
tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)),
tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
# endif
float t, u, v;
if (!ray_triangle_intersect(P,
@ -167,8 +170,8 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
/* Record intersection. */
Intersection *isect = &local_isect->hits[hit];
isect->prim = prim_addr;
isect->object = object;
isect->prim = prim;
isect->object = local_object;
isect->type = PRIMITIVE_TRIANGLE;
isect->u = u;
isect->v = v;
@ -176,9 +179,9 @@ ccl_device_inline bool triangle_intersect_local(const KernelGlobals *kg,
/* Record geometric normal. */
# if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0)),
tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1)),
tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
# endif
local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
@ -206,7 +209,7 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
const int isect_prim)
{
#ifdef __INTERSECTION_REFINE__
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
if (UNLIKELY(t == 0.0f)) {
return P;
}
@ -219,10 +222,10 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
P = P + D * t;
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect_prim);
const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@ -239,7 +242,7 @@ ccl_device_inline float3 triangle_refine(const KernelGlobals *kg,
P = P + D * rt;
}
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}
@ -265,7 +268,7 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
/* t is always in world space with OptiX. */
return triangle_refine(kg, sd, P, D, t, isect_object, isect_prim);
#else
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_inverse_transform(kg, sd);
P = transform_point(&tfm, P);
@ -276,10 +279,10 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
P = P + D * t;
# ifdef __INTERSECTION_REFINE__
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect_prim);
const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2);
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@ -297,7 +300,7 @@ ccl_device_inline float3 triangle_refine_local(const KernelGlobals *kg,
}
# endif /* __INTERSECTION_REFINE__ */
if (isect_object != OBJECT_NONE) {
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
const Transform tfm = object_get_transform(kg, sd);
P = transform_point(&tfm, P);
}

25
intern/cycles/kernel/integrator/integrator_init_from_bake.h
View File

@ -109,9 +109,17 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
}
/* Position and normal on triangle. */
const int object = kernel_data.bake.object_index;
float3 P, Ng;
int shader;
triangle_point_normal(kg, kernel_data.bake.object_index, prim, u, v, &P, &Ng, &shader);
triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
const int object_flag = kernel_tex_fetch(__object_flag, object);
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM);
P = transform_point_auto(&tfm, P);
}
if (kernel_data.film.pass_background != PASS_UNUSED) {
/* Environment baking. */
@ -130,8 +138,13 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
}
else {
/* Surface baking. */
const float3 N = (shader & SHADER_SMOOTH_NORMAL) ? triangle_smooth_normal(kg, Ng, prim, u, v) :
Ng;
float3 N = (shader & SHADER_SMOOTH_NORMAL) ? triangle_smooth_normal(kg, Ng, prim, u, v) : Ng;
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
Transform itfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
N = normalize(transform_direction_transposed(&itfm, N));
Ng = normalize(transform_direction_transposed(&itfm, Ng));
}
/* Setup ray. */
Ray ray ccl_optional_struct_init;
@ -143,6 +156,12 @@ ccl_device bool integrator_init_from_bake(INTEGRATOR_STATE_ARGS,
/* Setup differentials. */
float3 dPdu, dPdv;
triangle_dPdudv(kg, prim, &dPdu, &dPdv);
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM);
dPdu = transform_direction(&tfm, dPdu);
dPdv = transform_direction(&tfm, dPdv);
}
differential3 dP;
dP.dx = dPdu * dudx + dPdv * dvdx;
dP.dy = dPdu * dudy + dPdv * dvdy;

7
intern/cycles/kernel/integrator/integrator_intersect_closest.h
View File

@ -123,7 +123,7 @@ ccl_device_forceinline void integrator_intersect_shader_next_kernel(
#ifdef __SHADOW_CATCHER__
const int object_flags = intersection_get_object_flags(kg, isect);
if (kernel_shadow_catcher_split(INTEGRATOR_STATE_PASS, object_flags)) {
if (kernel_data.film.use_approximate_shadow_catcher && !kernel_data.background.transparent) {
if (kernel_data.film.pass_background != PASS_UNUSED && !kernel_data.background.transparent) {
INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_SHADOW_CATCHER_BACKGROUND;
if (use_raytrace_kernel) {
@ -160,10 +160,7 @@ ccl_device void integrator_intersect_closest(INTEGRATOR_STATE_ARGS)
if (path_state_ao_bounce(INTEGRATOR_STATE_PASS)) {
ray.t = kernel_data.integrator.ao_bounces_distance;
const int last_object = last_isect_object != OBJECT_NONE ?
last_isect_object :
kernel_tex_fetch(__prim_object, last_isect_prim);
const float object_ao_distance = kernel_tex_fetch(__objects, last_object).ao_distance;
const float object_ao_distance = kernel_tex_fetch(__objects, last_isect_object).ao_distance;
if (object_ao_distance != 0.0f) {
ray.t = object_ao_distance;
}

27
intern/cycles/kernel/integrator/integrator_intersect_volume_stack.h
View File

@ -38,10 +38,13 @@ ccl_device void integrator_volume_stack_update_for_subsurface(INTEGRATOR_STATE_A
volume_ray.P = from_P;
volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t);
/* Store to avoid global fetches on every intersection step. */
const uint volume_stack_size = kernel_data.volume_stack_size;
#ifdef __VOLUME_RECORD_ALL__
Intersection hits[2 * VOLUME_STACK_SIZE + 1];
Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
uint num_hits = scene_intersect_volume_all(
kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, PATH_RAY_ALL_VISIBILITY);
kg, &volume_ray, hits, 2 * volume_stack_size, PATH_RAY_ALL_VISIBILITY);
if (num_hits > 0) {
Intersection *isect = hits;
@ -55,7 +58,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(INTEGRATOR_STATE_A
#else
Intersection isect;
int step = 0;
while (step < 2 * VOLUME_STACK_SIZE &&
while (step < 2 * volume_stack_size &&
scene_intersect_volume(kg, &volume_ray, &isect, PATH_RAY_ALL_VISIBILITY)) {
shader_setup_from_ray(kg, stack_sd, &volume_ray, &isect);
volume_stack_enter_exit(INTEGRATOR_STATE_PASS, stack_sd);
@ -91,12 +94,15 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
stack_index++;
}
/* Store to avoid global fetches on every intersection step. */
const uint volume_stack_size = kernel_data.volume_stack_size;
#ifdef __VOLUME_RECORD_ALL__
Intersection hits[2 * VOLUME_STACK_SIZE + 1];
Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1];
uint num_hits = scene_intersect_volume_all(
kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, visibility);
kg, &volume_ray, hits, 2 * volume_stack_size, visibility);
if (num_hits > 0) {
int enclosed_volumes[VOLUME_STACK_SIZE];
int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
Intersection *isect = hits;
qsort(hits, num_hits, sizeof(Intersection), intersections_compare);
@ -121,7 +127,7 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
break;
}
}
if (need_add && stack_index < VOLUME_STACK_SIZE - 1) {
if (need_add && stack_index < volume_stack_size - 1) {
const VolumeStack new_entry = {stack_sd->object, stack_sd->shader};
integrator_state_write_volume_stack(INTEGRATOR_STATE_PASS, stack_index, new_entry);
++stack_index;
@ -136,11 +142,12 @@ ccl_device void integrator_intersect_volume_stack(INTEGRATOR_STATE_ARGS)
}
}
#else
int enclosed_volumes[VOLUME_STACK_SIZE];
/* CUDA does not support defintion of a variable size arrays, so use the maximum possible. */
int enclosed_volumes[MAX_VOLUME_STACK_SIZE];
int step = 0;
while (stack_index < VOLUME_STACK_SIZE - 1 && enclosed_index < VOLUME_STACK_SIZE - 1 &&
step < 2 * VOLUME_STACK_SIZE) {
while (stack_index < volume_stack_size - 1 && enclosed_index < volume_stack_size - 1 &&
step < 2 * volume_stack_size) {
Intersection isect;
if (!scene_intersect_volume(kg, &volume_ray, &isect, visibility)) {
break;

3
intern/cycles/kernel/integrator/integrator_shade_background.h
View File

@ -192,7 +192,8 @@ ccl_device void integrator_shade_background(INTEGRATOR_STATE_ARGS,
INTEGRATOR_STATE_WRITE(path, flag) &= ~PATH_RAY_SHADOW_CATCHER_BACKGROUND;
const int isect_prim = INTEGRATOR_STATE(isect, prim);
const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim);
const int isect_type = INTEGRATOR_STATE(isect, type);
const int shader = intersection_get_shader_from_isect_prim(kg, isect_prim, isect_type);
const int shader_flags = kernel_tex_fetch(__shaders, shader).flags;
if ((shader_flags & SD_HAS_RAYTRACE) || (kernel_data.film.pass_ao != PASS_UNUSED)) {

6
intern/cycles/kernel/integrator/integrator_state.h
View File

@ -59,8 +59,6 @@ CCL_NAMESPACE_BEGIN
*
* TODO: these could be made dynamic depending on the features used in the scene. */
#define INTEGRATOR_VOLUME_STACK_SIZE VOLUME_STACK_SIZE
#define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024
#define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4
@ -85,12 +83,14 @@ typedef struct IntegratorStateCPU {
#define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
} \
name[cpu_size];
#define KERNEL_STRUCT_VOLUME_STACK_SIZE MAX_VOLUME_STACK_SIZE
#include "kernel/integrator/integrator_state_template.h"
#undef KERNEL_STRUCT_BEGIN
#undef KERNEL_STRUCT_MEMBER
#undef KERNEL_STRUCT_ARRAY_MEMBER
#undef KERNEL_STRUCT_END
#undef KERNEL_STRUCT_END_ARRAY
#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
} IntegratorStateCPU;
/* Path Queue
@ -114,12 +114,14 @@ typedef struct IntegratorStateGPU {
#define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
} \
name[gpu_size];
#define KERNEL_STRUCT_VOLUME_STACK_SIZE MAX_VOLUME_STACK_SIZE
#include "kernel/integrator/integrator_state_template.h"
#undef KERNEL_STRUCT_BEGIN
#undef KERNEL_STRUCT_MEMBER
#undef KERNEL_STRUCT_ARRAY_MEMBER
#undef KERNEL_STRUCT_END
#undef KERNEL_STRUCT_END_ARRAY
#undef KERNEL_STRUCT_VOLUME_STACK_SIZE
/* Count number of queued kernels. */
IntegratorQueueCounter *queue_counter;

8
intern/cycles/kernel/integrator/integrator_state_template.h
View File

@ -107,7 +107,9 @@ KERNEL_STRUCT_END(subsurface)
KERNEL_STRUCT_BEGIN(volume_stack)
KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, object, KERNEL_FEATURE_VOLUME)
KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
KERNEL_STRUCT_END_ARRAY(volume_stack, INTEGRATOR_VOLUME_STACK_SIZE, INTEGRATOR_VOLUME_STACK_SIZE)
KERNEL_STRUCT_END_ARRAY(volume_stack,
KERNEL_STRUCT_VOLUME_STACK_SIZE,
KERNEL_STRUCT_VOLUME_STACK_SIZE)
/********************************* Shadow Path State **************************/
@ -163,5 +165,5 @@ KERNEL_STRUCT_BEGIN(shadow_volume_stack)
KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, object, KERNEL_FEATURE_VOLUME)
KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
KERNEL_STRUCT_END_ARRAY(shadow_volume_stack,
INTEGRATOR_VOLUME_STACK_SIZE,
INTEGRATOR_VOLUME_STACK_SIZE)
KERNEL_STRUCT_VOLUME_STACK_SIZE,
KERNEL_STRUCT_VOLUME_STACK_SIZE)

5
intern/cycles/kernel/integrator/integrator_state_util.h
View File

@ -155,7 +155,7 @@ ccl_device_forceinline void integrator_state_read_shadow_isect(INTEGRATOR_STATE_
ccl_device_forceinline void integrator_state_copy_volume_stack_to_shadow(INTEGRATOR_STATE_ARGS)
{
if (kernel_data.kernel_features & KERNEL_FEATURE_VOLUME) {
for (int i = 0; i < INTEGRATOR_VOLUME_STACK_SIZE; i++) {
for (int i = 0; i < kernel_data.volume_stack_size; i++) {
INTEGRATOR_STATE_ARRAY_WRITE(shadow_volume_stack, i, object) = INTEGRATOR_STATE_ARRAY(
volume_stack, i, object);
INTEGRATOR_STATE_ARRAY_WRITE(shadow_volume_stack, i, shader) = INTEGRATOR_STATE_ARRAY(
@ -223,6 +223,8 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
while (index < gpu_array_size) \
;
# define KERNEL_STRUCT_VOLUME_STACK_SIZE kernel_data.volume_stack_size
# include "kernel/integrator/integrator_state_template.h"
# undef KERNEL_STRUCT_BEGIN
@ -230,6 +232,7 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
# undef KERNEL_STRUCT_ARRAY_MEMBER
# undef KERNEL_STRUCT_END
# undef KERNEL_STRUCT_END_ARRAY
# undef KERNEL_STRUCT_VOLUME_STACK_SIZE
}
ccl_device_inline void integrator_state_move(const IntegratorState to_state,

2
intern/cycles/kernel/integrator/integrator_subsurface.h
View File

@ -577,7 +577,7 @@ ccl_device_inline bool subsurface_scatter(INTEGRATOR_STATE_ARGS)
# ifdef __VOLUME__
/* Update volume stack if needed. */
if (kernel_data.integrator.use_volumes) {
const int object = intersection_get_object(kg, &ss_isect.hits[0]);
const int object = ss_isect.hits[0].object;
const int object_flag = kernel_tex_fetch(__object_flag, object);
if (object_flag & SD_OBJECT_INTERSECTS_VOLUME) {

2
intern/cycles/kernel/integrator/integrator_volume_stack.h
View File

@ -72,7 +72,7 @@ ccl_device void volume_stack_enter_exit(INTEGRATOR_STATE_ARGS,
}
/* If we exceed the stack limit, ignore. */
if (i >= VOLUME_STACK_SIZE - 1) {
if (i >= kernel_data.volume_stack_size - 1) {
return;
}

8
intern/cycles/kernel/kernel_shader.h
View File

@ -186,8 +186,8 @@ ccl_device_inline float _shader_bsdf_multi_eval(const KernelGlobals *kg,
float sum_sample_weight,
const uint light_shader_flags)
{
/* this is the veach one-sample model with balance heuristic, some pdf
* factors drop out when using balance heuristic weighting */
/* This is the veach one-sample model with balance heuristic,
* some PDF factors drop out when using balance heuristic weighting. */
for (int i = 0; i < sd->num_closure; i++) {
const ShaderClosure *sc = &sd->closure[i];
@ -780,8 +780,8 @@ ccl_device_inline void shader_eval_volume(INTEGRATOR_STATE_CONST_ARGS,
break;
}
/* setup shaderdata from stack. it's mostly setup already in
* shader_setup_from_volume, this switching should be quick */
/* Setup shader-data from stack. it's mostly setup already in
* shader_setup_from_volume, this switching should be quick. */
sd->object = entry.object;
sd->lamp = LAMP_NONE;
sd->shader = entry.shader;

8
intern/cycles/kernel/kernel_textures.h
View File

@ -18,11 +18,9 @@
# define KERNEL_TEX(type, name)
#endif
/* bvh */
/* BVH2, not used for OptiX or Embree. */
KERNEL_TEX(float4, __bvh_nodes)
KERNEL_TEX(float4, __bvh_leaf_nodes)
KERNEL_TEX(float4, __prim_tri_verts)
KERNEL_TEX(uint, __prim_tri_index)
KERNEL_TEX(uint, __prim_type)
KERNEL_TEX(uint, __prim_visibility)
KERNEL_TEX(uint, __prim_index)
@ -46,10 +44,12 @@ KERNEL_TEX(float4, __tri_vnormal)
KERNEL_TEX(uint4, __tri_vindex)
KERNEL_TEX(uint, __tri_patch)
KERNEL_TEX(float2, __tri_patch_uv)
KERNEL_TEX(float4, __tri_verts)
/* curves */
KERNEL_TEX(float4, __curves)
KERNEL_TEX(KernelCurve, __curves)
KERNEL_TEX(float4, __curve_keys)
KERNEL_TEX(KernelCurveSegment, __curve_segments)
/* patches */
KERNEL_TEX(uint, __patches)

30
intern/cycles/kernel/kernel_types.h
View File

@ -61,8 +61,6 @@ CCL_NAMESPACE_BEGIN
#define ID_NONE (0.0f)
#define PASS_UNUSED (~0)
#define VOLUME_STACK_SIZE 4
/* Kernel features */
#define __SOBOL__
#define __DPDU__
@ -190,7 +188,7 @@ enum SamplingPattern {
SAMPLING_NUM_PATTERNS,
};
/* these flags values correspond to raytypes in osl.cpp, so keep them in sync! */
/* These flags values correspond to `raytypes` in `osl.cpp`, so keep them in sync! */
enum PathRayFlag {
/* --------------------------------------------------------------------
@ -360,7 +358,6 @@ typedef enum PassType {
PASS_MATERIAL_ID,
PASS_MOTION,
PASS_MOTION_WEIGHT,
PASS_RENDER_TIME,
PASS_CRYPTOMATTE,
PASS_AOV_COLOR,
PASS_AOV_VALUE,
@ -609,6 +606,12 @@ typedef struct AttributeDescriptor {
# define MAX_CLOSURE __MAX_CLOSURE__
#endif
#ifndef __MAX_VOLUME_STACK_SIZE__
# define MAX_VOLUME_STACK_SIZE 32
#else
# define MAX_VOLUME_STACK_SIZE __MAX_VOLUME_STACK_SIZE__
#endif
#define MAX_VOLUME_CLOSURE 8
/* This struct is the base class for all closures. The common members are
@ -1224,7 +1227,7 @@ typedef struct KernelData {
uint kernel_features;
uint max_closures;
uint max_shaders;
uint pad;
uint volume_stack_size;
KernelCamera cam;
KernelFilm film;
@ -1267,10 +1270,25 @@ typedef struct KernelObject {
float ao_distance;
float pad1, pad2;
uint visibility;
int primitive_type;
} KernelObject;
static_assert_align(KernelObject, 16);
typedef struct KernelCurve {
int shader_id;
int first_key;
int num_keys;
int type;
} KernelCurve;
static_assert_align(KernelCurve, 16);
typedef struct KernelCurveSegment {
int prim;
int type;
} KernelCurveSegment;
static_assert_align(KernelCurveSegment, 8);
typedef struct KernelSpotLight {
float radius;
float invarea;

13
intern/cycles/kernel/osl/osl_services.cpp
View File

@ -110,6 +110,7 @@ ustring OSLRenderServices::u_curve_thickness("geom:curve_thickness");
ustring OSLRenderServices::u_curve_length("geom:curve_length");
ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
ustring OSLRenderServices::u_curve_random("geom:curve_random");
ustring OSLRenderServices::u_normal_map_normal("geom:normal_map_normal");
ustring OSLRenderServices::u_path_ray_length("path:ray_length");
ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
ustring OSLRenderServices::u_path_diffuse_depth("path:diffuse_depth");
@ -985,8 +986,18 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobals *kg,
float3 f = curve_tangent_normal(kg, sd);
return set_attribute_float3(f, type, derivatives, val);
}
else
else if (name == u_normal_map_normal) {
if (sd->type & PRIMITIVE_ALL_TRIANGLE) {
float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v);
return set_attribute_float3(f, type, derivatives, val);
}
else {
return false;
}
}
else {
return false;
}
}
bool OSLRenderServices::get_background_attribute(const KernelGlobals *kg,

1
intern/cycles/kernel/osl/osl_services.h
View File

@ -297,6 +297,7 @@ class OSLRenderServices : public OSL::RendererServices {
static ustring u_curve_length;
static ustring u_curve_tangent_normal;
static ustring u_curve_random;
static ustring u_normal_map_normal;
static ustring u_path_ray_length;
static ustring u_path_ray_depth;
static ustring u_path_diffuse_depth;

2
intern/cycles/kernel/shaders/node_normal_map.osl
View File

@ -45,7 +45,7 @@ shader node_normal_map(normal NormalIn = N,
// get _unnormalized_ interpolated normal and tangent
if (getattribute(attr_name, tangent) && getattribute(attr_sign_name, tangent_sign) &&
(!is_smooth || getattribute("geom:N", ninterp))) {
(!is_smooth || getattribute("geom:normal_map_normal", ninterp))) {
// apply normal map
vector B = tangent_sign * cross(ninterp, tangent);
Normal = normalize(mcolor[0] * tangent + mcolor[1] * B + mcolor[2] * ninterp);

9
intern/cycles/kernel/svm/svm_bevel.h
View File

@ -206,8 +206,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
# ifdef __OBJECT_MOTION__
else if (sd->type & PRIMITIVE_MOTION_TRIANGLE) {
float3 verts[3];
motion_triangle_vertices(
kg, sd->object, kernel_tex_fetch(__prim_index, isect.hits[hit].prim), sd->time, verts);
motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
hit_P = motion_triangle_refine_local(
kg, sd, ray->P, ray->D, ray->t, isect.hits[hit].object, isect.hits[hit].prim, verts);
}
@ -215,9 +214,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
/* Get geometric normal. */
float3 hit_Ng = isect.Ng[hit];
int object = (isect.hits[hit].object == OBJECT_NONE) ?
kernel_tex_fetch(__prim_object, isect.hits[hit].prim) :
isect.hits[hit].object;
int object = isect.hits[hit].object;
int object_flag = kernel_tex_fetch(__object_flag, object);
if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
hit_Ng = -hit_Ng;
@ -225,7 +222,7 @@ ccl_device float3 svm_bevel(INTEGRATOR_STATE_CONST_ARGS,
/* Compute smooth normal. */
float3 N = hit_Ng;
int prim = kernel_tex_fetch(__prim_index, isect.hits[hit].prim);
int prim = isect.hits[hit].prim;
int shader = kernel_tex_fetch(__tri_shader, prim);
if (shader & SHADER_SMOOTH_NORMAL) {

6
intern/cycles/kernel/svm/svm_closure.h
View File

@ -85,6 +85,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
}
float3 N = stack_valid(data_node.x) ? stack_load_float3(stack, data_node.x) : sd->N;
if (!(sd->type & PRIMITIVE_ALL_CURVE)) {
N = ensure_valid_reflection(sd->Ng, sd->I, N);
}
float param1 = (stack_valid(param1_offset)) ? stack_load_float(stack, param1_offset) :
__uint_as_float(node.z);
@ -166,6 +169,9 @@ ccl_device_noinline int svm_node_closure_bsdf(
float3 clearcoat_normal = stack_valid(data_cn_ssr.x) ?
stack_load_float3(stack, data_cn_ssr.x) :
sd->N;
if (!(sd->type & PRIMITIVE_ALL_CURVE)) {
clearcoat_normal = ensure_valid_reflection(sd->Ng, sd->I, clearcoat_normal);
}
float3 subsurface_radius = stack_valid(data_cn_ssr.y) ?
stack_load_float3(stack, data_cn_ssr.y) :
make_float3(1.0f, 1.0f, 1.0f);

66
intern/cycles/kernel/svm/svm_math_util.h
View File

@ -200,43 +200,43 @@ ccl_device float svm_math(NodeMathType type, float a, float b, float c)
}
}
/* Calculate color in range 800..12000 using an approximation
* a/x+bx+c for R and G and ((at + b)t + c)t + d) for B
* Max absolute error for RGB is (0.00095, 0.00077, 0.00057),
* which is enough to get the same 8 bit/channel color.
*/
ccl_static_constant float blackbody_table_r[6][3] = {
{2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f},
{3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f},
{4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f},
{4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f},
{4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f},
{3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f},
};
ccl_static_constant float blackbody_table_g[6][3] = {
{-7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f},
{-1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f},
{-1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f},
{-1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f},
{-1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f},
{-5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f},
};
ccl_static_constant float blackbody_table_b[6][4] = {
{0.0f, 0.0f, 0.0f, 0.0f}, /* zeros should be optimized by compiler */
{0.0f, 0.0f, 0.0f, 0.0f},
{0.0f, 0.0f, 0.0f, 0.0f},
{-2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f},
{-2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f},
{6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f},
};
ccl_device float3 svm_math_blackbody_color(float t)
{
/* TODO(lukas): Reimplement in XYZ. */
/* Calculate color in range 800..12000 using an approximation
* a/x+bx+c for R and G and ((at + b)t + c)t + d) for B
* Max absolute error for RGB is (0.00095, 0.00077, 0.00057),
* which is enough to get the same 8 bit/channel color.
*/
const float blackbody_table_r[6][3] = {
{2.52432244e+03f, -1.06185848e-03f, 3.11067539e+00f},
{3.37763626e+03f, -4.34581697e-04f, 1.64843306e+00f},
{4.10671449e+03f, -8.61949938e-05f, 6.41423749e-01f},
{4.66849800e+03f, 2.85655028e-05f, 1.29075375e-01f},
{4.60124770e+03f, 2.89727618e-05f, 1.48001316e-01f},
{3.78765709e+03f, 9.36026367e-06f, 3.98995841e-01f},
};
const float blackbody_table_g[6][3] = {
{-7.50343014e+02f, 3.15679613e-04f, 4.73464526e-01f},
{-1.00402363e+03f, 1.29189794e-04f, 9.08181524e-01f},
{-1.22075471e+03f, 2.56245413e-05f, 1.20753416e+00f},
{-1.42546105e+03f, -4.01730887e-05f, 1.44002695e+00f},
{-1.18134453e+03f, -2.18913373e-05f, 1.30656109e+00f},
{-5.00279505e+02f, -4.59745390e-06f, 1.09090465e+00f},
};
const float blackbody_table_b[6][4] = {
{0.0f, 0.0f, 0.0f, 0.0f}, /* zeros should be optimized by compiler */
{0.0f, 0.0f, 0.0f, 0.0f},
{0.0f, 0.0f, 0.0f, 0.0f},
{-2.02524603e-11f, 1.79435860e-07f, -2.60561875e-04f, -1.41761141e-02f},
{-2.22463426e-13f, -1.55078698e-08f, 3.81675160e-04f, -7.30646033e-01f},
{6.72595954e-13f, -2.73059993e-08f, 4.24068546e-04f, -7.52204323e-01f},
};
if (t >= 12000.0f) {
return make_float3(0.826270103f, 0.994478524f, 1.56626022f);
}

2
intern/cycles/kernel/svm/svm_tex_coord.h
View File

@ -347,8 +347,6 @@ ccl_device_noinline void svm_node_normal_map(const KernelGlobals *kg,
N = safe_normalize(sd->N + (N - sd->N) * strength);
}
N = ensure_valid_reflection(sd->Ng, sd->I, N);
if (is_zero(N)) {
N = sd->N;
}

70
intern/cycles/kernel/svm/svm_wavelength.h
View File

@ -34,44 +34,44 @@ CCL_NAMESPACE_BEGIN
/* Wavelength to RGB */
// CIE colour matching functions xBar, yBar, and zBar for
// wavelengths from 380 through 780 nanometers, every 5
// nanometers. For a wavelength lambda in this range:
// cie_colour_match[(lambda - 380) / 5][0] = xBar
// cie_colour_match[(lambda - 380) / 5][1] = yBar
// cie_colour_match[(lambda - 380) / 5][2] = zBar
ccl_static_constant float cie_colour_match[81][3] = {
{0.0014f, 0.0000f, 0.0065f}, {0.0022f, 0.0001f, 0.0105f}, {0.0042f, 0.0001f, 0.0201f},
{0.0076f, 0.0002f, 0.0362f}, {0.0143f, 0.0004f, 0.0679f}, {0.0232f, 0.0006f, 0.1102f},
{0.0435f, 0.0012f, 0.2074f}, {0.0776f, 0.0022f, 0.3713f}, {0.1344f, 0.0040f, 0.6456f},
{0.2148f, 0.0073f, 1.0391f}, {0.2839f, 0.0116f, 1.3856f}, {0.3285f, 0.0168f, 1.6230f},
{0.3483f, 0.0230f, 1.7471f}, {0.3481f, 0.0298f, 1.7826f}, {0.3362f, 0.0380f, 1.7721f},
{0.3187f, 0.0480f, 1.7441f}, {0.2908f, 0.0600f, 1.6692f}, {0.2511f, 0.0739f, 1.5281f},
{0.1954f, 0.0910f, 1.2876f}, {0.1421f, 0.1126f, 1.0419f}, {0.0956f, 0.1390f, 0.8130f},
{0.0580f, 0.1693f, 0.6162f}, {0.0320f, 0.2080f, 0.4652f}, {0.0147f, 0.2586f, 0.3533f},
{0.0049f, 0.3230f, 0.2720f}, {0.0024f, 0.4073f, 0.2123f}, {0.0093f, 0.5030f, 0.1582f},
{0.0291f, 0.6082f, 0.1117f}, {0.0633f, 0.7100f, 0.0782f}, {0.1096f, 0.7932f, 0.0573f},
{0.1655f, 0.8620f, 0.0422f}, {0.2257f, 0.9149f, 0.0298f}, {0.2904f, 0.9540f, 0.0203f},
{0.3597f, 0.9803f, 0.0134f}, {0.4334f, 0.9950f, 0.0087f}, {0.5121f, 1.0000f, 0.0057f},
{0.5945f, 0.9950f, 0.0039f}, {0.6784f, 0.9786f, 0.0027f}, {0.7621f, 0.9520f, 0.0021f},
{0.8425f, 0.9154f, 0.0018f}, {0.9163f, 0.8700f, 0.0017f}, {0.9786f, 0.8163f, 0.0014f},
{1.0263f, 0.7570f, 0.0011f}, {1.0567f, 0.6949f, 0.0010f}, {1.0622f, 0.6310f, 0.0008f},
{1.0456f, 0.5668f, 0.0006f}, {1.0026f, 0.5030f, 0.0003f}, {0.9384f, 0.4412f, 0.0002f},
{0.8544f, 0.3810f, 0.0002f}, {0.7514f, 0.3210f, 0.0001f}, {0.6424f, 0.2650f, 0.0000f},
{0.5419f, 0.2170f, 0.0000f}, {0.4479f, 0.1750f, 0.0000f}, {0.3608f, 0.1382f, 0.0000f},
{0.2835f, 0.1070f, 0.0000f}, {0.2187f, 0.0816f, 0.0000f}, {0.1649f, 0.0610f, 0.0000f},
{0.1212f, 0.0446f, 0.0000f}, {0.0874f, 0.0320f, 0.0000f}, {0.0636f, 0.0232f, 0.0000f},
{0.0468f, 0.0170f, 0.0000f}, {0.0329f, 0.0119f, 0.0000f}, {0.0227f, 0.0082f, 0.0000f},
{0.0158f, 0.0057f, 0.0000f}, {0.0114f, 0.0041f, 0.0000f}, {0.0081f, 0.0029f, 0.0000f},
{0.0058f, 0.0021f, 0.0000f}, {0.0041f, 0.0015f, 0.0000f}, {0.0029f, 0.0010f, 0.0000f},
{0.0020f, 0.0007f, 0.0000f}, {0.0014f, 0.0005f, 0.0000f}, {0.0010f, 0.0004f, 0.0000f},
{0.0007f, 0.0002f, 0.0000f}, {0.0005f, 0.0002f, 0.0000f}, {0.0003f, 0.0001f, 0.0000f},
{0.0002f, 0.0001f, 0.0000f}, {0.0002f, 0.0001f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f},
{0.0001f, 0.0000f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f}, {0.0000f, 0.0000f, 0.0000f}};
ccl_device_noinline void svm_node_wavelength(
const KernelGlobals *kg, ShaderData *sd, float *stack, uint wavelength, uint color_out)
{
// CIE colour matching functions xBar, yBar, and zBar for
// wavelengths from 380 through 780 nanometers, every 5
// nanometers. For a wavelength lambda in this range:
// cie_colour_match[(lambda - 380) / 5][0] = xBar
// cie_colour_match[(lambda - 380) / 5][1] = yBar
// cie_colour_match[(lambda - 380) / 5][2] = zBar
const float cie_colour_match[81][3] = {
{0.0014f, 0.0000f, 0.0065f}, {0.0022f, 0.0001f, 0.0105f}, {0.0042f, 0.0001f, 0.0201f},
{0.0076f, 0.0002f, 0.0362f}, {0.0143f, 0.0004f, 0.0679f}, {0.0232f, 0.0006f, 0.1102f},
{0.0435f, 0.0012f, 0.2074f}, {0.0776f, 0.0022f, 0.3713f}, {0.1344f, 0.0040f, 0.6456f},
{0.2148f, 0.0073f, 1.0391f}, {0.2839f, 0.0116f, 1.3856f}, {0.3285f, 0.0168f, 1.6230f},
{0.3483f, 0.0230f, 1.7471f}, {0.3481f, 0.0298f, 1.7826f}, {0.3362f, 0.0380f, 1.7721f},
{0.3187f, 0.0480f, 1.7441f}, {0.2908f, 0.0600f, 1.6692f}, {0.2511f, 0.0739f, 1.5281f},
{0.1954f, 0.0910f, 1.2876f}, {0.1421f, 0.1126f, 1.0419f}, {0.0956f, 0.1390f, 0.8130f},
{0.0580f, 0.1693f, 0.6162f}, {0.0320f, 0.2080f, 0.4652f}, {0.0147f, 0.2586f, 0.3533f},
{0.0049f, 0.3230f, 0.2720f}, {0.0024f, 0.4073f, 0.2123f}, {0.0093f, 0.5030f, 0.1582f},
{0.0291f, 0.6082f, 0.1117f}, {0.0633f, 0.7100f, 0.0782f}, {0.1096f, 0.7932f, 0.0573f},
{0.1655f, 0.8620f, 0.0422f}, {0.2257f, 0.9149f, 0.0298f}, {0.2904f, 0.9540f, 0.0203f},
{0.3597f, 0.9803f, 0.0134f}, {0.4334f, 0.9950f, 0.0087f}, {0.5121f, 1.0000f, 0.0057f},
{0.5945f, 0.9950f, 0.0039f}, {0.6784f, 0.9786f, 0.0027f}, {0.7621f, 0.9520f, 0.0021f},
{0.8425f, 0.9154f, 0.0018f}, {0.9163f, 0.8700f, 0.0017f}, {0.9786f, 0.8163f, 0.0014f},
{1.0263f, 0.7570f, 0.0011f}, {1.0567f, 0.6949f, 0.0010f}, {1.0622f, 0.6310f, 0.0008f},
{1.0456f, 0.5668f, 0.0006f}, {1.0026f, 0.5030f, 0.0003f}, {0.9384f, 0.4412f, 0.0002f},
{0.8544f, 0.3810f, 0.0002f}, {0.7514f, 0.3210f, 0.0001f}, {0.6424f, 0.2650f, 0.0000f},
{0.5419f, 0.2170f, 0.0000f}, {0.4479f, 0.1750f, 0.0000f}, {0.3608f, 0.1382f, 0.0000f},
{0.2835f, 0.1070f, 0.0000f}, {0.2187f, 0.0816f, 0.0000f}, {0.1649f, 0.0610f, 0.0000f},
{0.1212f, 0.0446f, 0.0000f}, {0.0874f, 0.0320f, 0.0000f}, {0.0636f, 0.0232f, 0.0000f},
{0.0468f, 0.0170f, 0.0000f}, {0.0329f, 0.0119f, 0.0000f}, {0.0227f, 0.0082f, 0.0000f},
{0.0158f, 0.0057f, 0.0000f}, {0.0114f, 0.0041f, 0.0000f}, {0.0081f, 0.0029f, 0.0000f},
{0.0058f, 0.0021f, 0.0000f}, {0.0041f, 0.0015f, 0.0000f}, {0.0029f, 0.0010f, 0.0000f},
{0.0020f, 0.0007f, 0.0000f}, {0.0014f, 0.0005f, 0.0000f}, {0.0010f, 0.0004f, 0.0000f},
{0.0007f, 0.0002f, 0.0000f}, {0.0005f, 0.0002f, 0.0000f}, {0.0003f, 0.0001f, 0.0000f},
{0.0002f, 0.0001f, 0.0000f}, {0.0002f, 0.0001f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f},
{0.0001f, 0.0000f, 0.0000f}, {0.0001f, 0.0000f, 0.0000f}, {0.0000f, 0.0000f, 0.0000f}};
float lambda_nm = stack_load_float(stack, wavelength);
float ii = (lambda_nm - 380.0f) * (1.0f / 5.0f); // scaled 0..80
int i = float_to_int(ii);

37
intern/cycles/render/buffers.cpp
View File

@ -97,6 +97,11 @@ NODE_DEFINE(BufferParams)
SOCKET_INT(width, "Width", 0);
SOCKET_INT(height, "Height", 0);
SOCKET_INT(window_x, "Window X", 0);
SOCKET_INT(window_y, "Window Y", 0);
SOCKET_INT(window_width, "Window Width", 0);
SOCKET_INT(window_height, "Window Height", 0);
SOCKET_INT(full_x, "Full X", 0);
SOCKET_INT(full_y, "Full Y", 0);
SOCKET_INT(full_width, "Full Width", 0);
@ -233,13 +238,31 @@ void BufferParams::update_offset_stride()
bool BufferParams::modified(const BufferParams &other) const
{
if (!(width == other.width && height == other.height && full_x == other.full_x &&
full_y == other.full_y && full_width == other.full_width &&
full_height == other.full_height && offset == other.offset && stride == other.stride &&
pass_stride == other.pass_stride && layer == other.layer && view == other.view &&
exposure == other.exposure &&
use_approximate_shadow_catcher == other.use_approximate_shadow_catcher &&
use_transparent_background == other.use_transparent_background)) {
if (width != other.width || height != other.height) {
return true;
}
if (full_x != other.full_x || full_y != other.full_y || full_width != other.full_width ||
full_height != other.full_height) {
return true;
}
if (window_x != other.window_x || window_y != other.window_y ||
window_width != other.window_width || window_height != other.window_height) {
return true;
}
if (offset != other.offset || stride != other.stride || pass_stride != other.pass_stride) {
return true;
}
if (layer != other.layer || view != other.view) {
return false;
}
if (exposure != other.exposure ||
use_approximate_shadow_catcher != other.use_approximate_shadow_catcher ||
use_transparent_background != other.use_transparent_background) {
return true;
}

9
intern/cycles/render/buffers.h
View File

@ -82,6 +82,15 @@ class BufferParams : public Node {
int width = 0;
int height = 0;
/* Windows defines which part of the buffers is visible. The part outside of the window is
* considered an "overscan".
*
* Window X and Y are relative to the position of the buffer in the full buffer. */
int window_x = 0;
int window_y = 0;
int window_width = 0;
int window_height = 0;
/* Offset into and width/height of the full buffer. */
int full_x = 0;
int full_y = 0;

2
intern/cycles/render/film.cpp
View File

@ -326,8 +326,6 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
kfilm->pass_bake_differential = kfilm->pass_stride;
break;
case PASS_RENDER_TIME:
break;
case PASS_CRYPTOMATTE:
kfilm->pass_cryptomatte = have_cryptomatte ?
min(kfilm->pass_cryptomatte, kfilm->pass_stride) :

244
intern/cycles/render/geometry.cpp
View File

@ -46,12 +46,6 @@ CCL_NAMESPACE_BEGIN
/* Geometry */
PackFlags operator|=(PackFlags &pack_flags, uint32_t value)
{
pack_flags = (PackFlags)((uint32_t)pack_flags | value);
return pack_flags;
}
NODE_ABSTRACT_DEFINE(Geometry)
{
NodeType *type = NodeType::add("geometry_base", NULL);
@ -79,7 +73,6 @@ Geometry::Geometry(const NodeType *node_type, const Type type)
bvh = NULL;
attr_map_offset = 0;
optix_prim_offset = 0;
prim_offset = 0;
}
@ -707,9 +700,9 @@ void GeometryManager::update_attribute_element_offset(Geometry *geom,
if (element == ATTR_ELEMENT_CURVE)
offset -= hair->prim_offset;
else if (element == ATTR_ELEMENT_CURVE_KEY)
offset -= hair->curvekey_offset;
offset -= hair->curve_key_offset;
else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION)
offset -= hair->curvekey_offset;
offset -= hair->curve_key_offset;
}
}
else {
@ -972,28 +965,22 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
size_t vert_size = 0;
size_t tri_size = 0;
size_t curve_key_size = 0;
size_t curve_size = 0;
size_t curve_key_size = 0;
size_t curve_segment_size = 0;
size_t patch_size = 0;
size_t face_size = 0;
size_t corner_size = 0;
size_t optix_prim_size = 0;
foreach (Geometry *geom, scene->geometry) {
if (geom->optix_prim_offset != optix_prim_size) {
/* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
const bool has_optix_bvh = bvh_layout == BVH_LAYOUT_OPTIX ||
bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE;
geom->need_update_rebuild |= has_optix_bvh;
geom->need_update_bvh_for_offset = true;
}
bool prim_offset_changed = false;
if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
Mesh *mesh = static_cast<Mesh *>(geom);
prim_offset_changed = (mesh->prim_offset != tri_size);
mesh->vert_offset = vert_size;
mesh->prim_offset = tri_size;
@ -1017,27 +1004,35 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
face_size += mesh->get_num_subd_faces();
corner_size += mesh->subd_face_corners.size();
mesh->optix_prim_offset = optix_prim_size;
optix_prim_size += mesh->num_triangles();
}
else if (geom->is_hair()) {
Hair *hair = static_cast<Hair *>(geom);
hair->curvekey_offset = curve_key_size;
prim_offset_changed = (hair->curve_segment_offset != curve_segment_size);
hair->curve_key_offset = curve_key_size;
hair->curve_segment_offset = curve_segment_size;
hair->prim_offset = curve_size;
curve_key_size += hair->get_curve_keys().size();
curve_size += hair->num_curves();
curve_key_size += hair->get_curve_keys().size();
curve_segment_size += hair->num_segments();
}
hair->optix_prim_offset = optix_prim_size;
optix_prim_size += hair->num_segments();
if (prim_offset_changed) {
/* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
const bool has_optix_bvh = bvh_layout == BVH_LAYOUT_OPTIX ||
bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE;
geom->need_update_rebuild |= has_optix_bvh;
geom->need_update_bvh_for_offset = true;
}
}
}
void GeometryManager::device_update_mesh(
Device *, DeviceScene *dscene, Scene *scene, bool for_displacement, Progress &progress)
void GeometryManager::device_update_mesh(Device *,
DeviceScene *dscene,
Scene *scene,
Progress &progress)
{
/* Count. */
size_t vert_size = 0;
@ -1045,6 +1040,7 @@ void GeometryManager::device_update_mesh(
size_t curve_key_size = 0;
size_t curve_size = 0;
size_t curve_segment_size = 0;
size_t patch_size = 0;
@ -1071,31 +1067,7 @@ void GeometryManager::device_update_mesh(
curve_key_size += hair->get_curve_keys().size();
curve_size += hair->num_curves();
}
}
/* Create mapping from triangle to primitive triangle array. */
vector<uint> tri_prim_index(tri_size);
if (for_displacement) {
/* For displacement kernels we do some trickery to make them believe
* we've got all required data ready. However, that data is different
* from final render kernels since we don't have BVH yet, so can't
* really use same semantic of arrays.
*/
foreach (Geometry *geom, scene->geometry) {
if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
Mesh *mesh = static_cast<Mesh *>(geom);
for (size_t i = 0; i < mesh->num_triangles(); ++i) {
tri_prim_index[i + mesh->prim_offset] = 3 * (i + mesh->prim_offset);
}
}
}
}
else {
for (size_t i = 0; i < dscene->prim_index.size(); ++i) {
if ((dscene->prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
tri_prim_index[dscene->prim_index[i]] = dscene->prim_tri_index[i];
}
curve_segment_size += hair->num_segments();
}
}
@ -1104,6 +1076,7 @@ void GeometryManager::device_update_mesh(
/* normals */
progress.set_status("Updating Mesh", "Computing normals");
float4 *tri_verts = dscene->tri_verts.alloc(tri_size * 3);
uint *tri_shader = dscene->tri_shader.alloc(tri_size);
float4 *vnormal = dscene->tri_vnormal.alloc(vert_size);
uint4 *tri_vindex = dscene->tri_vindex.alloc(tri_size);
@ -1129,13 +1102,12 @@ void GeometryManager::device_update_mesh(
mesh->pack_normals(&vnormal[mesh->vert_offset]);
}
if (mesh->triangles_is_modified() || mesh->vert_patch_uv_is_modified() || copy_all_data) {
mesh->pack_verts(tri_prim_index,
if (mesh->verts_is_modified() || mesh->triangles_is_modified() ||
mesh->vert_patch_uv_is_modified() || copy_all_data) {
mesh->pack_verts(&tri_verts[mesh->prim_offset * 3],
&tri_vindex[mesh->prim_offset],
&tri_patch[mesh->prim_offset],
&tri_patch_uv[mesh->vert_offset],
mesh->vert_offset,
mesh->prim_offset);
&tri_patch_uv[mesh->vert_offset]);
}
if (progress.get_cancel())
@ -1146,6 +1118,7 @@ void GeometryManager::device_update_mesh(
/* vertex coordinates */
progress.set_status("Updating Mesh", "Copying Mesh to device");
dscene->tri_verts.copy_to_device_if_modified();
dscene->tri_shader.copy_to_device_if_modified();
dscene->tri_vnormal.copy_to_device_if_modified();
dscene->tri_vindex.copy_to_device_if_modified();
@ -1153,13 +1126,16 @@ void GeometryManager::device_update_mesh(
dscene->tri_patch_uv.copy_to_device_if_modified();
}
if (curve_size != 0) {
progress.set_status("Updating Mesh", "Copying Strands to device");
if (curve_segment_size != 0) {
progress.set_status("Updating Mesh", "Copying Curves to device");
float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
float4 *curves = dscene->curves.alloc(curve_size);
KernelCurve *curves = dscene->curves.alloc(curve_size);
KernelCurveSegment *curve_segments = dscene->curve_segments.alloc(curve_segment_size);
const bool copy_all_data = dscene->curve_keys.need_realloc() || dscene->curves.need_realloc();
const bool copy_all_data = dscene->curve_keys.need_realloc() ||
dscene->curves.need_realloc() ||
dscene->curve_segments.need_realloc();
foreach (Geometry *geom, scene->geometry) {
if (geom->is_hair()) {
@ -1175,9 +1151,9 @@ void GeometryManager::device_update_mesh(
}
hair->pack_curves(scene,
&curve_keys[hair->curvekey_offset],
&curve_keys[hair->curve_key_offset],
&curves[hair->prim_offset],
hair->curvekey_offset);
&curve_segments[hair->curve_segment_offset]);
if (progress.get_cancel())
return;
}
@ -1185,6 +1161,7 @@ void GeometryManager::device_update_mesh(
dscene->curve_keys.copy_to_device_if_modified();
dscene->curves.copy_to_device_if_modified();
dscene->curve_segments.copy_to_device_if_modified();
}
if (patch_size != 0 && dscene->patches.need_realloc()) {
@ -1195,10 +1172,7 @@ void GeometryManager::device_update_mesh(
foreach (Geometry *geom, scene->geometry) {
if (geom->is_mesh()) {
Mesh *mesh = static_cast<Mesh *>(geom);
mesh->pack_patches(&patch_data[mesh->patch_offset],
mesh->vert_offset,
mesh->face_offset,
mesh->corner_offset);
mesh->pack_patches(&patch_data[mesh->patch_offset]);
if (mesh->patch_table) {
mesh->patch_table->copy_adjusting_offsets(&patch_data[mesh->patch_table_offset],
@ -1212,23 +1186,6 @@ void GeometryManager::device_update_mesh(
dscene->patches.copy_to_device();
}
if (for_displacement) {
float4 *prim_tri_verts = dscene->prim_tri_verts.alloc(tri_size * 3);
foreach (Geometry *geom, scene->geometry) {
if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
Mesh *mesh = static_cast<Mesh *>(geom);
for (size_t i = 0; i < mesh->num_triangles(); ++i) {
Mesh::Triangle t = mesh->get_triangle(i);
size_t offset = 3 * (i + mesh->prim_offset);
prim_tri_verts[offset + 0] = float3_to_float4(mesh->verts[t.v[0]]);
prim_tri_verts[offset + 1] = float3_to_float4(mesh->verts[t.v[1]]);
prim_tri_verts[offset + 2] = float3_to_float4(mesh->verts[t.v[2]]);
}
}
}
dscene->prim_tri_verts.copy_to_device();
}
}
void GeometryManager::device_update_bvh(Device *device,
@ -1256,16 +1213,6 @@ void GeometryManager::device_update_bvh(Device *device,
const bool can_refit = scene->bvh != nullptr &&
(bparams.bvh_layout == BVHLayout::BVH_LAYOUT_OPTIX);
PackFlags pack_flags = PackFlags::PACK_NONE;
if (scene->bvh == nullptr) {
pack_flags |= PackFlags::PACK_ALL;
}
if (dscene->prim_visibility.is_modified()) {
pack_flags |= PackFlags::PACK_VISIBILITY;
}
BVH *bvh = scene->bvh;
if (!scene->bvh) {
bvh = scene->bvh = BVH::create(bparams, scene->geometry, scene->objects, device);
@ -1284,77 +1231,7 @@ void GeometryManager::device_update_bvh(Device *device,
pack = std::move(static_cast<BVH2 *>(bvh)->pack);
}
else {
progress.set_status("Updating Scene BVH", "Packing BVH primitives");
size_t num_prims = 0;
size_t num_tri_verts = 0;
foreach (Geometry *geom, scene->geometry) {
if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
Mesh *mesh = static_cast<Mesh *>(geom);
num_prims += mesh->num_triangles();
num_tri_verts += 3 * mesh->num_triangles();
}
else if (geom->is_hair()) {
Hair *hair = static_cast<Hair *>(geom);
num_prims += hair->num_segments();
}
}
pack.root_index = -1;
if (pack_flags != PackFlags::PACK_ALL) {
/* if we do not need to recreate the BVH, then only the vertices are updated, so we can
* safely retake the memory */
dscene->prim_tri_verts.give_data(pack.prim_tri_verts);
if ((pack_flags & PackFlags::PACK_VISIBILITY) != 0) {
dscene->prim_visibility.give_data(pack.prim_visibility);
}
}
else {
/* It is not strictly necessary to skip those resizes we if do not have to repack, as the OS
* will not allocate pages if we do not touch them, however it does help catching bugs. */
pack.prim_tri_index.resize(num_prims);
pack.prim_tri_verts.resize(num_tri_verts);
pack.prim_type.resize(num_prims);
pack.prim_index.resize(num_prims);
pack.prim_object.resize(num_prims);
pack.prim_visibility.resize(num_prims);
}
// Merge visibility flags of all objects and find object index for non-instanced geometry
unordered_map<const Geometry *, pair<int, uint>> geometry_to_object_info;
geometry_to_object_info.reserve(scene->geometry.size());
foreach (Object *ob, scene->objects) {
const Geometry *const geom = ob->get_geometry();
pair<int, uint> &info = geometry_to_object_info[geom];
info.second |= ob->visibility_for_tracing();
if (!geom->is_instanced()) {
info.first = ob->get_device_index();
}
}
TaskPool pool;
// Iterate over scene mesh list instead of objects, since 'optix_prim_offset' was calculated
// based on that list, which may be ordered differently from the object list.
foreach (Geometry *geom, scene->geometry) {
/* Make a copy of the pack_flags so the current geometry's flags do not pollute the others'.
*/
PackFlags geom_pack_flags = pack_flags;
if (geom->is_modified()) {
geom_pack_flags |= PackFlags::PACK_VERTICES;
}
if (geom_pack_flags == PACK_NONE) {
continue;
}
const pair<int, uint> &info = geometry_to_object_info[geom];
pool.push(function_bind(
&Geometry::pack_primitives, geom, &pack, info.first, info.second, geom_pack_flags));
}
pool.wait_work();
}
/* copy to device */
@ -1375,31 +1252,23 @@ void GeometryManager::device_update_bvh(Device *device,
dscene->object_node.steal_data(pack.object_node);
dscene->object_node.copy_to_device();
}
if (pack.prim_tri_index.size() && (dscene->prim_tri_index.need_realloc() || has_bvh2_layout)) {
dscene->prim_tri_index.steal_data(pack.prim_tri_index);
dscene->prim_tri_index.copy_to_device();
}
if (pack.prim_tri_verts.size()) {
dscene->prim_tri_verts.steal_data(pack.prim_tri_verts);
dscene->prim_tri_verts.copy_to_device();
}
if (pack.prim_type.size() && (dscene->prim_type.need_realloc() || has_bvh2_layout)) {
if (pack.prim_type.size()) {
dscene->prim_type.steal_data(pack.prim_type);
dscene->prim_type.copy_to_device();
}
if (pack.prim_visibility.size() && (dscene->prim_visibility.is_modified() || has_bvh2_layout)) {
if (pack.prim_visibility.size()) {
dscene->prim_visibility.steal_data(pack.prim_visibility);
dscene->prim_visibility.copy_to_device();
}
if (pack.prim_index.size() && (dscene->prim_index.need_realloc() || has_bvh2_layout)) {
if (pack.prim_index.size()) {
dscene->prim_index.steal_data(pack.prim_index);
dscene->prim_index.copy_to_device();
}
if (pack.prim_object.size() && (dscene->prim_object.need_realloc() || has_bvh2_layout)) {
if (pack.prim_object.size()) {
dscene->prim_object.steal_data(pack.prim_object);
dscene->prim_object.copy_to_device();
}
if (pack.prim_time.size() && (dscene->prim_time.need_realloc() || has_bvh2_layout)) {
if (pack.prim_time.size()) {
dscene->prim_time.steal_data(pack.prim_time);
dscene->prim_time.copy_to_device();
}
@ -1629,8 +1498,6 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
dscene->bvh_nodes.tag_realloc();
dscene->bvh_leaf_nodes.tag_realloc();
dscene->object_node.tag_realloc();
dscene->prim_tri_verts.tag_realloc();
dscene->prim_tri_index.tag_realloc();
dscene->prim_type.tag_realloc();
dscene->prim_visibility.tag_realloc();
dscene->prim_index.tag_realloc();
@ -1649,6 +1516,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
if (device_update_flags & DEVICE_CURVE_DATA_NEEDS_REALLOC) {
dscene->curves.tag_realloc();
dscene->curve_keys.tag_realloc();
dscene->curve_segments.tag_realloc();
}
}
@ -1691,6 +1559,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
if (device_update_flags & DEVICE_MESH_DATA_MODIFIED) {
/* if anything else than vertices or shaders are modified, we would need to reallocate, so
* these are the only arrays that can be updated */
dscene->tri_verts.tag_modified();
dscene->tri_vnormal.tag_modified();
dscene->tri_shader.tag_modified();
}
@ -1698,6 +1567,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
if (device_update_flags & DEVICE_CURVE_DATA_MODIFIED) {
dscene->curve_keys.tag_modified();
dscene->curves.tag_modified();
dscene->curve_segments.tag_modified();
}
need_flags_update = false;
@ -1906,7 +1776,7 @@ void GeometryManager::device_update(Device *device,
{"device_update (displacement: copy meshes to device)", time});
}
});
device_update_mesh(device, dscene, scene, true, progress);
device_update_mesh(device, dscene, scene, progress);
}
if (progress.get_cancel()) {
return;
@ -2058,7 +1928,7 @@ void GeometryManager::device_update(Device *device,
{"device_update (copy meshes to device)", time});
}
});
device_update_mesh(device, dscene, scene, false, progress);
device_update_mesh(device, dscene, scene, progress);
if (progress.get_cancel()) {
return;
}
@ -2091,13 +1961,12 @@ void GeometryManager::device_update(Device *device,
dscene->bvh_nodes.clear_modified();
dscene->bvh_leaf_nodes.clear_modified();
dscene->object_node.clear_modified();
dscene->prim_tri_verts.clear_modified();
dscene->prim_tri_index.clear_modified();
dscene->prim_type.clear_modified();
dscene->prim_visibility.clear_modified();
dscene->prim_index.clear_modified();
dscene->prim_object.clear_modified();
dscene->prim_time.clear_modified();
dscene->tri_verts.clear_modified();
dscene->tri_shader.clear_modified();
dscene->tri_vindex.clear_modified();
dscene->tri_patch.clear_modified();
@ -2105,6 +1974,7 @@ void GeometryManager::device_update(Device *device,
dscene->tri_patch_uv.clear_modified();
dscene->curves.clear_modified();
dscene->curve_keys.clear_modified();
dscene->curve_segments.clear_modified();
dscene->patches.clear_modified();
dscene->attributes_map.clear_modified();
dscene->attributes_float.clear_modified();
@ -2118,13 +1988,12 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
dscene->bvh_nodes.free_if_need_realloc(force_free);
dscene->bvh_leaf_nodes.free_if_need_realloc(force_free);
dscene->object_node.free_if_need_realloc(force_free);
dscene->prim_tri_verts.free_if_need_realloc(force_free);
dscene->prim_tri_index.free_if_need_realloc(force_free);
dscene->prim_type.free_if_need_realloc(force_free);
dscene->prim_visibility.free_if_need_realloc(force_free);
dscene->prim_index.free_if_need_realloc(force_free);
dscene->prim_object.free_if_need_realloc(force_free);
dscene->prim_time.free_if_need_realloc(force_free);
dscene->tri_verts.free_if_need_realloc(force_free);
dscene->tri_shader.free_if_need_realloc(force_free);
dscene->tri_vnormal.free_if_need_realloc(force_free);
dscene->tri_vindex.free_if_need_realloc(force_free);
@ -2132,6 +2001,7 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
dscene->tri_patch_uv.free_if_need_realloc(force_free);
dscene->curves.free_if_need_realloc(force_free);
dscene->curve_keys.free_if_need_realloc(force_free);
dscene->curve_segments.free_if_need_realloc(force_free);
dscene->patches.free_if_need_realloc(force_free);
dscene->attributes_map.free_if_need_realloc(force_free);
dscene->attributes_float.free_if_need_realloc(force_free);

30
intern/cycles/render/geometry.h
View File

@ -43,24 +43,6 @@ class Shader;
class Volume;
struct PackedBVH;
/* Flags used to determine which geometry data need to be packed. */
enum PackFlags : uint32_t {
PACK_NONE = 0u,
/* Pack the geometry information (e.g. triangle or curve keys indices). */
PACK_GEOMETRY = (1u << 0),
/* Pack the vertices, for Meshes and Volumes' bounding meshes. */
PACK_VERTICES = (1u << 1),
/* Pack the visibility flags for each triangle or curve. */
PACK_VISIBILITY = (1u << 2),
PACK_ALL = (PACK_GEOMETRY | PACK_VERTICES | PACK_VISIBILITY),
};
PackFlags operator|=(PackFlags &pack_flags, uint32_t value);
/* Geometry
*
* Base class for geometric types like Mesh and Hair. */
@ -100,7 +82,6 @@ class Geometry : public Node {
BVH *bvh;
size_t attr_map_offset;
size_t prim_offset;
size_t optix_prim_offset;
/* Shader Properties */
bool has_volume; /* Set in the device_update_flags(). */
@ -144,10 +125,7 @@ class Geometry : public Node {
int n,
int total);
virtual void pack_primitives(PackedBVH *pack,
int object,
uint visibility,
PackFlags pack_flags) = 0;
virtual PrimitiveType primitive_type() const = 0;
/* Check whether the geometry should have own BVH built separately. Briefly,
* own BVH is needed for geometry, if:
@ -260,11 +238,7 @@ class GeometryManager {
void device_update_object(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
void device_update_mesh(Device *device,
DeviceScene *dscene,
Scene *scene,
bool for_displacement,
Progress &progress);
void device_update_mesh(Device *device, DeviceScene *dscene, Scene *scene, Progress &progress);
void device_update_attributes(Device *device,
DeviceScene *dscene,

4
intern/cycles/render/graph.cpp
View File

@ -1149,7 +1149,9 @@ int ShaderGraph::get_num_closures()
num_closures += 8;
}
else if (CLOSURE_IS_VOLUME(closure_type)) {
num_closures += VOLUME_STACK_SIZE;
/* TODO(sergey): Verify this is still needed, since we have special minimized volume storage
* for the volume steps. */
num_closures += MAX_VOLUME_STACK_SIZE;
}
else if (closure_type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
num_closures += 4;

70
intern/cycles/render/hair.cpp
View File

@ -295,7 +295,8 @@ NODE_DEFINE(Hair)
Hair::Hair() : Geometry(get_node_type(), Geometry::HAIR)
{
curvekey_offset = 0;
curve_key_offset = 0;
curve_segment_offset = 0;
curve_shape = CURVE_RIBBON;
}
@ -462,8 +463,8 @@ void Hair::apply_transform(const Transform &tfm, const bool apply_to_motion)
void Hair::pack_curves(Scene *scene,
float4 *curve_key_co,
float4 *curve_data,
size_t curvekey_offset)
KernelCurve *curves,
KernelCurveSegment *curve_segments)
{
size_t curve_keys_size = curve_keys.size();
@ -477,7 +478,10 @@ void Hair::pack_curves(Scene *scene,
}
/* pack curve segments */
const PrimitiveType type = primitive_type();
size_t curve_num = num_curves();
size_t index = 0;
for (size_t i = 0; i < curve_num; i++) {
Curve curve = get_curve(i);
@ -487,56 +491,24 @@ void Hair::pack_curves(Scene *scene,
scene->default_surface;
shader_id = scene->shader_manager->get_shader_id(shader, false);
curve_data[i] = make_float4(__int_as_float(curve.first_key + curvekey_offset),
__int_as_float(curve.num_keys),
__int_as_float(shader_id),
0.0f);
curves[i].shader_id = shader_id;
curves[i].first_key = curve_key_offset + curve.first_key;
curves[i].num_keys = curve.num_keys;
curves[i].type = type;
for (int k = 0; k < curve.num_segments(); ++k, ++index) {
curve_segments[index].prim = prim_offset + i;
curve_segments[index].type = PRIMITIVE_PACK_SEGMENT(type, k);
}
}
}
void Hair::pack_primitives(PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
PrimitiveType Hair::primitive_type() const
{
if (curve_first_key.empty())
return;
/* Separate loop as other arrays are not initialized if their packing is not required. */
if ((pack_flags & PACK_VISIBILITY) != 0) {
unsigned int *prim_visibility = &pack->prim_visibility[optix_prim_offset];
size_t index = 0;
for (size_t j = 0; j < num_curves(); ++j) {
Curve curve = get_curve(j);
for (size_t k = 0; k < curve.num_segments(); ++k, ++index) {
prim_visibility[index] = visibility;
}
}
}
if ((pack_flags & PACK_GEOMETRY) != 0) {
unsigned int *prim_tri_index = &pack->prim_tri_index[optix_prim_offset];
int *prim_type = &pack->prim_type[optix_prim_offset];
int *prim_index = &pack->prim_index[optix_prim_offset];
int *prim_object = &pack->prim_object[optix_prim_offset];
// 'pack->prim_time' is unused by Embree and OptiX
uint type = has_motion_blur() ?
((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
PRIMITIVE_MOTION_CURVE_THICK) :
((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON :
PRIMITIVE_CURVE_THICK);
size_t index = 0;
for (size_t j = 0; j < num_curves(); ++j) {
Curve curve = get_curve(j);
for (size_t k = 0; k < curve.num_segments(); ++k, ++index) {
prim_tri_index[index] = -1;
prim_type[index] = PRIMITIVE_PACK_SEGMENT(type, k);
// Each curve segment points back to its curve index
prim_index[index] = j + prim_offset;
prim_object[index] = object;
}
}
}
return has_motion_blur() ?
((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON :
PRIMITIVE_MOTION_CURVE_THICK) :
((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK);
}
CCL_NAMESPACE_END

15
intern/cycles/render/hair.h
View File

@ -21,6 +21,8 @@
CCL_NAMESPACE_BEGIN
struct KernelCurveSegment;
class Hair : public Geometry {
public:
NODE_DECLARE
@ -95,7 +97,8 @@ class Hair : public Geometry {
NODE_SOCKET_API_ARRAY(array<int>, curve_shader)
/* BVH */
size_t curvekey_offset;
size_t curve_key_offset;
size_t curve_segment_offset;
CurveShapeType curve_shape;
/* Constructor/Destructor */
@ -144,12 +147,12 @@ class Hair : public Geometry {
void get_uv_tiles(ustring map, unordered_set<int> &tiles) override;
/* BVH */
void pack_curves(Scene *scene, float4 *curve_key_co, float4 *curve_data, size_t curvekey_offset);
void pack_curves(Scene *scene,
float4 *curve_key_co,
KernelCurve *curve,
KernelCurveSegment *curve_segments);
void pack_primitives(PackedBVH *pack,
int object,
uint visibility,
PackFlags pack_flags) override;
PrimitiveType primitive_type() const override;
};
CCL_NAMESPACE_END

69
intern/cycles/render/mesh.cpp
View File

@ -729,12 +729,7 @@ void Mesh::pack_normals(float4 *vnormal)
}
}
void Mesh::pack_verts(const vector<uint> &tri_prim_index,
uint4 *tri_vindex,
uint *tri_patch,
float2 *tri_patch_uv,
size_t vert_offset,
size_t tri_offset)
void Mesh::pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv)
{
size_t verts_size = verts.size();
@ -749,17 +744,19 @@ void Mesh::pack_verts(const vector<uint> &tri_prim_index,
size_t triangles_size = num_triangles();
for (size_t i = 0; i < triangles_size; i++) {
Triangle t = get_triangle(i);
tri_vindex[i] = make_uint4(t.v[0] + vert_offset,
t.v[1] + vert_offset,
t.v[2] + vert_offset,
tri_prim_index[i + tri_offset]);
const Triangle t = get_triangle(i);
tri_vindex[i] = make_uint4(
t.v[0] + vert_offset, t.v[1] + vert_offset, t.v[2] + vert_offset, 3 * (prim_offset + i));
tri_patch[i] = (!get_num_subd_faces()) ? -1 : (triangle_patch[i] * 8 + patch_offset);
tri_verts[i * 3] = float3_to_float4(verts[t.v[0]]);
tri_verts[i * 3 + 1] = float3_to_float4(verts[t.v[1]]);
tri_verts[i * 3 + 2] = float3_to_float4(verts[t.v[2]]);
}
}
void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset)
void Mesh::pack_patches(uint *patch_data)
{
size_t num_faces = get_num_subd_faces();
int ngons = 0;
@ -805,53 +802,9 @@ void Mesh::pack_patches(uint *patch_data, uint vert_offset, uint face_offset, ui
}
}
void Mesh::pack_primitives(ccl::PackedBVH *pack, int object, uint visibility, PackFlags pack_flags)
PrimitiveType Mesh::primitive_type() const
{
if (triangles.empty())
return;
const size_t num_prims = num_triangles();
/* Use prim_offset for indexing as it is computed per geometry type, and prim_tri_verts does not
* contain data for Hair geometries. */
float4 *prim_tri_verts = &pack->prim_tri_verts[prim_offset * 3];
// 'pack->prim_time' is unused by Embree and OptiX
uint type = has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
/* Separate loop as other arrays are not initialized if their packing is not required. */
if ((pack_flags & PackFlags::PACK_VISIBILITY) != 0) {
unsigned int *prim_visibility = &pack->prim_visibility[optix_prim_offset];
for (size_t k = 0; k < num_prims; ++k) {
prim_visibility[k] = visibility;
}
}
if ((pack_flags & PackFlags::PACK_GEOMETRY) != 0) {
/* Use optix_prim_offset for indexing as those arrays also contain data for Hair geometries. */
unsigned int *prim_tri_index = &pack->prim_tri_index[optix_prim_offset];
int *prim_type = &pack->prim_type[optix_prim_offset];
int *prim_index = &pack->prim_index[optix_prim_offset];
int *prim_object = &pack->prim_object[optix_prim_offset];
for (size_t k = 0; k < num_prims; ++k) {
if ((pack_flags & PackFlags::PACK_GEOMETRY) != 0) {
prim_tri_index[k] = (prim_offset + k) * 3;
prim_type[k] = type;
prim_index[k] = prim_offset + k;
prim_object[k] = object;
}
}
}
if ((pack_flags & PackFlags::PACK_VERTICES) != 0) {
for (size_t k = 0; k < num_prims; ++k) {
const Mesh::Triangle t = get_triangle(k);
prim_tri_verts[k * 3] = float3_to_float4(verts[t.v[0]]);
prim_tri_verts[k * 3 + 1] = float3_to_float4(verts[t.v[1]]);
prim_tri_verts[k * 3 + 2] = float3_to_float4(verts[t.v[2]]);
}
}
return has_motion_blur() ? PRIMITIVE_MOTION_TRIANGLE : PRIMITIVE_TRIANGLE;
}
CCL_NAMESPACE_END

14
intern/cycles/render/mesh.h
View File

@ -224,18 +224,10 @@ class Mesh : public Geometry {
void pack_shaders(Scene *scene, uint *shader);
void pack_normals(float4 *vnormal);
void pack_verts(const vector<uint> &tri_prim_index,
uint4 *tri_vindex,
uint *tri_patch,
float2 *tri_patch_uv,
size_t vert_offset,
size_t tri_offset);
void pack_patches(uint *patch_data, uint vert_offset, uint face_offset, uint corner_offset);
void pack_verts(float4 *tri_verts, uint4 *tri_vindex, uint *tri_patch, float2 *tri_patch_uv);
void pack_patches(uint *patch_data);
void pack_primitives(PackedBVH *pack,
int object,
uint visibility,
PackFlags pack_flags) override;
PrimitiveType primitive_type() const override;
void tessellate(DiagSplit *split);

20
intern/cycles/render/object.cpp
View File

@ -60,6 +60,7 @@ struct UpdateObjectTransformState {
/* Packed object arrays. Those will be filled in. */
uint *object_flag;
uint *object_visibility;
KernelObject *objects;
Transform *object_motion_pass;
DecomposedTransform *object_motion;
@ -366,6 +367,22 @@ float Object::compute_volume_step_size() const
return step_size;
}
bool Object::check_is_volume() const
{
if (geometry->geometry_type == Geometry::VOLUME) {
return true;
}
for (Node *node : get_geometry()->get_used_shaders()) {
const Shader *shader = static_cast<const Shader *>(node);
if (shader->has_volume_connected) {
return true;
}
}
return false;
}
int Object::get_device_index() const
{
return index;
@ -512,6 +529,9 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
(1.0f - 0.5f * ob->shadow_terminator_shading_offset);
kobject.shadow_terminator_geometry_offset = ob->shadow_terminator_geometry_offset;
kobject.visibility = ob->visibility_for_tracing();
kobject.primitive_type = geom->primitive_type();
/* Object flag. */
if (ob->use_holdout) {
flag |= SD_OBJECT_HOLDOUT_MASK;

7
intern/cycles/render/object.h
View File

@ -109,6 +109,13 @@ class Object : public Node {
/* Compute step size from attributes, shaders, transforms. */
float compute_volume_step_size() const;
/* Check whether this object requires volume sampling (and hence might require space in the
* volume stack).
*
* Note that this is a naive iteration over sharders, which allows to access information prior
* to `scene_update()`. */
bool check_is_volume() const;
protected:
/* Specifies the position of the object in scene->objects and
* in the device vectors. Gets set in device_update. */

5
intern/cycles/render/pass.cpp
View File

@ -89,7 +89,6 @@ const NodeEnum *Pass::get_type_enum()
pass_type_enum.insert("material_id", PASS_MATERIAL_ID);
pass_type_enum.insert("motion", PASS_MOTION);
pass_type_enum.insert("motion_weight", PASS_MOTION_WEIGHT);
pass_type_enum.insert("render_time", PASS_RENDER_TIME);
pass_type_enum.insert("cryptomatte", PASS_CRYPTOMATTE);
pass_type_enum.insert("aov_color", PASS_AOV_COLOR);
pass_type_enum.insert("aov_value", PASS_AOV_VALUE);
@ -217,10 +216,6 @@ PassInfo Pass::get_info(const PassType type, const bool include_albedo)
pass_info.num_components = 3;
pass_info.use_exposure = false;
break;
case PASS_RENDER_TIME:
/* This pass is handled entirely on the host side. */
pass_info.num_components = 0;
break;
case PASS_DIFFUSE_COLOR:
case PASS_GLOSSY_COLOR:

33
intern/cycles/render/scene.cpp
View File

@ -49,13 +49,12 @@ DeviceScene::DeviceScene(Device *device)
: bvh_nodes(device, "__bvh_nodes", MEM_GLOBAL),
bvh_leaf_nodes(device, "__bvh_leaf_nodes", MEM_GLOBAL),
object_node(device, "__object_node", MEM_GLOBAL),
prim_tri_index(device, "__prim_tri_index", MEM_GLOBAL),
prim_tri_verts(device, "__prim_tri_verts", MEM_GLOBAL),
prim_type(device, "__prim_type", MEM_GLOBAL),
prim_visibility(device, "__prim_visibility", MEM_GLOBAL),
prim_index(device, "__prim_index", MEM_GLOBAL),
prim_object(device, "__prim_object", MEM_GLOBAL),
prim_time(device, "__prim_time", MEM_GLOBAL),
tri_verts(device, "__tri_verts", MEM_GLOBAL),
tri_shader(device, "__tri_shader", MEM_GLOBAL),
tri_vnormal(device, "__tri_vnormal", MEM_GLOBAL),
tri_vindex(device, "__tri_vindex", MEM_GLOBAL),
@ -63,6 +62,7 @@ DeviceScene::DeviceScene(Device *device)
tri_patch_uv(device, "__tri_patch_uv", MEM_GLOBAL),
curves(device, "__curves", MEM_GLOBAL),
curve_keys(device, "__curve_keys", MEM_GLOBAL),
curve_segments(device, "__curve_segments", MEM_GLOBAL),
patches(device, "__patches", MEM_GLOBAL),
objects(device, "__objects", MEM_GLOBAL),
object_motion_pass(device, "__object_motion_pass", MEM_GLOBAL),
@ -527,6 +527,8 @@ void Scene::update_kernel_features()
const uint max_closures = (params.background) ? get_max_closure_count() : MAX_CLOSURE;
dscene.data.max_closures = max_closures;
dscene.data.max_shaders = shaders.size();
dscene.data.volume_stack_size = get_volume_stack_size();
}
bool Scene::update(Progress &progress)
@ -642,6 +644,33 @@ int Scene::get_max_closure_count()
return max_closure_global;
}
int Scene::get_volume_stack_size() const
{
/* Quick non-expensive check. Can over-estimate maximum possible nested level, but does not
* require expensive calculation during pre-processing. */
int num_volume_objects = 0;
for (const Object *object : objects) {
if (object->check_is_volume()) {
++num_volume_objects;
}
if (num_volume_objects == MAX_VOLUME_STACK_SIZE) {
break;
}
}
/* Count background world for the stack. */
const Shader *background_shader = background->get_shader(this);
if (background_shader && background_shader->has_volume_connected) {
++num_volume_objects;
}
/* Space for terminator. */
++num_volume_objects;
return min(num_volume_objects, MAX_VOLUME_STACK_SIZE);
}
bool Scene::has_shadow_catcher()
{
if (shadow_catcher_modified_) {

9
intern/cycles/render/scene.h
View File

@ -74,8 +74,6 @@ class DeviceScene {
device_vector<int4> bvh_nodes;
device_vector<int4> bvh_leaf_nodes;
device_vector<int> object_node;
device_vector<uint> prim_tri_index;
device_vector<float4> prim_tri_verts;
device_vector<int> prim_type;
device_vector<uint> prim_visibility;
device_vector<int> prim_index;
@ -83,14 +81,16 @@ class DeviceScene {
device_vector<float2> prim_time;
/* mesh */
device_vector<float4> tri_verts;
device_vector<uint> tri_shader;
device_vector<float4> tri_vnormal;
device_vector<uint4> tri_vindex;
device_vector<uint> tri_patch;
device_vector<float2> tri_patch_uv;
device_vector<float4> curves;
device_vector<KernelCurve> curves;
device_vector<float4> curve_keys;
device_vector<KernelCurveSegment> curve_segments;
device_vector<uint> patches;
@ -344,6 +344,9 @@ class Scene : public NodeOwner {
/* Get maximum number of closures to be used in kernel. */
int get_max_closure_count();
/* Get size of a volume stack needed to render this scene. */
int get_volume_stack_size() const;
template<typename T> void delete_node_impl(T *node)
{
delete node;

20
intern/cycles/render/session.cpp
View File

@ -157,6 +157,13 @@ void Session::run_main_render_loop()
continue;
}
/* Stop rendering if error happened during scene update or other step of preparing scene
* for render. */
if (device->have_error()) {
progress.set_error(device->error_message());
break;
}
{
/* buffers mutex is locked entirely while rendering each
* sample, and released/reacquired on each iteration to allow
@ -172,10 +179,9 @@ void Session::run_main_render_loop()
/* update status and timing */
update_status_time();
/* Stop rendering if error happened during path tracing. */
if (device->have_error()) {
const string &error_message = device->error_message();
progress.set_error(error_message);
progress.set_cancel(error_message);
progress.set_error(device->error_message());
break;
}
}
@ -280,12 +286,20 @@ RenderWork Session::run_update_for_next_iteration()
BufferParams tile_params = buffer_params_;
const Tile &tile = tile_manager_.get_current_tile();
tile_params.width = tile.width;
tile_params.height = tile.height;
tile_params.window_x = tile.window_x;
tile_params.window_y = tile.window_y;
tile_params.window_width = tile.window_width;
tile_params.window_height = tile.window_height;
tile_params.full_x = tile.x + buffer_params_.full_x;
tile_params.full_y = tile.y + buffer_params_.full_y;
tile_params.full_width = buffer_params_.full_width;
tile_params.full_height = buffer_params_.full_height;
tile_params.update_offset_stride();
path_trace_->reset(buffer_params_, tile_params);

70
intern/cycles/render/tile.cpp
View File

@ -372,8 +372,17 @@ void TileManager::update(const BufferParams &params, const Scene *scene)
configure_image_spec_from_buffer(&write_state_.image_spec, buffer_params_, tile_size_);
const DenoiseParams denoise_params = scene->integrator->get_denoise_params();
const AdaptiveSampling adaptive_sampling = scene->integrator->get_adaptive_sampling();
node_to_image_spec_atttributes(
&write_state_.image_spec, &denoise_params, ATTR_DENOISE_SOCKET_PREFIX);
if (adaptive_sampling.use) {
overscan_ = 4;
}
else {
overscan_ = 0;
}
}
bool TileManager::done()
@ -399,19 +408,29 @@ Tile TileManager::get_tile_for_index(int index) const
/* TODO(sergey): Consider using hilbert spiral, or. maybe, even configurable. Not sure this
* brings a lot of value since this is only applicable to BIG tiles. */
const int tile_y = index / tile_state_.num_tiles_x;
const int tile_x = index - tile_y * tile_state_.num_tiles_x;
const int tile_index_y = index / tile_state_.num_tiles_x;
const int tile_index_x = index - tile_index_y * tile_state_.num_tiles_x;
const int tile_x = tile_index_x * tile_size_.x;
const int tile_y = tile_index_y * tile_size_.y;
Tile tile;
tile.x = tile_x * tile_size_.x;
tile.y = tile_y * tile_size_.y;
tile.width = tile_size_.x;
tile.height = tile_size_.y;
tile.x = tile_x - overscan_;
tile.y = tile_y - overscan_;
tile.width = tile_size_.x + 2 * overscan_;
tile.height = tile_size_.y + 2 * overscan_;
tile.x = max(tile.x, 0);
tile.y = max(tile.y, 0);
tile.width = min(tile.width, buffer_params_.width - tile.x);
tile.height = min(tile.height, buffer_params_.height - tile.y);
tile.window_x = tile_x - tile.x;
tile.window_y = tile_y - tile.y;
tile.window_width = min(tile_size_.x, buffer_params_.width - (tile.x + tile.window_x));
tile.window_height = min(tile_size_.y, buffer_params_.height - (tile.y + tile.window_y));
return tile;
}
@ -483,11 +502,22 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
DCHECK_EQ(tile_buffers.params.pass_stride, buffer_params_.pass_stride);
vector<float> pixel_storage;
const BufferParams &tile_params = tile_buffers.params;
const float *pixels = tile_buffers.buffer.data();
const int tile_x = tile_params.full_x - buffer_params_.full_x;
const int tile_y = tile_params.full_y - buffer_params_.full_y;
const int tile_x = tile_params.full_x - buffer_params_.full_x + tile_params.window_x;
const int tile_y = tile_params.full_y - buffer_params_.full_y + tile_params.window_y;
const int64_t pass_stride = tile_params.pass_stride;
const int64_t tile_row_stride = tile_params.width * pass_stride;
const int64_t xstride = pass_stride * sizeof(float);
const int64_t ystride = xstride * tile_params.width;
const int64_t zstride = ystride * tile_params.height;
const float *pixels = tile_buffers.buffer.data() + tile_params.window_x * pass_stride +
tile_params.window_y * tile_row_stride;
VLOG(3) << "Write tile at " << tile_x << ", " << tile_y;
@ -499,13 +529,16 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
* The only thing we have to ensure is that the tile_x and tile_y are a multiple of the
* image tile size, which happens in compute_render_tile_size. */
if (!write_state_.tile_out->write_tiles(tile_x,
tile_x + tile_params.width,
tile_x + tile_params.window_width,
tile_y,
tile_y + tile_params.height,
tile_y + tile_params.window_height,
0,
1,
TypeDesc::FLOAT,
pixels)) {
pixels,
xstride,
ystride,
zstride)) {
LOG(ERROR) << "Error writing tile " << write_state_.tile_out->geterror();
return false;
}
@ -531,12 +564,15 @@ void TileManager::finish_write_tiles()
++tile_index) {
const Tile tile = get_tile_for_index(tile_index);
VLOG(3) << "Write dummy tile at " << tile.x << ", " << tile.y;
const int tile_x = tile.x + tile.window_x;
const int tile_y = tile.y + tile.window_y;
write_state_.tile_out->write_tiles(tile.x,
tile.x + tile.width,
tile.y,
tile.y + tile.height,
VLOG(3) << "Write dummy tile at " << tile_x << ", " << tile_y;
write_state_.tile_out->write_tiles(tile_x,
tile_x + tile.window_width,
tile_y,
tile_y + tile.window_height,
0,
1,
TypeDesc::FLOAT,

6
intern/cycles/render/tile.h
View File

@ -35,6 +35,9 @@ class Tile {
int x = 0, y = 0;
int width = 0, height = 0;
int window_x = 0, window_y = 0;
int window_width = 0, window_height = 0;
Tile()
{
}
@ -128,6 +131,9 @@ class TileManager {
int2 tile_size_ = make_int2(0, 0);
/* Number of extra pixels around the actual tile to render. */
int overscan_ = 0;
BufferParams buffer_params_;
/* Tile scheduling state. */

3
intern/ghost/intern/GHOST_DisplayManagerSDL.cpp
View File

@ -101,8 +101,7 @@ GHOST_TSuccess GHOST_DisplayManagerSDL::setCurrentDisplaySetting(
uint8_t display, const GHOST_DisplaySetting &setting)
{
/*
* Mode switching code ported from Quake 2 version 3.21 and bzflag version
* 2.4.0:
* Mode switching code ported from Quake 2 version 3.21 and BZFLAG version 2.4.0:
* ftp://ftp.idsoftware.com/idstuff/source/q2source-3.21.zip
* See linux/gl_glx.c:GLimp_SetMode
* http://wiki.bzflag.org/BZFlag_Source

4
intern/ghost/intern/GHOST_SystemPathsUnix.cpp
View File

@ -142,7 +142,7 @@ const char *GHOST_SystemPathsUnix::getUserSpecialDir(GHOST_TUserSpecialDirTypes
}
static string path = "";
/* Pipe stderr to /dev/null to avoid error prints. We will fail gracefully still. */
/* Pipe `stderr` to `/dev/null` to avoid error prints. We will fail gracefully still. */
string command = string("xdg-user-dir ") + type_str + " 2> /dev/null";
FILE *fstream = popen(command.c_str(), "r");
@ -152,7 +152,7 @@ const char *GHOST_SystemPathsUnix::getUserSpecialDir(GHOST_TUserSpecialDirTypes
std::stringstream path_stream;
while (!feof(fstream)) {
char c = fgetc(fstream);
/* xdg-user-dir ends the path with '\n'. */
/* `xdg-user-dir` ends the path with '\n'. */
if (c == '\n') {
break;
}

16
intern/ghost/intern/GHOST_SystemX11.cpp
View File

@ -1044,7 +1044,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
#ifdef USE_NON_LATIN_KB_WORKAROUND
/* XXX: Code below is kinda awfully convoluted... Issues are:
* - In keyboards like latin ones, numbers need a 'Shift' to be accessed but key_sym
* - In keyboards like Latin ones, numbers need a 'Shift' to be accessed but key_sym
* is unmodified (or anyone swapping the keys with `xmodmap`).
* - #XLookupKeysym seems to always use first defined key-map (see T47228), which generates
* key-codes unusable by ghost_key_from_keysym for non-Latin-compatible key-maps.
@ -1131,7 +1131,7 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
}
}
#else
/* In keyboards like latin ones,
/* In keyboards like Latin ones,
* numbers needs a 'Shift' to be accessed but key_sym
* is unmodified (or anyone swapping the keys with xmodmap).
*
@ -1528,13 +1528,13 @@ void GHOST_SystemX11::processEvent(XEvent *xe)
window->GetTabletData().Pressure = axis_value / ((float)xtablet.PressureLevels);
}
/* the (short) cast and the & 0xffff is bizarre and unexplained anywhere,
* but I got garbage data without it. Found it in the xidump.c source --matt
/* NOTE(@broken): the (short) cast and the & 0xffff is bizarre and unexplained anywhere,
* but I got garbage data without it. Found it in the `xidump.c` source.
*
* The '& 0xffff' just truncates the value to its two lowest bytes, this probably means
* some drivers do not properly set the whole int value? Since we convert to float
* afterward, I don't think we need to cast to short here, but do not have a device to
* check this. --mont29
* NOTE(@mont29): The '& 0xffff' just truncates the value to its two lowest bytes,
* this probably means some drivers do not properly set the whole int value?
* Since we convert to float afterward,
* I don't think we need to cast to short here, but do not have a device to check this.
*/
if (AXIS_VALUE_GET(3, axis_value)) {
window->GetTabletData().Xtilt = (short)(axis_value & 0xffff) /

6
intern/ghost/intern/GHOST_WindowX11.cpp
View File

@ -1092,9 +1092,9 @@ GHOST_TSuccess GHOST_WindowX11::setOrder(GHOST_TWindowOrder order)
XWindowAttributes attr;
Atom atom;
/* We use both XRaiseWindow and _NET_ACTIVE_WINDOW, since some
* window managers ignore the former (e.g. kwin from kde) and others
* don't implement the latter (e.g. fluxbox pre 0.9.9) */
/* We use both #XRaiseWindow and #_NET_ACTIVE_WINDOW, since some
* window managers ignore the former (e.g. KWIN from KDE) and others
* don't implement the latter (e.g. FLUXBOX before 0.9.9). */
XRaiseWindow(m_display, m_window);

7
intern/guardedalloc/intern/mallocn_guarded_impl.c
View File

@ -89,7 +89,7 @@ typedef struct localListBase {
void *first, *last;
} localListBase;
/* note: keep this struct aligned (e.g., irix/gcc) - Hos */
/* NOTE(@hos): keep this struct aligned (e.g., IRIX/GCC). */
typedef struct MemHead {
int tag1;
size_t len;
@ -98,9 +98,8 @@ typedef struct MemHead {
const char *nextname;
int tag2;
short pad1;
short alignment; /* if non-zero aligned alloc was used
* and alignment is stored here.
*/
/* if non-zero aligned allocation was used and alignment is stored here. */
short alignment;
#ifdef DEBUG_MEMCOUNTER
int _count;
#endif

4
intern/mikktspace/mikktspace.c
View File

@ -1112,7 +1112,7 @@ static tbool AssignRecur(const int piTriListIn[],
static tbool CompareSubGroups(const SSubGroup *pg1, const SSubGroup *pg2);
static void QuickSort(int *pSortBuffer, int iLeft, int iRight, unsigned int uSeed);
static STSpace EvalTspace(int face_indices[],
static STSpace EvalTspace(const int face_indices[],
const int iFaces,
const int piTriListIn[],
const STriInfo pTriInfos[],
@ -1292,7 +1292,7 @@ static tbool GenerateTSpaces(STSpace psTspace[],
return TTRUE;
}
static STSpace EvalTspace(int face_indices[],
static STSpace EvalTspace(const int face_indices[],
const int iFaces,
const int piTriListIn[],
const STriInfo pTriInfos[],

10
release/scripts/modules/bl_i18n_utils/bl_extract_messages.py
View File

@ -378,7 +378,15 @@ def dump_rna_messages(msgs, reports, settings, verbose=False):
if cls in blacklist_rna_class:
return cls.__name__
cls_id = ""
bl_rna = cls.bl_rna
bl_rna = getattr(cls, "bl_rna", None)
# It seems that py-defined 'wrappers' RNA classes (like `MeshEdge` in `bpy_types.py`) need to be accessed
# once from `bpy.types` before they have a valid `bl_rna` member.
# Weirdly enough, this is only triggered on release builds, debug builds somehow do not have that issue.
if bl_rna is None:
if getattr(bpy.types, cls.__name__, None) is not None:
bl_rna = getattr(cls, "bl_rna", None)
if bl_rna is None:
raise TypeError("Unknown RNA class")
while bl_rna:
cls_id = bl_rna.identifier + "." + cls_id
bl_rna = bl_rna.base

2
release/scripts/modules/bl_i18n_utils/utils.py
View File

@ -219,7 +219,7 @@ def enable_addons(addons=None, support=None, disable=False, check_only=False):
try:
import bpy
except ModuleNotFoundError:
print("Could not import bpy, enable_addons must be run from whithin Blender.")
print("Could not import bpy, enable_addons must be run from within Blender.")
return
if addons is None:

3
release/scripts/modules/bl_i18n_utils/utils_spell_check.py
View File

@ -176,6 +176,7 @@ class SpellChecker:
"precalculate",
"precomputing",
"prefetch",
"prefilter", "prefiltering",
"preload",
"premultiply", "premultiplied",
"prepass",
@ -225,6 +226,7 @@ class SpellChecker:
"subpath",
"subsize",
"substep", "substeps",
"substring",
"targetless",
"textbox", "textboxes",
"tilemode",
@ -731,6 +733,7 @@ class SpellChecker:
"tma",
"ui",
"unix",
"uuid",
"vbo", "vbos",
"vr",
"wxyz",

10
release/scripts/modules/bl_keymap_utils/keymap_hierarchy.py
View File

@ -168,9 +168,13 @@ _km_hierarchy = [
('Node Editor', 'NODE_EDITOR', 'WINDOW', [
('Node Generic', 'NODE_EDITOR', 'WINDOW', []),
]),
('Sequencer', 'SEQUENCE_EDITOR', 'WINDOW', [
('SequencerCommon', 'SEQUENCE_EDITOR', 'WINDOW', []),
('SequencerPreview', 'SEQUENCE_EDITOR', 'WINDOW', []),
('SequencerCommon', 'SEQUENCE_EDITOR', 'WINDOW', [
('Sequencer', 'SEQUENCE_EDITOR', 'WINDOW', [
_km_expand_from_toolsystem('SEQUENCE_EDITOR', 'SEQUENCER'),
]),
('SequencerPreview', 'SEQUENCE_EDITOR', 'WINDOW', [
_km_expand_from_toolsystem('SEQUENCE_EDITOR', 'PREVIEW'),
]),
]),
('File Browser', 'FILE_BROWSER', 'WINDOW', [

215
release/scripts/presets/keyconfig/keymap_data/blender_default.py
View File

@ -70,14 +70,14 @@ class Params:
# (derived from other settings).
#
# This case needs to be checked often,
# Shorthand for: `(params.use_fallback_tool if params.select_mouse ==
# 'RIGHT' else False)`.
# Shorthand for: `(params.use_fallback_tool if params.select_mouse == 'RIGHTMOUSE' else False)`.
"use_fallback_tool_rmb",
# Shorthand for: `('CLICK' if params.use_fallback_tool_rmb else
# params.select_mouse_value)`.
"select_mouse_value_fallback",
# Shorthand for: `('CLICK_DRAG' if params.use_pie_click_drag else
# 'PRESS')`
# Shorthand for: `{"type": params.select_tweak, "value": 'ANY'}`.
"select_tweak_event",
# Shorthand for: `('CLICK_DRAG' if params.use_pie_click_drag else 'PRESS')`
"pie_value",
# Shorthand for: `{"type": params.tool_tweak, "value": 'ANY'}`.
"tool_tweak_event",
@ -199,8 +199,9 @@ class Params:
self.use_file_single_click = use_file_single_click
# Convenience variables:
self.use_fallback_tool_rmb = self.use_fallback_tool if self.select_mouse == 'RIGHT' else False
self.use_fallback_tool_rmb = self.use_fallback_tool if select_mouse == 'RIGHT' else False
self.select_mouse_value_fallback = 'CLICK' if self.use_fallback_tool_rmb else self.select_mouse_value
self.select_tweak_event = {"type": self.select_tweak, "value": 'ANY'}
self.pie_value = 'CLICK_DRAG' if use_pie_click_drag else 'PRESS'
self.tool_tweak_event = {"type": self.tool_tweak, "value": 'ANY'}
self.tool_maybe_tweak_event = {"type": self.tool_maybe_tweak, "value": self.tool_maybe_tweak_value}
@ -2355,9 +2356,7 @@ def km_sequencercommon(params):
("wm.context_toggle_enum", {"type": 'TAB', "value": 'PRESS', "ctrl": True},
{"properties": [("data_path", 'space_data.view_type'), ("value_1", 'SEQUENCER'), ("value_2", 'PREVIEW')]}),
("sequencer.refresh_all", {"type": 'R', "value": 'PRESS', "ctrl": True}, None),
("sequencer.select", {"type": params.select_mouse, "value": 'PRESS'}, None),
("sequencer.select", {"type": params.select_mouse, "value": 'PRESS', "shift": True},
{"properties": [("extend", True)]}),])
])
if params.select_mouse == 'LEFTMOUSE' and not params.legacy:
# Quick switch to select tool, since left select can't easily
@ -2431,17 +2430,13 @@ def km_sequencer(params):
{"type": NUMBERS_1[i], "value": 'PRESS'},
{"properties": [("camera", i + 1)]})
for i in range(10)
)),
("sequencer.select", {"type": params.select_mouse, "value": 'PRESS', "alt": True},
{"properties": [("linked_handle", True)]}),
("sequencer.select", {"type": params.select_mouse, "value": 'PRESS', "shift": True, "alt": True},
{"properties": [("extend", True), ("linked_handle", True)]}),
("sequencer.select",
{"type": params.select_mouse, "value": 'PRESS' if params.legacy else 'CLICK', "ctrl": True},
{"properties": [("side_of_frame", True), ("linked_time", True)]}),
("sequencer.select",
{"type": params.select_mouse, "value": 'PRESS' if params.legacy else 'CLICK', "ctrl": True, "shift": True},
{"properties": [("side_of_frame", True), ("linked_time", True), ("extend", True)]}),
)
),
*_template_sequencer_timeline_select(
type=params.select_mouse,
value=params.select_mouse_value_fallback,
legacy=params.legacy,
),
("sequencer.select_more", {"type": 'NUMPAD_PLUS', "value": 'PRESS', "ctrl": True, "repeat": True}, None),
("sequencer.select_less", {"type": 'NUMPAD_MINUS', "value": 'PRESS', "ctrl": True, "repeat": True}, None),
("sequencer.select_linked_pick", {"type": 'L', "value": 'PRESS'}, None),
@ -2486,7 +2481,14 @@ def km_sequencerpreview(params):
{"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
{"items": items},)
items.extend([("sequencer.view_all_preview", {"type": 'HOME', "value": 'PRESS'}, None),
items.extend([
# Selection.
*_template_sequencer_preview_select(
type=params.select_mouse,
value=params.select_mouse_value_fallback,
legacy=params.legacy,
),
("sequencer.view_all_preview", {"type": 'HOME', "value": 'PRESS'}, None),
("sequencer.view_all_preview", {"type": 'NDOF_BUTTON_FIT', "value": 'PRESS'}, None),
("sequencer.view_ghost_border", {"type": 'O', "value": 'PRESS'}, None),
("sequencer.view_zoom_ratio", {"type": 'NUMPAD_8', "value": 'PRESS', "ctrl": True},
@ -2503,8 +2505,9 @@ def km_sequencerpreview(params):
{"properties": [("ratio", 0.25)]}),
("sequencer.view_zoom_ratio", {"type": 'NUMPAD_8', "value": 'PRESS'},
{"properties": [("ratio", 0.125)]}),
("sequencer.sample", {"type": params.action_mouse, "value": 'PRESS'}, None),
op_tool_optional(("transform.translate", {"type": 'G', "value": 'PRESS'}, None),
("transform.translate", {"type": params.select_tweak, "value": 'ANY'}, None),
op_tool_optional(
("transform.translate", {"type": 'G', "value": 'PRESS'}, None),
(op_tool_cycle, "builtin.move"), params),
op_tool_optional(("transform.rotate", {"type": 'R', "value": 'PRESS'}, None),
(op_tool_cycle, "builtin.rotate"), params),
@ -2515,7 +2518,9 @@ def km_sequencerpreview(params):
("sequencer.strip_transform_clear", {"type": 'S', "alt": True, "value": 'PRESS'},
{"properties": [("property", 'SCALE')]}),
("sequencer.strip_transform_clear", {"type": 'R', "alt": True, "value": 'PRESS'},
{"properties": [("property", 'ROTATION')]}),])
{"properties": [("property", 'ROTATION')]}),
*_template_items_context_menu("SEQUENCER_MT_preview_context_menu", params.context_menu_event),
])
return keymap
@ -3756,7 +3761,7 @@ def km_pose(params):
("pose.push", {"type": 'E', "value": 'PRESS', "ctrl": True}, None),
("pose.relax", {"type": 'E', "value": 'PRESS', "alt": True}, None),
("pose.breakdown", {"type": 'E', "value": 'PRESS', "shift": True}, None),
("pose.blend_to_neighbour", {"type": 'E', "value": 'PRESS', "shift": True, "alt": True}, None),
("pose.blend_to_neighbor", {"type": 'E', "value": 'PRESS', "shift": True, "alt": True}, None),
op_menu("VIEW3D_MT_pose_propagate", {"type": 'P', "value": 'PRESS', "alt": True}),
*((("object.hide_collection",
{"type": NUMBERS_1[i], "value": 'PRESS', "any": True},
@ -4077,6 +4082,62 @@ def _template_uv_select_for_fallback(params, fallback):
legacy=params.legacy,)
return []
def _template_sequencer_generic_select(*, type, value, legacy):
return [(
"sequencer.select",
{"type": type, "value": value, **{m: True for m in mods}},
{"properties": [(c, True) for c in props]},
) for props, mods in (
(("deselect_all",) if not legacy else (), ()),
(("toggle",), ("shift",)),
)]
def _template_sequencer_preview_select(*, type, value, legacy):
return _template_sequencer_generic_select(
type=type, value=value, legacy=legacy,
) + [(
"sequencer.select",
{"type": type, "value": value, **{m: True for m in mods}},
{"properties": [(c, True) for c in props]},
) for props, mods in (
(("center",), ("ctrl",)),
# TODO:
# (("enumerate",), ("alt",)),
(("toggle", "center"), ("shift", "ctrl")),
# (("center", "enumerate"), ("ctrl", "alt")),
# (("toggle", "enumerate"), ("shift", "alt")),
# (("toggle", "center", "enumerate"), ("shift", "ctrl", "alt")),
)]
def _template_sequencer_timeline_select(*, type, value, legacy):
return _template_sequencer_generic_select(
type=type, value=value, legacy=legacy,
) + [(
"sequencer.select",
{"type": type, "value": value, **{m: True for m in mods}},
{"properties": [(c, True) for c in props]},
) for props, mods in (
(("linked_handle",), ("alt",)),
(("linked_handle", "extend"), ("shift", "alt",)),
(("side_of_frame", "linked_time"), ("ctrl",)),
(("side_of_frame", "linked_time", "extend"), ("ctrl", "shift")),
)]
def _template_sequencer_select_for_fallback(params, fallback):
if (not fallback) and params.use_fallback_tool_rmb:
# Needed so we have immediate select+tweak when the default select tool is active.
return _template_sequencer_generic_select(
type=params.select_mouse,
value=params.select_mouse_value,
legacy=params.legacy,
)
return []
def km_image_paint(params):
items = []
keymap = ("Image Paint",
@ -5399,19 +5460,23 @@ def km_image_editor_tool_uv_select(params, *, fallback):
def km_image_editor_tool_uv_select_box(params, *, fallback):
return (_fallback_id("Image Editor Tool: Uv, Select Box", fallback),
{"space_type": 'IMAGE_EDITOR', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple("uv.select_box",
# Don't use `tool_maybe_tweak_event`, see comment for this
# slot.
**({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
*_template_uv_select_for_fallback(params, fallback),]},)
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
"uv.select_box",
# Don't use `tool_maybe_tweak_event`, see comment for this slot.
**(params.select_tweak_event if fallback else params.tool_tweak_event))),
*_template_uv_select_for_fallback(params, fallback),
]},
)
def km_image_editor_tool_uv_select_circle(params, *, fallback):
return (_fallback_id("Image Editor Tool: Uv, Select Circle", fallback),
{"space_type": 'IMAGE_EDITOR', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple("uv.select_circle",
type=params.select_tweak if fallback else params.tool_mouse,
value='ANY' if fallback else 'PRESS',
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
"uv.select_circle",
**(params.select_tweak_event if fallback else {"type": params.tool_mouse, "value": 'PRESS'}),
properties=[("wait_for_input", False)])),
# No selection fallback since this operates on press.
]},)
@ -5421,9 +5486,13 @@ def km_image_editor_tool_uv_select_lasso(params, *, fallback):
return (_fallback_id("Image Editor Tool: Uv, Select Lasso", fallback),
{"space_type": 'IMAGE_EDITOR', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple("uv.select_lasso",
**({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
*_template_uv_select_for_fallback(params, fallback),]},)
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
"uv.select_lasso",
**(params.select_tweak_event if fallback else params.tool_tweak_event))),
*_template_uv_select_for_fallback(params, fallback),
]},
)
def km_image_editor_tool_uv_rip_region(params):
@ -5524,11 +5593,14 @@ def km_3d_view_tool_select(params, *, fallback):
def km_3d_view_tool_select_box(params, *, fallback):
return (_fallback_id("3D View Tool: Select Box", fallback),
{"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions("view3d.select_box",
# Don't use `tool_maybe_tweak_event`, see comment for this
# slot.
**({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
*_template_view3d_select_for_fallback(params, fallback),]},)
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
"view3d.select_box",
# Don't use `tool_maybe_tweak_event`, see comment for this slot.
**(params.select_tweak_event if fallback else params.tool_tweak_event))),
*_template_view3d_select_for_fallback(params, fallback),
]},
)
def km_3d_view_tool_select_circle(params, *, fallback):
@ -5548,9 +5620,13 @@ def km_3d_view_tool_select_circle(params, *, fallback):
def km_3d_view_tool_select_lasso(params, *, fallback):
return (_fallback_id("3D View Tool: Select Lasso", fallback),
{"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions("view3d.select_lasso",
**({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
*_template_view3d_select_for_fallback(params, fallback),]})
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
"view3d.select_lasso",
**(params.select_tweak_event if fallback else params.tool_tweak_event))),
*_template_view3d_select_for_fallback(params, fallback),
]}
)
def km_3d_view_tool_transform(params):
@ -6127,11 +6203,14 @@ def km_3d_view_tool_edit_gpencil_select(params, *, fallback):
def km_3d_view_tool_edit_gpencil_select_box(params, *, fallback):
return (_fallback_id("3D View Tool: Edit Gpencil, Select Box", fallback),
{"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions("gpencil.select_box",
# Don't use `tool_maybe_tweak_event`, see comment for this
# slot.
**({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
*_template_view3d_gpencil_select_for_fallback(params, fallback),]},)
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
"gpencil.select_box",
# Don't use `tool_maybe_tweak_event`, see comment for this slot.
**(params.select_tweak_event if fallback else params.tool_tweak_event))),
*_template_view3d_gpencil_select_for_fallback(params, fallback),
]},
)
def km_3d_view_tool_edit_gpencil_select_circle(params, *, fallback):
@ -6151,9 +6230,13 @@ def km_3d_view_tool_edit_gpencil_select_circle(params, *, fallback):
def km_3d_view_tool_edit_gpencil_select_lasso(params, *, fallback):
return (_fallback_id("3D View Tool: Edit Gpencil, Select Lasso", fallback),
{"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
{"items": [*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions("gpencil.select_lasso",
**({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
*_template_view3d_gpencil_select_for_fallback(params, fallback),]})
{"items": [
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
"gpencil.select_lasso",
**(params.select_tweak_event if fallback else params.tool_tweak_event))),
*_template_view3d_gpencil_select_for_fallback(params, fallback),
]}
)
def km_3d_view_tool_edit_gpencil_extrude(params):
@ -6235,20 +6318,34 @@ def km_3d_view_tool_sculpt_gpencil_select_lasso(params):
def km_sequencer_editor_tool_select(params, *, fallback):
return (# TODO, fall-back tool support.
_fallback_id("Sequencer Tool: Select", fallback),
return (
_fallback_id("Sequencer Tool: Tweak", fallback),
{"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
{"items": [("sequencer.select", {"type": params.select_mouse, "value": 'PRESS'}, None),
*_template_items_change_frame(params),]},)
{"items": [
# TODO: Use 2D cursor for preview region (currently `sequencer.sample`).
*([] if fallback else
_template_items_tool_select(params, "sequencer.select", "sequencer.sample", extend="toggle")
),
*([] if (not params.use_fallback_tool_rmb) else _template_sequencer_generic_select(
type=params.select_mouse, value=params.select_mouse_value, legacy=params.legacy)),
*_template_items_change_frame(params),
]},
)
def km_sequencer_editor_tool_select_box(params, *, fallback):
return (# TODO, fall-back tool support.
return (
_fallback_id("Sequencer Tool: Select Box", fallback),
{"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
{"items": [# Don't use `tool_maybe_tweak_event`, see comment for this slot.
*_template_items_tool_select_actions_simple("sequencer.select_box", **params.tool_tweak_event,
properties=[("tweak", params.select_mouse == 'LEFTMOUSE')],),
{"items": [
# Don't use `tool_maybe_tweak_event`, see comment for this slot.
*([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
"sequencer.select_box",
**(params.select_tweak_event if fallback else params.tool_tweak_event),
properties=[("tweak", params.select_mouse == 'LEFTMOUSE')])),
*_template_sequencer_select_for_fallback(params, fallback),
# RMB select can already set the frame, match the tweak tool.
*(_template_items_change_frame(params)
if params.select_mouse == 'LEFTMOUSE' else []),]},)

2
release/scripts/startup/bl_operators/rigidbody.py
View File

@ -176,7 +176,7 @@ class BakeToKeyframes(Operator):
# NOTE: assume that on first frame, the starting rotation is appropriate
obj.rotation_euler = mat.to_euler(rot_mode, obj.rotation_euler)
bpy.ops.anim.keyframe_insert(type='BUILTIN_KSI_LocRot', confirm_success=False)
bpy.ops.anim.keyframe_insert(type='BUILTIN_KSI_LocRot')
# remove baked objects from simulation
bpy.ops.rigidbody.objects_remove()

89
release/scripts/startup/bl_ui/space_sequencer.py
View File

@ -1021,6 +1021,27 @@ class SEQUENCER_MT_context_menu(Menu):
layout.menu("SEQUENCER_MT_strip_lock_mute")
class SEQUENCER_MT_preview_context_menu(Menu):
bl_label = "Sequencer Preview Context Menu"
def draw(self, context):
layout = self.layout
layout.operator_context = 'INVOKE_REGION_WIN'
props = layout.operator("wm.call_panel", text="Rename...")
props.name = "TOPBAR_PT_name"
props.keep_open = False
# TODO: support in preview.
# layout.operator("sequencer.delete", text="Delete")
strip = context.active_sequence_strip
if strip:
pass
class SequencerButtonsPanel:
bl_space_type = 'SEQUENCE_EDITOR'
bl_region_type = 'UI'
@ -1048,16 +1069,23 @@ class SequencerButtonsPanel_Output:
return cls.has_preview(context)
class SEQUENCER_PT_color_tag_picker(Panel):
bl_label = "Color Tag"
class SequencerColorTagPicker:
bl_space_type = 'SEQUENCE_EDITOR'
bl_region_type = 'UI'
bl_category = "Strip"
bl_options = {'HIDE_HEADER', 'INSTANCED'}
@staticmethod
def has_sequencer(context):
return (context.space_data.view_type in {'SEQUENCER', 'SEQUENCER_PREVIEW'})
@classmethod
def poll(cls, context):
return context.active_sequence_strip is not None
return cls.has_sequencer(context) and context.active_sequence_strip is not None
class SEQUENCER_PT_color_tag_picker(SequencerColorTagPicker, Panel):
bl_label = "Color Tag"
bl_category = "Strip"
bl_options = {'HIDE_HEADER', 'INSTANCED'}
def draw(self, context):
layout = self.layout
@ -1069,13 +1097,9 @@ class SEQUENCER_PT_color_tag_picker(Panel):
row.operator("sequencer.strip_color_tag_set", icon=icon).color = 'COLOR_%02d' % i
class SEQUENCER_MT_color_tag_picker(Menu):
class SEQUENCER_MT_color_tag_picker(SequencerColorTagPicker, Menu):
bl_label = "Set Color Tag"
@classmethod
def poll(cls, context):
return context.active_sequence_strip is not None
def draw(self, context):
layout = self.layout
@ -1755,7 +1779,6 @@ class SEQUENCER_PT_adjust_sound(SequencerButtonsPanel, Panel):
def draw(self, context):
layout = self.layout
layout.use_property_split = False
st = context.space_data
overlay_settings = st.timeline_overlay
@ -1765,16 +1788,7 @@ class SEQUENCER_PT_adjust_sound(SequencerButtonsPanel, Panel):
layout.active = not strip.mute
if sound is not None:
col = layout.column()
split = col.split(factor=0.4)
split.label(text="")
split.prop(sound, "use_mono")
if overlay_settings.waveform_display_type == 'DEFAULT_WAVEFORMS':
split = col.split(factor=0.4)
split.label(text="")
split.prop(strip, "show_waveform")
layout.use_property_split = True
col = layout.column()
split = col.split(factor=0.4)
@ -1787,15 +1801,37 @@ class SEQUENCER_PT_adjust_sound(SequencerButtonsPanel, Panel):
split.label(text="Pitch")
split.prop(strip, "pitch", text="")
audio_channels = context.scene.render.ffmpeg.audio_channels
pan_enabled = sound.use_mono and audio_channels != 'MONO'
pan_text = "%.2f°" % (strip.pan * 90)
split = col.split(factor=0.4)
split.alignment = 'RIGHT'
split.label(text="Pan")
audio_channels = context.scene.render.ffmpeg.audio_channels
pan_text = ""
split.prop(strip, "pan", text="")
split.enabled = pan_enabled
if audio_channels != 'MONO' and audio_channels != 'STEREO':
pan_text = "%.2f°" % (strip.pan * 90)
split.prop(strip, "pan", text=pan_text)
split.enabled = sound.use_mono and audio_channels != 'MONO'
split = col.split(factor=0.4)
split.alignment = 'RIGHT'
split.label(text="Pan Angle")
split.enabled = pan_enabled
subsplit = split.row()
subsplit.alignment = 'CENTER'
subsplit.label(text=pan_text)
subsplit.label(text=" ") # Compensate for no decorate.
subsplit.enabled = pan_enabled
layout.use_property_split = False
col = layout.column()
split = col.split(factor=0.4)
split.label(text="")
split.prop(sound, "use_mono")
if overlay_settings.waveform_display_type == 'DEFAULT_WAVEFORMS':
split = col.split(factor=0.4)
split.label(text="")
split.prop(strip, "show_waveform")
@ -2427,6 +2463,7 @@ classes = (
SEQUENCER_MT_strip_lock_mute,
SEQUENCER_MT_color_tag_picker,
SEQUENCER_MT_context_menu,
SEQUENCER_MT_preview_context_menu,
SEQUENCER_PT_color_tag_picker,

8
release/scripts/startup/bl_ui/space_toolsystem_toolbar.py
View File

@ -2232,11 +2232,13 @@ class _defs_sequencer_generic:
class _defs_sequencer_select:
@ToolDef.from_fn
def select():
return dict(idname="builtin.select",
label="Select",
return dict(
idname="builtin.select",
label="Tweak",
icon="ops.generic.select",
widget=None,
keymap="Sequencer Tool: Select",)
keymap="Sequencer Tool: Tweak",
)
@ToolDef.from_fn
def box():

4
release/scripts/startup/bl_ui/space_view3d.py
View File

@ -3469,7 +3469,7 @@ class VIEW3D_MT_pose_slide(Menu):
layout.operator("pose.push")
layout.operator("pose.relax")
layout.operator("pose.breakdown")
layout.operator("pose.blend_to_neighbour")
layout.operator("pose.blend_to_neighbor")
class VIEW3D_MT_pose_propagate(Menu):
@ -3622,7 +3622,7 @@ class VIEW3D_MT_pose_context_menu(Menu):
layout.operator("pose.push")
layout.operator("pose.relax")
layout.operator("pose.breakdown")
layout.operator("pose.blend_to_neighbour")
layout.operator("pose.blend_to_neighbor")
layout.separator()

3
release/scripts/startup/nodeitems_builtins.py
View File

@ -532,6 +532,7 @@ geometry_node_categories = [
NodeItem("GeometryNodeCurveSetHandles"),
NodeItem("GeometryNodeInputTangent"),
NodeItem("GeometryNodeCurveSample"),
NodeItem("GeometryNodeCurveHandleTypeSelection"),
NodeItem("GeometryNodeCurveFillet"),
NodeItem("GeometryNodeCurveReverse"),
]),
@ -624,6 +625,7 @@ geometry_node_categories = [
NodeItem("ShaderNodeClamp"),
NodeItem("ShaderNodeMath"),
NodeItem("FunctionNodeBooleanMath"),
NodeItem("FunctionNodeRotateEuler"),
NodeItem("FunctionNodeFloatCompare"),
NodeItem("FunctionNodeFloatToInt"),
NodeItem("GeometryNodeSwitch"),
@ -645,6 +647,7 @@ geometry_node_categories = [
GeometryNodeCategory("GEO_VOLUME", "Volume", items=[
NodeItem("GeometryNodeLegacyPointsToVolume", poll=geometry_nodes_legacy_poll),
NodeItem("GeometryNodePointsToVolume"),
NodeItem("GeometryNodeVolumeToMesh"),
]),
GeometryNodeCategory("GEO_GROUP", "Group", items=node_group_items),

2
source/blender/blenfont/intern/blf_internal_types.h
View File

@ -152,7 +152,7 @@ typedef struct FontBufInfoBLF {
struct ColorManagedDisplay *display;
/* and the color, the alphas is get from the glyph!
* color is srgb space */
* color is sRGB space */
float col_init[4];
/* cached conversion from 'col_init' */
unsigned char col_char[4];

Some files were not shown because too many files have changed in this diff Show More