diff options
Diffstat (limited to 'intern/cycles/kernel')
22 files changed, 128 insertions, 67 deletions
diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp index 656e90eb526..d9ba237d964 100644 --- a/intern/cycles/kernel/kernel.cpp +++ b/intern/cycles/kernel/kernel.cpp @@ -129,18 +129,6 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t kg->__filter_table.data = (float*)mem; kg->__filter_table.width = width; } - else if(strcmp(name, "__response_curve_R") == 0) { - kg->__response_curve_R.data = (float*)mem; - kg->__response_curve_R.width = width; - } - else if(strcmp(name, "__response_curve_B") == 0) { - kg->__response_curve_B.data = (float*)mem; - kg->__response_curve_B.width = width; - } - else if(strcmp(name, "__response_curve_G") == 0) { - kg->__response_curve_G.data = (float*)mem; - kg->__response_curve_G.width = width; - } else if(strcmp(name, "__sobol_directions") == 0) { kg->__sobol_directions.data = (uint*)mem; kg->__sobol_directions.width = width; diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h index b7b29d46323..3a0eff5210c 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -33,6 +33,7 @@ CCL_NAMESPACE_BEGIN #define __device __device__ __inline__ #define __device_inline __device__ __inline__ +#define __device_noinline __device__ __noinline__ #define __global #define __shared __shared__ #define __constant diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h index 3d493c61fe4..287bf320881 100644 --- a/intern/cycles/kernel/kernel_compat_opencl.h +++ b/intern/cycles/kernel/kernel_compat_opencl.h @@ -30,6 +30,7 @@ /* in opencl all functions are device functions, so leave this empty */ #define __device #define __device_inline +#define __device_noinline /* no assert in opencl */ #define kernel_assert(cond) diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h index 0bbd6c202a1..f6351a73295 100644 --- a/intern/cycles/kernel/kernel_film.h +++ b/intern/cycles/kernel/kernel_film.h @@ -20,21 +20,17 @@ CCL_NAMESPACE_BEGIN __device float4 film_map(KernelGlobals *kg, float4 irradiance, int pass) { - float scale = kernel_data.film.exposure*(1.0f/(pass+1)); + float scale = 1.0f/(float)(pass+1); + float exposure = kernel_data.film.exposure; float4 result = irradiance*scale; - if(kernel_data.film.use_response_curve) { - /* camera response curve */ - result.x = kernel_tex_interp(__response_curve_R, result.x); - result.y = kernel_tex_interp(__response_curve_G, result.y); - result.z = kernel_tex_interp(__response_curve_B, result.z); - } - else { - /* conversion to srgb */ - result.x = color_scene_linear_to_srgb(result.x); - result.y = color_scene_linear_to_srgb(result.y); - result.z = color_scene_linear_to_srgb(result.z); - } + /* conversion to srgb */ + result.x = color_scene_linear_to_srgb(result.x*exposure); + result.y = color_scene_linear_to_srgb(result.y*exposure); + result.z = color_scene_linear_to_srgb(result.z*exposure); + + /* clamp since alpha might be > 1.0 due to russian roulette */ + result.w = clamp(result.w, 0.0f, 1.0f); return result; } @@ -47,7 +43,7 @@ __device uchar4 film_float_to_byte(float4 color) result.x = (uchar)clamp(color.x*255.0f, 0.0f, 255.0f); result.y = (uchar)clamp(color.y*255.0f, 0.0f, 255.0f); result.z = (uchar)clamp(color.z*255.0f, 0.0f, 255.0f); - result.w = 255; + result.w = (uchar)clamp(color.w*255.0f, 0.0f, 255.0f); return result; } diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 47aeea4f5c5..2cbb4ae306e 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -105,10 +105,11 @@ __device int path_flag_from_label(int path_flag, int label) return path_flag; } -__device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray ray, float3 throughput) +__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray ray, float3 throughput) { /* initialize */ float3 L = make_float3(0.0f, 0.0f, 0.0f); + float Ltransparent = 0.0f; #ifdef __EMISSION__ float ray_pdf = 0.0f; @@ -123,14 +124,20 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray if(!scene_intersect(kg, &ray, false, &isect)) { /* eval background shader if nothing hit */ + if(kernel_data.background.transparent && (path_flag & PATH_RAY_CAMERA)) { + Ltransparent += average(throughput); + } + else { #ifdef __BACKGROUND__ - ShaderData sd; - shader_setup_from_background(kg, &sd, &ray); - L += throughput*shader_eval_background(kg, &sd, path_flag); - shader_release(kg, &sd); + ShaderData sd; + shader_setup_from_background(kg, &sd, &ray); + L += throughput*shader_eval_background(kg, &sd, path_flag); + shader_release(kg, &sd); #else - L += make_float3(0.8f, 0.8f, 0.8f); + L += make_float3(0.8f, 0.8f, 0.8f); #endif + } + break; } @@ -140,6 +147,22 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray float rbsdf = path_rng(kg, rng, pass, rng_offset + PRNG_BSDF); shader_eval_surface(kg, &sd, rbsdf, path_flag); +#ifdef __HOLDOUT__ + if((sd.flag & SD_HOLDOUT) && (path_flag & PATH_RAY_CAMERA)) { + float3 holdout_weight = shader_holdout_eval(kg, &sd); + + if(kernel_data.background.transparent) { + Ltransparent += average(holdout_weight*throughput); + } + else { + ShaderData sd; + shader_setup_from_background(kg, &sd, &ray); + L += holdout_weight*throughput*shader_eval_background(kg, &sd, path_flag); + shader_release(kg, &sd); + } + } +#endif + #ifdef __EMISSION__ /* emission */ if(kernel_data.integrator.use_emission) { @@ -166,6 +189,12 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray } #endif + /* no BSDF? we can stop here */ + if(!(sd.flag & SD_BSDF)) { + path_flag &= ~PATH_RAY_CAMERA; + break; + } + /* sample BSDF */ float bsdf_pdf; float3 bsdf_eval; @@ -180,8 +209,10 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray shader_release(kg, &sd); - if(bsdf_pdf == 0.0f || is_zero(bsdf_eval)) + if(bsdf_pdf == 0.0f || is_zero(bsdf_eval)) { + path_flag &= ~PATH_RAY_CAMERA; break; + } /* modify throughput */ throughput *= bsdf_eval/bsdf_pdf; @@ -197,8 +228,10 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray float probability = path_terminate_probability(kg, bounce, throughput); float terminate = path_rng(kg, rng, pass, rng_offset + PRNG_TERMINATE); - if(terminate >= probability) + if(terminate >= probability) { + path_flag &= ~PATH_RAY_CAMERA; break; + } throughput /= probability; @@ -212,7 +245,7 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray #endif } - return L; + return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent); } __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int pass, int x, int y) @@ -236,25 +269,19 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl /* integrate */ #ifdef __MODIFY_TP__ float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, pass); - float3 L = kernel_path_integrate(kg, &rng, pass, ray, throughput)/throughput; + float4 L = kernel_path_integrate(kg, &rng, pass, ray, throughput)/throughput; #else float3 throughput = make_float3(1.0f, 1.0f, 1.0f); - float3 L = kernel_path_integrate(kg, &rng, pass, ray, throughput); + float4 L = kernel_path_integrate(kg, &rng, pass, ray, throughput); #endif /* accumulate result in output buffer */ int index = x + y*kernel_data.cam.width; - float4 result; - result.x = L.x; - result.y = L.y; - result.z = L.z; - result.w = 1.0f; - if(pass == 0) - buffer[index] = result; + buffer[index] = L; else - buffer[index] += result; + buffer[index] += L; path_rng_end(kg, rng_state, rng, x, y); } diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index c4fc65596b2..ed3aec21207 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -342,6 +342,26 @@ __device void shader_emissive_sample(KernelGlobals *kg, ShaderData *sd, } } +/* Holdout */ + +__device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef WITH_OSL + if(kg->osl.use) { + return OSLShader::holdout_eval(sd); + } + else +#endif + { +#ifdef __SVM__ + if(sd->svm_closure == CLOSURE_HOLDOUT_ID) + return make_float3(1.0f, 1.0f, 1.0f); + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif + } +} + /* Surface Evaluation */ __device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd, diff --git a/intern/cycles/kernel/kernel_textures.h b/intern/cycles/kernel/kernel_textures.h index bd44ed7eee5..dd4566ec184 100644 --- a/intern/cycles/kernel/kernel_textures.h +++ b/intern/cycles/kernel/kernel_textures.h @@ -38,9 +38,6 @@ KERNEL_TEX(uint4, texture_uint4, __svm_nodes) /* camera/film */ KERNEL_TEX(float, texture_float, __filter_table) -KERNEL_TEX(float, texture_float, __response_curve_R) -KERNEL_TEX(float, texture_float, __response_curve_G) -KERNEL_TEX(float, texture_float, __response_curve_B) /* sobol */ KERNEL_TEX(uint, texture_uint, __sobol_directions) diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index c0cb3fc8a09..06581d593e4 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -37,6 +37,7 @@ CCL_NAMESPACE_BEGIN #ifndef __KERNEL_OPENCL__ #define __SVM__ #define __TEXTURES__ +#define __HOLDOUT__ #endif #define __RAY_DIFFERENTIALS__ #define __CAMERA_CLIPPING__ @@ -194,8 +195,10 @@ struct FlatClosure { enum ShaderDataFlag { SD_BACKFACING = 1, /* backside of surface? */ SD_EMISSION = 2, /* have emissive closure? */ - SD_BSDF_HAS_EVAL = 4, /* have non-singular bsdf closure? */ - SD_BSDF_GLOSSY = 8 /* have glossy bsdf */ + SD_BSDF = 4, /* have bsdf closure? */ + SD_BSDF_HAS_EVAL = 8, /* have non-singular bsdf closure? */ + SD_BSDF_GLOSSY = 16, /* have glossy bsdf */ + SD_HOLDOUT = 32 /* have holdout closure? */ }; typedef struct ShaderData { @@ -257,6 +260,8 @@ typedef struct ShaderData { float emissive_sample_sum; float volume_sample_sum; + float3 holdout_weight; + float randb; } osl_closure; @@ -313,14 +318,14 @@ typedef struct KernelCamera { typedef struct KernelFilm { float exposure; - int use_response_curve; - int pad1, pad2; + int pad1, pad2, pad3; } KernelFilm; typedef struct KernelBackground { /* only shader index */ int shader; - int pad1, pad2, pad3; + int transparent; + int pad1, pad2; } KernelBackground; typedef struct KernelSunSky { diff --git a/intern/cycles/kernel/osl/osl_shader.cpp b/intern/cycles/kernel/osl/osl_shader.cpp index a86946a680e..007d14b526b 100644 --- a/intern/cycles/kernel/osl/osl_shader.cpp +++ b/intern/cycles/kernel/osl/osl_shader.cpp @@ -145,9 +145,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, bool no_glossy, /* scattering flags */ if(scattering == OSL::Labels::DIFFUSE) - sd->flag |= SD_BSDF_HAS_EVAL; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL; else if(scattering == OSL::Labels::GLOSSY) - sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + else + sd->flag |= SD_BSDF; /* add */ sd->osl_closure.bsdf[sd->osl_closure.num_bsdf++] = flat; @@ -170,8 +172,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, bool no_glossy, sd->osl_closure.emissive[sd->osl_closure.num_emissive++] = flat; break; } - case ClosurePrimitive::BSSRDF: case ClosurePrimitive::Holdout: + sd->osl_closure.holdout_weight += weight; + sd->flag |= SD_HOLDOUT; + break; + case ClosurePrimitive::BSSRDF: case ClosurePrimitive::Debug: break; /* not implemented */ case ClosurePrimitive::Background: @@ -212,6 +217,7 @@ void OSLShader::eval_surface(KernelGlobals *kg, ShaderData *sd, float randb, int sd->osl_closure.emissive_sample_sum = 0.0f; sd->osl_closure.num_bsdf = 0; sd->osl_closure.num_emissive = 0; + sd->osl_closure.holdout_weight = make_float3(0.0f, 0.0f, 0.0f); sd->osl_closure.randb = randb; if(globals->Ci) { @@ -555,5 +561,12 @@ float3 OSLShader::volume_eval_phase(const ShaderData *sd, const float3 omega_in, return eval; } +/* Holdout Closure */ + +float3 OSLShader::holdout_eval(const ShaderData *sd) +{ + return sd->osl_closure.holdout_weight; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/osl/osl_shader.h b/intern/cycles/kernel/osl/osl_shader.h index 2e5279c12ce..9b578b159ae 100644 --- a/intern/cycles/kernel/osl/osl_shader.h +++ b/intern/cycles/kernel/osl/osl_shader.h @@ -74,6 +74,7 @@ public: float3 *eval, float3 *I, float *pdf); static float3 volume_eval_phase(const ShaderData *sd, const float3 omega_in, const float3 omega_out); + static float3 holdout_eval(const ShaderData *sd); /* release */ static void release(KernelGlobals *kg, const ShaderData *sd); diff --git a/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h b/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h index 04e4ccb8313..a04f4e2b076 100644 --- a/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h +++ b/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h @@ -50,7 +50,7 @@ __device void bsdf_ashikhmin_velvet_setup(ShaderData *sd, float3 N, float sigma) self->m_invsigma2 = 1.0f/(sigma * sigma); sd->svm_closure = CLOSURE_BSDF_ASHIKHMIN_VELVET_ID; - sd->flag |= SD_BSDF_HAS_EVAL; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL; } __device void bsdf_ashikhmin_velvet_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_diffuse.h b/intern/cycles/kernel/svm/bsdf_diffuse.h index 00493e72203..dcd29534109 100644 --- a/intern/cycles/kernel/svm/bsdf_diffuse.h +++ b/intern/cycles/kernel/svm/bsdf_diffuse.h @@ -47,7 +47,7 @@ __device void bsdf_diffuse_setup(ShaderData *sd, float3 N) //self->m_N = N; sd->svm_closure = CLOSURE_BSDF_DIFFUSE_ID; - sd->flag |= SD_BSDF_HAS_EVAL; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL; } __device void bsdf_diffuse_blur(ShaderData *sd, float roughness) @@ -110,7 +110,7 @@ __device void bsdf_translucent_setup(ShaderData *sd, float3 N) //self->m_N = N; sd->svm_closure = CLOSURE_BSDF_TRANSLUCENT_ID; - sd->flag |= SD_BSDF_HAS_EVAL; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL; } __device void bsdf_translucent_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_microfacet.h b/intern/cycles/kernel/svm/bsdf_microfacet.h index a948ba06871..c98092b2f8f 100644 --- a/intern/cycles/kernel/svm/bsdf_microfacet.h +++ b/intern/cycles/kernel/svm/bsdf_microfacet.h @@ -58,7 +58,7 @@ __device void bsdf_microfacet_ggx_setup(ShaderData *sd, float3 N, float ag, floa else sd->svm_closure = CLOSURE_BSDF_MICROFACET_GGX_ID; - sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } __device void bsdf_microfacet_ggx_blur(ShaderData *sd, float roughness) @@ -278,7 +278,7 @@ __device void bsdf_microfacet_beckmann_setup(ShaderData *sd, float3 N, float ab, else sd->svm_closure = CLOSURE_BSDF_MICROFACET_BECKMANN_ID; - sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } __device void bsdf_microfacet_beckmann_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_reflection.h b/intern/cycles/kernel/svm/bsdf_reflection.h index 287cc9c2506..2c49eb7ef24 100644 --- a/intern/cycles/kernel/svm/bsdf_reflection.h +++ b/intern/cycles/kernel/svm/bsdf_reflection.h @@ -47,6 +47,7 @@ __device void bsdf_reflection_setup(ShaderData *sd, float3 N) //self->m_N = N; sd->svm_closure = CLOSURE_BSDF_REFLECTION_ID; + sd->flag |= SD_BSDF; } __device void bsdf_reflection_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_refraction.h b/intern/cycles/kernel/svm/bsdf_refraction.h index 55a914f8334..534945f4f0b 100644 --- a/intern/cycles/kernel/svm/bsdf_refraction.h +++ b/intern/cycles/kernel/svm/bsdf_refraction.h @@ -48,6 +48,7 @@ __device void bsdf_refraction_setup(ShaderData *sd, float3 N, float eta) self->m_eta = eta; sd->svm_closure = CLOSURE_BSDF_REFRACTION_ID; + sd->flag |= SD_BSDF; } __device void bsdf_refraction_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_transparent.h b/intern/cycles/kernel/svm/bsdf_transparent.h index e689e3db357..30288bf251d 100644 --- a/intern/cycles/kernel/svm/bsdf_transparent.h +++ b/intern/cycles/kernel/svm/bsdf_transparent.h @@ -38,6 +38,7 @@ CCL_NAMESPACE_BEGIN __device void bsdf_transparent_setup(ShaderData *sd) { sd->svm_closure = CLOSURE_BSDF_TRANSPARENT_ID; + sd->flag |= SD_BSDF; } __device void bsdf_transparent_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_ward.h b/intern/cycles/kernel/svm/bsdf_ward.h index c54418afa77..6680644a1a7 100644 --- a/intern/cycles/kernel/svm/bsdf_ward.h +++ b/intern/cycles/kernel/svm/bsdf_ward.h @@ -54,7 +54,7 @@ __device void bsdf_ward_setup(ShaderData *sd, float3 N, float3 T, float ax, floa self->m_ay = clamp(ay, 1e-5f, 1.0f); sd->svm_closure = CLOSURE_BSDF_WARD_ID; - sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } __device void bsdf_ward_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/bsdf_westin.h b/intern/cycles/kernel/svm/bsdf_westin.h index 6031012d0ca..55fc8286529 100644 --- a/intern/cycles/kernel/svm/bsdf_westin.h +++ b/intern/cycles/kernel/svm/bsdf_westin.h @@ -51,7 +51,7 @@ __device void bsdf_westin_backscatter_setup(ShaderData *sd, float3 N, float roug self->m_invroughness = 1.0f/roughness; sd->svm_closure = CLOSURE_BSDF_WESTIN_BACKSCATTER_ID; - sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } __device void bsdf_westin_backscatter_blur(ShaderData *sd, float roughness) @@ -146,7 +146,7 @@ __device void bsdf_westin_sheen_setup(ShaderData *sd, float3 N, float edginess) self->m_edginess = edginess; sd->svm_closure = CLOSURE_BSDF_WESTIN_SHEEN_ID; - sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; + sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } __device void bsdf_westin_sheen_blur(ShaderData *sd, float roughness) diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h index 68e54d9752e..3a2b6bfb1ae 100644 --- a/intern/cycles/kernel/svm/svm.h +++ b/intern/cycles/kernel/svm/svm.h @@ -148,7 +148,7 @@ CCL_NAMESPACE_BEGIN /* Main Interpreter Loop */ -__device void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag) +__device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag) { float stack[SVM_STACK_SIZE]; float closure_weight = 1.0f; @@ -172,6 +172,8 @@ __device void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, svm_node_closure_emission(sd); else if(node.x == NODE_CLOSURE_BACKGROUND) svm_node_closure_background(sd); + else if(node.x == NODE_CLOSURE_HOLDOUT) + svm_node_closure_holdout(sd); else if(node.x == NODE_CLOSURE_SET_WEIGHT) svm_node_closure_set_weight(sd, node.y, node.z, node.w); else if(node.x == NODE_CLOSURE_WEIGHT) diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h index a7462e2c358..1efadcefbd5 100644 --- a/intern/cycles/kernel/svm/svm_closure.h +++ b/intern/cycles/kernel/svm/svm_closure.h @@ -136,6 +136,12 @@ __device void svm_node_closure_background(ShaderData *sd) sd->svm_closure = CLOSURE_BACKGROUND_ID; } +__device void svm_node_closure_holdout(ShaderData *sd) +{ + sd->svm_closure = CLOSURE_HOLDOUT_ID; + sd->flag |= SD_HOLDOUT; +} + /* Closure Nodes */ __device void svm_node_closure_set_weight(ShaderData *sd, uint r, uint g, uint b) diff --git a/intern/cycles/kernel/svm/svm_convert.h b/intern/cycles/kernel/svm/svm_convert.h index 14925569bfb..188b0489d9e 100644 --- a/intern/cycles/kernel/svm/svm_convert.h +++ b/intern/cycles/kernel/svm/svm_convert.h @@ -30,7 +30,7 @@ __device void svm_node_convert(ShaderData *sd, float *stack, uint type, uint fro } case NODE_CONVERT_CF: { float3 f = stack_load_float3(stack, from); - float g = f.x*0.2126f + f.y*0.7152f + f.z*0.0722f; + float g = linear_rgb_to_gray(f); stack_store_float(stack, to, g); break; } diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h index 5cbd05e4400..786478c0c03 100644 --- a/intern/cycles/kernel/svm/svm_types.h +++ b/intern/cycles/kernel/svm/svm_types.h @@ -81,7 +81,8 @@ typedef enum NodeType { NODE_EMISSION_SET_WEIGHT_TOTAL = 4300, NODE_ATTR_BUMP_DX = 4400, NODE_ATTR_BUMP_DY = 4500, - NODE_TEX_ENVIRONMENT = 4600 + NODE_TEX_ENVIRONMENT = 4600, + NODE_CLOSURE_HOLDOUT = 4700 } NodeType; typedef enum NodeAttributeType { |