Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern/cycles/kernel
diff options
context:
space:
mode:
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r--intern/cycles/kernel/CMakeLists.txt4
-rw-r--r--intern/cycles/kernel/bvh/bvh.h23
-rw-r--r--intern/cycles/kernel/geom/geom_curve_intersect.h9
-rw-r--r--intern/cycles/kernel/kernel_camera.h2
-rw-r--r--intern/cycles/kernel/kernel_light_background.h2
-rw-r--r--intern/cycles/kernel/kernel_path.h14
-rw-r--r--intern/cycles/kernel/kernel_shader.h31
-rw-r--r--intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h2
-rw-r--r--intern/cycles/kernel/kernels/optix/kernel_optix.cu48
-rw-r--r--intern/cycles/kernel/shaders/node_sky_texture.osl50
-rw-r--r--intern/cycles/kernel/svm/svm_sky.h21
11 files changed, 123 insertions, 83 deletions
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 7cc0d32d521..db146226dc7 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -539,7 +539,7 @@ if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES)
${SRC_UTIL_HEADERS}
COMMAND ${CUBIN_CC_ENV}
"$<TARGET_FILE:cycles_cubin_cc>"
- -target 30
+ -target 50
-ptx
-i ${CMAKE_CURRENT_SOURCE_DIR}/${input}
${cuda_flags}
@@ -563,7 +563,7 @@ if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES)
COMMAND
${CUDA_NVCC_EXECUTABLE}
--ptx
- -arch=sm_30
+ -arch=sm_50
${cuda_flags}
${input}
WORKING_DIRECTORY
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index 80b58f46329..3049f243ae9 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -172,11 +172,11 @@ ccl_device_intersect bool scene_intersect(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
- 0xFF,
+ 0xF,
OPTIX_RAY_FLAG_NONE,
+ 0, // SBT offset for PG_HITD
0,
0,
- 0, // SBT offset for PG_HITD
p0,
p1,
p2,
@@ -264,12 +264,13 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
+ // Skip curves
+ 0x3,
// Need to always call into __anyhit__kernel_optix_local_hit
- 0xFF,
OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
- 1,
+ 2, // SBT offset for PG_HITL
+ 0,
0,
- 0, // SBT offset for PG_HITL
p0,
p1,
p2,
@@ -374,12 +375,12 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
+ 0xF,
// Need to always call into __anyhit__kernel_optix_shadow_all_hit
- 0xFF,
OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
- 2,
+ 1, // SBT offset for PG_HITS
+ 0,
0,
- 0, // SBT offset for PG_HITS
p0,
p1,
*num_hits,
@@ -458,12 +459,12 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
- // Visibility mask set to only intersect objects with volumes
- 0x02,
+ // Skip everything but volumes
+ 0x2,
OPTIX_RAY_FLAG_NONE,
+ 0, // SBT offset for PG_HITD
0,
0,
- 0, // SBT offset for PG_HITD
p0,
p1,
p2,
diff --git a/intern/cycles/kernel/geom/geom_curve_intersect.h b/intern/cycles/kernel/geom/geom_curve_intersect.h
index c04dbee52cc..06d2c016f5b 100644
--- a/intern/cycles/kernel/geom/geom_curve_intersect.h
+++ b/intern/cycles/kernel/geom/geom_curve_intersect.h
@@ -734,7 +734,6 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
}
sd->u = isect->u;
- sd->v = isect->v;
P = P + D * t;
@@ -750,6 +749,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
sd->N = normalize(sine * bitangent - cosine * normalize(cross(tangent, bitangent)));
sd->Ng = -D;
+ sd->v = isect->v;
# if 0
/* This approximates the position and geometric normal of a thick curve too,
@@ -764,8 +764,11 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
* This could be optimized by recording the normal in the intersection,
* however for Optix this would go beyond the size of the payload. */
const float3 P_inside = float4_to_float3(catmull_rom_basis_eval(P_curve, isect->u));
- sd->Ng = normalize(P - P_inside);
- sd->N = sd->Ng;
+ const float3 Ng = normalize(P - P_inside);
+
+ sd->N = Ng;
+ sd->Ng = Ng;
+ sd->v = 0.0f;
}
# ifdef __DPDU__
diff --git a/intern/cycles/kernel/kernel_camera.h b/intern/cycles/kernel/kernel_camera.h
index 445cf9eb44b..efe46d5b0dd 100644
--- a/intern/cycles/kernel/kernel_camera.h
+++ b/intern/cycles/kernel/kernel_camera.h
@@ -379,7 +379,7 @@ ccl_device_inline void camera_sample(KernelGlobals *kg,
const int shutter_table_offset = kernel_data.cam.shutter_table_offset;
ray->time = lookup_table_read(kg, time, shutter_table_offset, SHUTTER_TABLE_SIZE);
/* TODO(sergey): Currently single rolling shutter effect type only
- * where scan-lines are acquired from top to bottom and whole scanline
+ * where scan-lines are acquired from top to bottom and whole scan-line
* is acquired at once (no delay in acquisition happens between pixels
* of single scan-line).
*
diff --git a/intern/cycles/kernel/kernel_light_background.h b/intern/cycles/kernel/kernel_light_background.h
index 30e336f0f80..5fa25069fcd 100644
--- a/intern/cycles/kernel/kernel_light_background.h
+++ b/intern/cycles/kernel/kernel_light_background.h
@@ -445,4 +445,4 @@ ccl_device float background_light_pdf(KernelGlobals *kg, float3 P, float3 direct
#endif
-CCL_NAMESPACE_END \ No newline at end of file
+CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index ba46d84d158..c332d5ad3ec 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -284,19 +284,11 @@ ccl_device_forceinline bool kernel_path_shader_apply(KernelGlobals *kg,
#ifdef __HOLDOUT__
if (((sd->flag & SD_HOLDOUT) || (sd->object_flag & SD_OBJECT_HOLDOUT_MASK)) &&
(state->flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
+ const float3 holdout_weight = shader_holdout_apply(kg, sd);
if (kernel_data.background.transparent) {
- float3 holdout_weight;
- if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
- holdout_weight = make_float3(1.0f, 1.0f, 1.0f);
- }
- else {
- holdout_weight = shader_holdout_eval(kg, sd);
- }
- /* any throughput is ok, should all be identical here */
L->transparent += average(holdout_weight * throughput);
}
-
- if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
+ if (isequal_float3(holdout_weight, make_float3(1.0f, 1.0f, 1.0f))) {
return false;
}
}
@@ -673,11 +665,9 @@ ccl_device void kernel_path_trace(
kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray);
-# ifndef __KERNEL_OPTIX__
if (ray.t == 0.0f) {
return;
}
-# endif
/* Initialize state. */
float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index 3d9f787f267..e461e1642b6 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -1017,15 +1017,36 @@ ccl_device float3 shader_emissive_eval(ShaderData *sd)
/* Holdout */
-ccl_device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd)
+ccl_device float3 shader_holdout_apply(KernelGlobals *kg, ShaderData *sd)
{
float3 weight = make_float3(0.0f, 0.0f, 0.0f);
- for (int i = 0; i < sd->num_closure; i++) {
- ShaderClosure *sc = &sd->closure[i];
+ /* For objects marked as holdout, preserve transparency and remove all other
+ * closures, replacing them with a holdout weight. */
+ if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
+ if ((sd->flag & SD_TRANSPARENT) && !(sd->flag & SD_HAS_ONLY_VOLUME)) {
+ weight = make_float3(1.0f, 1.0f, 1.0f) - sd->closure_transparent_extinction;
- if (CLOSURE_IS_HOLDOUT(sc->type))
- weight += sc->weight;
+ for (int i = 0; i < sd->num_closure; i++) {
+ ShaderClosure *sc = &sd->closure[i];
+ if (!CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) {
+ sc->type = NBUILTIN_CLOSURES;
+ }
+ }
+
+ sd->flag &= ~(SD_CLOSURE_FLAGS - (SD_TRANSPARENT | SD_BSDF));
+ }
+ else {
+ weight = make_float3(1.0f, 1.0f, 1.0f);
+ }
+ }
+ else {
+ for (int i = 0; i < sd->num_closure; i++) {
+ ShaderClosure *sc = &sd->closure[i];
+ if (CLOSURE_IS_HOLDOUT(sc->type)) {
+ weight += sc->weight;
+ }
+ }
}
return weight;
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h b/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h
index 89fcb0ae60f..9ab374d1fba 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h
+++ b/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h
@@ -77,7 +77,7 @@ ccl_device_inline float4 svm_image_texture_read(KernelGlobals *kg,
return make_float4(f, f, f, 1.0f);
}
/* Byte */
-#ifdef cl_khr_fp16
+#ifdef __KERNEL_CL_KHR_FP16__
/* half and half4 are optional in OpenCL */
else if (texture_type == IMAGE_DATA_TYPE_HALF) {
float f = tex_fetch(half, info, offset);
diff --git a/intern/cycles/kernel/kernels/optix/kernel_optix.cu b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
index c730d952ed4..3b166e59dfd 100644
--- a/intern/cycles/kernel/kernels/optix/kernel_optix.cu
+++ b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
@@ -15,6 +15,7 @@
* limitations under the License.
*/
+// clang-format off
#include "kernel/kernel_compat_optix.h"
#include "util/util_atomic.h"
#include "kernel/kernel_types.h"
@@ -23,6 +24,7 @@
#include "kernel/kernel_path.h"
#include "kernel/kernel_bake.h"
+// clang-format on
template<typename T> ccl_device_forceinline T *get_payload_ptr_0()
{
@@ -139,8 +141,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
}
else {
if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {
- // Record closest intersection only (do not terminate ray here, since there is no guarantee
- // about distance ordering in anyhit)
+ // Record closest intersection only
+ // Do not terminate ray here, since there is no guarantee about distance ordering in any-hit
return optixIgnoreIntersection();
}
@@ -153,15 +155,9 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
isect->object = get_object_id();
isect->type = kernel_tex_fetch(__prim_type, isect->prim);
- if (optixIsTriangleHit()) {
- const float2 barycentrics = optixGetTriangleBarycentrics();
- isect->u = 1.0f - barycentrics.y - barycentrics.x;
- isect->v = barycentrics.x;
- }
- else {
- isect->u = __uint_as_float(optixGetAttribute_0());
- isect->v = __uint_as_float(optixGetAttribute_1());
- }
+ 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);
@@ -198,10 +194,18 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
isect->u = 1.0f - barycentrics.y - barycentrics.x;
isect->v = barycentrics.x;
}
+# ifdef __HAIR__
else {
- isect->u = __uint_as_float(optixGetAttribute_0());
+ const float u = __uint_as_float(optixGetAttribute_0());
+ isect->u = u;
isect->v = __uint_as_float(optixGetAttribute_1());
+
+ // Filter out curve endcaps
+ if (u == 0.0f || u == 1.0f) {
+ return optixIgnoreIntersection();
+ }
}
+# endif
# ifdef __TRANSPARENT_SHADOWS__
// Detect if this surface has a shader with transparent shadows
@@ -213,7 +217,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
# ifdef __TRANSPARENT_SHADOWS__
}
- // TODO(pmours): Do we need REQUIRE_UNIQUE_ANYHIT for this to work?
optixSetPayload_2(optixGetPayload_2() + 1); // num_hits++
// Continue tracing
@@ -227,13 +230,25 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
uint visibility = optixGetPayload_4();
#ifdef __VISIBILITY_FLAG__
const uint prim = optixGetPrimitiveIndex();
- if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0)
+ if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
return optixIgnoreIntersection();
+ }
+#endif
+
+#ifdef __HAIR__
+ if (!optixIsTriangleHit()) {
+ // Filter out curve endcaps
+ const float u = __uint_as_float(optixGetAttribute_0());
+ if (u == 0.0f || u == 1.0f) {
+ return optixIgnoreIntersection();
+ }
+ }
#endif
// Shadow ray early termination
- if (visibility & PATH_RAY_SHADOW_OPAQUE)
+ if (visibility & PATH_RAY_SHADOW_OPAQUE) {
return optixTerminateRay();
+ }
}
extern "C" __global__ void __closesthit__kernel_optix_hit()
@@ -250,7 +265,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
optixSetPayload_2(__float_as_uint(barycentrics.x));
}
else {
- optixSetPayload_1(optixGetAttribute_0());
+ optixSetPayload_1(optixGetAttribute_0()); // Same as 'optixGetCurveParameter()'
optixSetPayload_2(optixGetAttribute_1());
}
}
@@ -286,7 +301,6 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
__float_as_int(isect.u), // Attribute_0
__float_as_int(isect.v)); // Attribute_1
}
-
}
extern "C" __global__ void __intersection__curve_ribbon()
diff --git a/intern/cycles/kernel/shaders/node_sky_texture.osl b/intern/cycles/kernel/shaders/node_sky_texture.osl
index 08bc8f85120..a12e7a9dc17 100644
--- a/intern/cycles/kernel/shaders/node_sky_texture.osl
+++ b/intern/cycles/kernel/shaders/node_sky_texture.osl
@@ -122,12 +122,18 @@ vector geographical_to_direction(float lat, float lon)
return vector(cos(lat) * cos(lon), cos(lat) * sin(lon), sin(lat));
}
-color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
+float precise_angle(vector a, vector b)
+{
+ return 2.0 * atan2(length(a - b), length(a + b));
+}
+
+color sky_radiance_nishita(vector dir, float nishita_data[10], string filename)
{
/* definitions */
float sun_elevation = nishita_data[6];
float sun_rotation = nishita_data[7];
float angular_diameter = nishita_data[8];
+ float sun_intensity = nishita_data[9];
int sun_disc = angular_diameter > 0;
float alpha = 1.0;
color xyz;
@@ -138,13 +144,13 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
if (dir[2] >= 0.0) {
/* definitions */
vector sun_dir = geographical_to_direction(sun_elevation, sun_rotation + M_PI_2);
- float sun_dir_angle = acos(dot(dir, sun_dir));
+ float sun_dir_angle = precise_angle(dir, sun_dir);
float half_angular = angular_diameter / 2.0;
float dir_elevation = M_PI_2 - direction[0];
/* if ray inside sun disc render it, otherwise render sky */
if (sun_dir_angle < half_angular && sun_disc == 1) {
- /* get 3 pixels data */
+ /* get 2 pixels data */
color pixel_bottom = color(nishita_data[0], nishita_data[1], nishita_data[2]);
color pixel_top = color(nishita_data[3], nishita_data[4], nishita_data[5]);
float y;
@@ -153,13 +159,13 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
if (sun_elevation - half_angular > 0.0) {
if ((sun_elevation + half_angular) > 0.0) {
y = ((dir_elevation - sun_elevation) / angular_diameter) + 0.5;
- xyz = mix(pixel_bottom, pixel_top, y);
+ xyz = mix(pixel_bottom, pixel_top, y) * sun_intensity;
}
}
else {
if (sun_elevation + half_angular > 0.0) {
y = dir_elevation / (sun_elevation + half_angular);
- xyz = mix(pixel_bottom, pixel_top, y);
+ xyz = mix(pixel_bottom, pixel_top, y) * sun_intensity;
}
}
/* limb darkening, coefficient is 0.6f */
@@ -171,7 +177,8 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
else {
/* sky interpolation */
float x = (direction[1] + M_PI + sun_rotation) / M_2PI;
- float y = 1.0 - (dir_elevation / M_PI_2);
+ /* more pixels toward horizon compensation */
+ float y = 1.0 - sqrt(dir_elevation / M_PI_2);
if (x > 1.0) {
x = x - 1.0;
}
@@ -197,23 +204,24 @@ color sky_radiance_nishita(vector dir, float nishita_data[9], string filename)
mul;
}
}
- /* convert to RGB and adjust strength */
- return xyz_to_rgb(xyz[0], xyz[1], xyz[2]) * 120000.0;
+ /* convert to RGB */
+ return xyz_to_rgb(xyz[0], xyz[1], xyz[2]);
}
-shader node_sky_texture(int use_mapping = 0,
- matrix mapping = matrix(0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0),
- vector Vector = P,
- string type = "hosek_wilkie",
- float theta = 0.0,
- float phi = 0.0,
- string filename = "",
- color radiance = color(0.0, 0.0, 0.0),
- float config_x[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
- float config_y[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
- float config_z[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
- float nishita_data[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
- output color Color = color(0.0, 0.0, 0.0))
+shader node_sky_texture(
+ int use_mapping = 0,
+ matrix mapping = matrix(0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0),
+ vector Vector = P,
+ string type = "hosek_wilkie",
+ float theta = 0.0,
+ float phi = 0.0,
+ string filename = "",
+ color radiance = color(0.0, 0.0, 0.0),
+ float config_x[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+ float config_y[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+ float config_z[9] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+ float nishita_data[10] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0},
+ output color Color = color(0.0, 0.0, 0.0))
{
vector p = Vector;
diff --git a/intern/cycles/kernel/svm/svm_sky.h b/intern/cycles/kernel/svm/svm_sky.h
index e877bd9a5c8..f824184c1d4 100644
--- a/intern/cycles/kernel/svm/svm_sky.h
+++ b/intern/cycles/kernel/svm/svm_sky.h
@@ -136,6 +136,7 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
float sun_elevation = nishita_data[6];
float sun_rotation = nishita_data[7];
float angular_diameter = nishita_data[8];
+ float sun_intensity = nishita_data[9];
bool sun_disc = (angular_diameter > 0.0f);
float3 xyz;
/* convert dir to spherical coordinates */
@@ -145,13 +146,13 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
if (dir.z >= 0.0f) {
/* definitions */
float3 sun_dir = geographical_to_direction(sun_elevation, sun_rotation + M_PI_2_F);
- float sun_dir_angle = acos(dot(dir, sun_dir));
+ float sun_dir_angle = precise_angle(dir, sun_dir);
float half_angular = angular_diameter / 2.0f;
float dir_elevation = M_PI_2_F - direction.x;
/* if ray inside sun disc render it, otherwise render sky */
if (sun_disc && sun_dir_angle < half_angular) {
- /* get 3 pixels data */
+ /* get 2 pixels data */
float3 pixel_bottom = make_float3(nishita_data[0], nishita_data[1], nishita_data[2]);
float3 pixel_top = make_float3(nishita_data[3], nishita_data[4], nishita_data[5]);
float y;
@@ -160,13 +161,13 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
if (sun_elevation - half_angular > 0.0f) {
if (sun_elevation + half_angular > 0.0f) {
y = ((dir_elevation - sun_elevation) / angular_diameter) + 0.5f;
- xyz = interp(pixel_bottom, pixel_top, y);
+ xyz = interp(pixel_bottom, pixel_top, y) * sun_intensity;
}
}
else {
if (sun_elevation + half_angular > 0.0f) {
y = dir_elevation / (sun_elevation + half_angular);
- xyz = interp(pixel_bottom, pixel_top, y);
+ xyz = interp(pixel_bottom, pixel_top, y) * sun_intensity;
}
}
/* limb darkening, coefficient is 0.6f */
@@ -178,7 +179,8 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
else {
/* sky interpolation */
float x = (direction.y + M_PI_F + sun_rotation) / M_2PI_F;
- float y = dir_elevation / M_PI_2_F;
+ /* more pixels toward horizon compensation */
+ float y = safe_sqrtf(dir_elevation / M_PI_2_F);
if (x > 1.0f) {
x -= 1.0f;
}
@@ -203,8 +205,8 @@ ccl_device float3 sky_radiance_nishita(KernelGlobals *kg,
}
}
- /* convert to rgb and adjust strength */
- return xyz_to_rgb(kg, xyz) * 120000.0f;
+ /* convert to RGB */
+ return xyz_to_rgb(kg, xyz);
}
ccl_device void svm_node_tex_sky(
@@ -301,7 +303,7 @@ ccl_device void svm_node_tex_sky(
/* Nishita */
else {
/* Define variables */
- float nishita_data[9];
+ float nishita_data[10];
float4 data = read_node_float(kg, offset);
nishita_data[0] = data.x;
@@ -317,7 +319,8 @@ ccl_device void svm_node_tex_sky(
data = read_node_float(kg, offset);
nishita_data[8] = data.x;
- uint texture_id = __float_as_uint(data.y);
+ nishita_data[9] = data.y;
+ uint texture_id = __float_as_uint(data.z);
/* Compute Sky */
f = sky_radiance_nishita(kg, dir, nishita_data, texture_id);
generated by cgit v1.2.3 (git 2.25.1) at 2024-08-28 06:53:39 +0300