Merge branch 'blender-v4.2-release'

This commit is contained in:
Harley Acheson
2024-06-13 10:55:09 -07:00
18 changed files with 180 additions and 135 deletions

View File

@@ -127,9 +127,6 @@ CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler,
cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
cuDevArchitecture = major * 100 + minor * 10;
/* Pop context set by cuCtxCreate. */
cuCtxPopCurrent(NULL);
}
CUDADevice::~CUDADevice()

View File

@@ -326,6 +326,7 @@ set(SRC_KERNEL_SAMPLE_HEADERS
set(SRC_KERNEL_UTIL_HEADERS
util/color.h
util/differential.h
util/ies.h
util/lookup_table.h
util/nanovdb.h
util/profiling.h

View File

@@ -802,7 +802,11 @@ ccl_device_intersect bool kernel_embree_intersect_local(KernelGlobals kg,
float3 P = ray->P;
float3 dir = ray->D;
float3 idir = ray->D;
# ifdef __OBJECT_MOTION__
bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir);
# else
bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
# endif
rtc_ray.org_x = P.x;
rtc_ray.org_y = P.y;

View File

@@ -12,15 +12,17 @@
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_raytrace()
{
const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
const int path_index = (kernel_params.path_index_array) ?
kernel_params.path_index_array[global_index] :
global_index;
integrator_shade_surface_raytrace(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_mnee()
{
const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
global_index;
const int path_index = (kernel_params.path_index_array) ?
kernel_params.path_index_array[global_index] :
global_index;
integrator_shade_surface_mnee(nullptr, path_index, kernel_params.render_buffer);
}

View File

@@ -10,12 +10,12 @@
CCL_NAMESPACE_BEGIN
#ifdef __VOLUME__
ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
IntegratorState state,
const float3 from_P,
const float3 to_P)
{
#ifdef __VOLUME__
PROFILING_INIT(kg, PROFILING_INTERSECT_VOLUME_STACK);
ShaderDataTinyStorage stack_sd_storage;
@@ -217,10 +217,12 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
/* Write terminator. */
const VolumeStack new_entry = {OBJECT_NONE, SHADER_NONE};
integrator_state_write_volume_stack(state, stack_index, new_entry);
#endif
}
ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorState state)
{
#ifdef __VOLUME__
integrator_volume_stack_init(kg, state);
# ifdef __SHADOW_CATCHER__
@@ -238,7 +240,7 @@ ccl_device void integrator_intersect_volume_stack(KernelGlobals kg, IntegratorSt
DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
}
#endif
}
#endif /* __VOLUME__ */
CCL_NAMESPACE_END

View File

@@ -865,9 +865,11 @@ ccl_device_forceinline void integrator_shade_surface_raytrace(
ccl_device_forceinline void integrator_shade_surface_mnee(
KernelGlobals kg, IntegratorState state, ccl_global float *ccl_restrict render_buffer)
{
#ifdef __MNEE__
integrator_shade_surface<(KERNEL_FEATURE_NODE_MASK_SURFACE & ~KERNEL_FEATURE_NODE_RAYTRACE) |
KERNEL_FEATURE_MNEE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE>(kg, state, render_buffer);
#endif
}
CCL_NAMESPACE_END

View File

@@ -442,7 +442,9 @@ ccl_device_inline IntegratorState integrator_state_shadow_catcher_split(KernelGl
to_state->path = state->path;
to_state->ray = state->ray;
to_state->isect = state->isect;
# ifdef __VOLUME__
integrator_state_copy_volume_stack(kg, to_state, state);
# endif
#endif
return to_state;

View File

@@ -6,6 +6,8 @@
CCL_NAMESPACE_BEGIN
#ifdef __SUBSURFACE__
/* BSSRDF using disk based importance sampling.
*
* BSSRDF Importance Sampling, SIGGRAPH 2013
@@ -198,4 +200,6 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
return false;
}
#endif /* __SUBSURFACE__ */
CCL_NAMESPACE_END

View File

@@ -10,6 +10,8 @@
CCL_NAMESPACE_BEGIN
#ifdef __SUBSURFACE__
/* Random walk subsurface scattering.
*
* "Practical and Controllable Subsurface Scattering for Production Path
@@ -160,7 +162,7 @@ ccl_device_forceinline Spectrum subsurface_random_walk_pdf(Spectrum sigma_t,
/* Define the below variable to get the similarity code active,
* and the value represents the cutoff level */
#define SUBSURFACE_RANDOM_WALK_SIMILARITY_LEVEL 9
# define SUBSURFACE_RANDOM_WALK_SIMILARITY_LEVEL 9
ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
IntegratorState state,
@@ -233,20 +235,20 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
/* Our heuristic, a compromise between guiding and classic. */
const float guided_fraction = 1.0f - fmaxf(0.5f, powf(fabsf(anisotropy), 0.125f));
#ifdef SUBSURFACE_RANDOM_WALK_SIMILARITY_LEVEL
# ifdef SUBSURFACE_RANDOM_WALK_SIMILARITY_LEVEL
Spectrum sigma_s_star = sigma_s * (1.0f - anisotropy);
Spectrum sigma_t_star = sigma_t - sigma_s + sigma_s_star;
Spectrum sigma_t_org = sigma_t;
Spectrum sigma_s_org = sigma_s;
const float anisotropy_org = anisotropy;
const float guided_fraction_org = guided_fraction;
#endif
# endif
for (int bounce = 0; bounce < BSSRDF_MAX_BOUNCES; bounce++) {
/* Advance random number offset. */
rng_state.rng_offset += PRNG_BOUNCE_NUM;
#ifdef SUBSURFACE_RANDOM_WALK_SIMILARITY_LEVEL
# ifdef SUBSURFACE_RANDOM_WALK_SIMILARITY_LEVEL
// shadow with local variables according to depth
float anisotropy, guided_fraction;
Spectrum sigma_s, sigma_t;
@@ -262,7 +264,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
sigma_t = sigma_t_star;
sigma_s = sigma_s_star;
}
#endif
# endif
/* Sample color channel, use MIS with balance heuristic. */
float rphase = path_state_rng_1D(kg, &rng_state, PRNG_SUBSURFACE_PHASE_CHANNEL);
@@ -442,4 +444,6 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
return hit;
}
#endif /* __SUBSURFACE__ */
CCL_NAMESPACE_END

View File

@@ -1298,6 +1298,7 @@ bool OSLRenderServices::texture(OSLUStringHash filename,
switch (texture_type) {
case OSLTextureHandle::BEVEL: {
#ifdef __SHADER_RAYTRACE__
/* Bevel shader hack. */
if (nchannels >= 3) {
const IntegratorStateCPU *state = sd->osl_path_state;
@@ -1311,9 +1312,11 @@ bool OSLRenderServices::texture(OSLUStringHash filename,
status = true;
}
}
#endif
break;
}
case OSLTextureHandle::AO: {
#ifdef __SHADER_RAYTRACE__
/* AO shader hack. */
const IntegratorStateCPU *state = sd->osl_path_state;
if (state) {
@@ -1333,6 +1336,7 @@ bool OSLRenderServices::texture(OSLUStringHash filename,
result[0] = svm_ao(kernel_globals, state, sd, N, radius, num_samples, flags);
status = true;
}
#endif
break;
}
case OSLTextureHandle::SVM: {

View File

@@ -1553,7 +1553,7 @@ OSL_NOISE_IMPL(osl_snoise, snoise)
/* Texturing */
#include "kernel/svm/ies.h"
#include "kernel/util/ies.h"
ccl_device_extern ccl_private OSLTextureOptions *osl_get_texture_options(
ccl_private ShaderGlobals *sg)

View File

@@ -126,7 +126,9 @@ ccl_device
float ior = fmaxf(stack_load_float(stack, eta_offset), 1e-5f);
ClosureType distribution = (ClosureType)data_node2.y;
#ifdef __SUBSURFACE__
ClosureType subsurface_method = (ClosureType)data_node2.z;
#endif
float3 valid_reflection_N = maybe_ensure_valid_specular_reflection(sd, N);
float3 coat_normal = stack_valid(coat_normal_offset) ?
@@ -145,7 +147,9 @@ ccl_device
const float3 clamped_base_color = min(base_color, one_float3());
// get the subsurface scattering data
#ifdef __SUBSURFACE__
uint4 data_subsurf = read_node(kg, &offset);
#endif
uint4 data_alpha_emission_thin = read_node(kg, &offset);
svm_unpack_node_uchar4(data_alpha_emission_thin.x,

View File

@@ -4,123 +4,10 @@
#pragma once
#include "kernel/util/ies.h"
CCL_NAMESPACE_BEGIN
/* IES Light */
ccl_device_inline float interpolate_ies_vertical(KernelGlobals kg,
int ofs,
const bool wrap_vlow,
const bool wrap_vhigh,
int v,
int v_num,
float v_frac,
int h)
{
/* Since lookups are performed in spherical coordinates, clamping the coordinates at the low end
* of v (corresponding to the north pole) would result in artifacts. The proper way of dealing
* with this would be to lookup the corresponding value on the other side of the pole, but since
* the horizontal coordinates might be nonuniform, this would require yet another interpolation.
* Therefore, the assumption is made that the light is going to be symmetrical, which means that
* we can just take the corresponding value at the current horizontal coordinate. */
#define IES_LOOKUP(v) kernel_data_fetch(ies, ofs + h * v_num + (v))
float a = 0.0f;
if (v > 0) {
a = IES_LOOKUP(v - 1);
}
else if (wrap_vlow) {
a = IES_LOOKUP(1);
}
float b = IES_LOOKUP(v);
float c = IES_LOOKUP(v + 1);
float d = 0.0f;
if (v + 2 < v_num) {
d = IES_LOOKUP(v + 2);
}
else if (wrap_vhigh) {
d = IES_LOOKUP(v_num - 2);
}
#undef IES_LOOKUP
return cubic_interp(a, b, c, d, v_frac);
}
ccl_device_inline float kernel_ies_interp(KernelGlobals kg, int slot, float h_angle, float v_angle)
{
/* Find offset of the IES data in the table. */
int ofs = __float_as_int(kernel_data_fetch(ies, slot));
if (ofs == -1) {
return 100.0f;
}
int h_num = __float_as_int(kernel_data_fetch(ies, ofs++));
int v_num = __float_as_int(kernel_data_fetch(ies, ofs++));
#define IES_LOOKUP_ANGLE_H(h) kernel_data_fetch(ies, ofs + (h))
#define IES_LOOKUP_ANGLE_V(v) kernel_data_fetch(ies, ofs + h_num + (v))
/* Check whether the angle is within the bounds of the IES texture. */
const float v_low = IES_LOOKUP_ANGLE_V(0), v_high = IES_LOOKUP_ANGLE_V(v_num - 1);
const float h_low = IES_LOOKUP_ANGLE_H(0), h_high = IES_LOOKUP_ANGLE_H(h_num - 1);
if (v_angle < v_low || v_angle >= v_high) {
return 0.0f;
}
if (h_angle < h_low || h_angle >= h_high) {
return 0.0f;
}
/* If the texture covers the full 360° range horizontally, wrap around the lookup
* to get proper cubic interpolation. Otherwise, just set the out-of-range values to zero.
* Similar logic for V, but there we check the lower and upper wrap separately. */
const bool wrap_h = (h_low < 1e-7f && h_high > M_2PI_F - 1e-7f);
const bool wrap_vlow = (v_low < 1e-7f);
const bool wrap_vhigh = (v_high > M_PI_F - 1e-7f);
/* Lookup the angles to find the table position. */
int h_i, v_i;
/* TODO(lukas): Consider using bisection.
* Probably not worth it for the vast majority of IES files. */
for (h_i = 0; IES_LOOKUP_ANGLE_H(h_i + 1) < h_angle; h_i++) {
;
}
for (v_i = 0; IES_LOOKUP_ANGLE_V(v_i + 1) < v_angle; v_i++) {
;
}
float h_frac = inverse_lerp(IES_LOOKUP_ANGLE_H(h_i), IES_LOOKUP_ANGLE_H(h_i + 1), h_angle);
float v_frac = inverse_lerp(IES_LOOKUP_ANGLE_V(v_i), IES_LOOKUP_ANGLE_V(v_i + 1), v_angle);
#undef IES_LOOKUP_ANGLE_H
#undef IES_LOOKUP_ANGLE_V
/* Skip forward to the actual intensity data. */
ofs += h_num + v_num;
float a = 0.0f;
if (h_i > 0) {
a = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i - 1);
}
else if (wrap_h) {
/* The last entry (360°) equals the first one, so we need to wrap around to the one before. */
a = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_num - 2);
}
float b = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i);
float c = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i + 1);
float d = 0.0f;
if (h_i + 2 < h_num) {
d = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i + 2);
}
else if (wrap_h) {
/* Same logic here, wrap around to the second element if necessary. */
d = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, 1);
}
/* Cubic interpolation can result in negative values, so get rid of them. */
return max(cubic_interp(a, b, c, d, h_frac), 0.0f);
}
#ifdef __SVM__
ccl_device_noinline void svm_node_ies(KernelGlobals kg,
ccl_private ShaderData *sd,
ccl_private float *stack,
@@ -142,6 +29,5 @@ ccl_device_noinline void svm_node_ies(KernelGlobals kg,
stack_store_float(stack, fac_offset, fac);
}
}
#endif
CCL_NAMESPACE_END

View File

@@ -180,6 +180,7 @@ CCL_NAMESPACE_BEGIN
#define __SHADOW_LINKING__
#define __LIGHT_TREE__
#define __OBJECT_MOTION__
#define __MNEE__
#define __PASSES__
#define __PATCH_EVAL__
#define __POINTCLOUD__
@@ -211,8 +212,8 @@ CCL_NAMESPACE_BEGIN
/* MNEE caused "Compute function exceeds available temporary registers" in macOS < 13 due to a bug
* in spill buffer allocation sizing. */
#if !defined(__KERNEL_METAL__) || (__KERNEL_METAL_MACOS__ >= 13)
# define __MNEE__
#if defined(__KERNEL_METAL__) && (__KERNEL_METAL_MACOS__ < 13)
# undef __MNEE__
#endif
#if defined(__KERNEL_METAL_AMD__)

View File

@@ -0,0 +1,123 @@
/* SPDX-FileCopyrightText: 2011-2022 Blender Foundation
*
* SPDX-License-Identifier: Apache-2.0 */
#pragma once
CCL_NAMESPACE_BEGIN
/* IES Light */
ccl_device_inline float interpolate_ies_vertical(KernelGlobals kg,
int ofs,
const bool wrap_vlow,
const bool wrap_vhigh,
int v,
int v_num,
float v_frac,
int h)
{
/* Since lookups are performed in spherical coordinates, clamping the coordinates at the low end
* of v (corresponding to the north pole) would result in artifacts. The proper way of dealing
* with this would be to lookup the corresponding value on the other side of the pole, but since
* the horizontal coordinates might be nonuniform, this would require yet another interpolation.
* Therefore, the assumption is made that the light is going to be symmetrical, which means that
* we can just take the corresponding value at the current horizontal coordinate. */
#define IES_LOOKUP(v) kernel_data_fetch(ies, ofs + h * v_num + (v))
float a = 0.0f;
if (v > 0) {
a = IES_LOOKUP(v - 1);
}
else if (wrap_vlow) {
a = IES_LOOKUP(1);
}
float b = IES_LOOKUP(v);
float c = IES_LOOKUP(v + 1);
float d = 0.0f;
if (v + 2 < v_num) {
d = IES_LOOKUP(v + 2);
}
else if (wrap_vhigh) {
d = IES_LOOKUP(v_num - 2);
}
#undef IES_LOOKUP
return cubic_interp(a, b, c, d, v_frac);
}
ccl_device_inline float kernel_ies_interp(KernelGlobals kg, int slot, float h_angle, float v_angle)
{
/* Find offset of the IES data in the table. */
int ofs = __float_as_int(kernel_data_fetch(ies, slot));
if (ofs == -1) {
return 100.0f;
}
int h_num = __float_as_int(kernel_data_fetch(ies, ofs++));
int v_num = __float_as_int(kernel_data_fetch(ies, ofs++));
#define IES_LOOKUP_ANGLE_H(h) kernel_data_fetch(ies, ofs + (h))
#define IES_LOOKUP_ANGLE_V(v) kernel_data_fetch(ies, ofs + h_num + (v))
/* Check whether the angle is within the bounds of the IES texture. */
const float v_low = IES_LOOKUP_ANGLE_V(0), v_high = IES_LOOKUP_ANGLE_V(v_num - 1);
const float h_low = IES_LOOKUP_ANGLE_H(0), h_high = IES_LOOKUP_ANGLE_H(h_num - 1);
if (v_angle < v_low || v_angle >= v_high) {
return 0.0f;
}
if (h_angle < h_low || h_angle >= h_high) {
return 0.0f;
}
/* If the texture covers the full 360° range horizontally, wrap around the lookup
* to get proper cubic interpolation. Otherwise, just set the out-of-range values to zero.
* Similar logic for V, but there we check the lower and upper wrap separately. */
const bool wrap_h = (h_low < 1e-7f && h_high > M_2PI_F - 1e-7f);
const bool wrap_vlow = (v_low < 1e-7f);
const bool wrap_vhigh = (v_high > M_PI_F - 1e-7f);
/* Lookup the angles to find the table position. */
int h_i, v_i;
/* TODO(lukas): Consider using bisection.
* Probably not worth it for the vast majority of IES files. */
for (h_i = 0; IES_LOOKUP_ANGLE_H(h_i + 1) < h_angle; h_i++) {
;
}
for (v_i = 0; IES_LOOKUP_ANGLE_V(v_i + 1) < v_angle; v_i++) {
;
}
float h_frac = inverse_lerp(IES_LOOKUP_ANGLE_H(h_i), IES_LOOKUP_ANGLE_H(h_i + 1), h_angle);
float v_frac = inverse_lerp(IES_LOOKUP_ANGLE_V(v_i), IES_LOOKUP_ANGLE_V(v_i + 1), v_angle);
#undef IES_LOOKUP_ANGLE_H
#undef IES_LOOKUP_ANGLE_V
/* Skip forward to the actual intensity data. */
ofs += h_num + v_num;
float a = 0.0f;
if (h_i > 0) {
a = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i - 1);
}
else if (wrap_h) {
/* The last entry (360°) equals the first one, so we need to wrap around to the one before. */
a = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_num - 2);
}
float b = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i);
float c = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i + 1);
float d = 0.0f;
if (h_i + 2 < h_num) {
d = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, h_i + 2);
}
else if (wrap_h) {
/* Same logic here, wrap around to the second element if necessary. */
d = interpolate_ies_vertical(kg, ofs, wrap_vlow, wrap_vhigh, v_i, v_num, v_frac, 1);
}
/* Cubic interpolation can result in negative values, so get rid of them. */
return max(cubic_interp(a, b, c, d, h_frac), 0.0f);
}
CCL_NAMESPACE_END

View File

@@ -16,7 +16,7 @@ bool is_front_face_hit(float stored_hit_depth)
void main()
{
float hit_depths[VOLUME_HIT_DEPTH_MAX];
float hit_ordered[VOLUME_HIT_DEPTH_MAX];
float hit_ordered[VOLUME_HIT_DEPTH_MAX + 1];
int hit_index[VOLUME_HIT_DEPTH_MAX];
ivec2 texel = ivec2(gl_FragCoord.xy);
@@ -57,9 +57,12 @@ void main()
/* True if last interface was a volume entry. */
/* Initialized to front facing if first hit is a backface to support camera inside the volume. */
bool last_frontfacing = !is_front_face_hit(hit_ordered[0]);
/* Add artificial backfacing hit to close volumes we entered but never exited.
* Fixes issues with non-manifold meshes or things like water planes. */
hit_ordered[hit_count] = -1.0;
/* Bit index of the last interface. */
int last_bit = 0;
for (int i = 0; i < hit_count; i++) {
for (int i = 0; i <= hit_count; i++) {
bool frontfacing = is_front_face_hit(hit_ordered[i]);
if (last_frontfacing == frontfacing) {
/* Same facing, do not treat as a volume interface. */

View File

@@ -433,6 +433,10 @@ static void file_main_region_init(wmWindowManager *wm, ARegion *region)
UI_view2d_region_reinit(&region->v2d, V2D_COMMONVIEW_LIST, region->winx, region->winy);
/* Truncate, otherwise these can be on ".5" and give fuzzy text. #77696. */
region->v2d.cur.ymin = trunc(region->v2d.cur.ymin);
region->v2d.cur.ymax = trunc(region->v2d.cur.ymax);
/* own keymaps */
keymap = WM_keymap_ensure(wm->defaultconf, "File Browser", SPACE_FILE, RGN_TYPE_WINDOW);
WM_event_add_keymap_handler_v2d_mask(&region->handlers, keymap);

View File

@@ -782,6 +782,8 @@ static int view_socket(const bContext &C,
if (viewer_bsocket == nullptr) {
return OPERATOR_CANCELLED;
}
viewer_bsocket->flag &= ~SOCK_HIDDEN;
bNodeLink *viewer_link = nullptr;
LISTBASE_FOREACH_MUTABLE (bNodeLink *, link, &btree.links) {
if (link->tosock == viewer_bsocket) {