[Bf-blender-cvs] [e832b66] cycles_kernel_split: replace sc_fetch with ccl_fetch_array
Martijn Berger
noreply at git.blender.org
Sat May 9 11:14:40 CEST 2015
Commit: e832b6685515a52b33ed894fe3d3c8e299837e41
Author: Martijn Berger
Date: Sat May 9 11:12:39 2015 +0200
Branches: cycles_kernel_split
https://developer.blender.org/rBe832b6685515a52b33ed894fe3d3c8e299837e41
replace sc_fetch with ccl_fetch_array
===================================================================
M intern/cycles/kernel/kernel_shader.h
M intern/cycles/kernel/kernel_types.h
M intern/cycles/kernel/svm/svm_closure.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index 6e34b52..82016ff 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -497,7 +497,7 @@ ccl_device_inline void _shader_bsdf_multi_eval(KernelGlobals *kg, const ShaderDa
if(i == skip_bsdf)
continue;
- const ShaderClosure *sc = sc_fetch(i);
+ const ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSDF(sc->type)) {
float bsdf_pdf = 0.0f;
@@ -534,7 +534,7 @@ ccl_device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
float sum = 0.0f;
for(sampled = 0; sampled < ccl_fetch(sd,num_closure); sampled++) {
- const ShaderClosure *sc = sc_fetch(sampled);
+ const ShaderClosure *sc = ccl_fetch_array(sd, closure, sampled);
if(CLOSURE_IS_BSDF(sc->type))
sum += sc->sample_weight;
@@ -544,7 +544,7 @@ ccl_device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
sum = 0.0f;
for(sampled = 0; sampled < ccl_fetch(sd,num_closure); sampled++) {
- const ShaderClosure *sc = sc_fetch(sampled);
+ const ShaderClosure *sc = ccl_fetch_array(sd, closure, sampled);
if(CLOSURE_IS_BSDF(sc->type)) {
sum += sc->sample_weight;
@@ -560,7 +560,7 @@ ccl_device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
}
}
- const ShaderClosure *sc = sc_fetch(sampled);
+ const ShaderClosure *sc = ccl_fetch_array(sd, closure, sampled);
int label;
float3 eval;
@@ -599,7 +599,7 @@ ccl_device int shader_bsdf_sample_closure(KernelGlobals *kg, const ShaderData *s
ccl_device void shader_bsdf_blur(KernelGlobals *kg, ShaderData *sd, float roughness)
{
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSDF(sc->type))
bsdf_blur(kg, sc, roughness);
@@ -614,7 +614,7 @@ ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd)
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(sc->type == CLOSURE_BSDF_TRANSPARENT_ID) // todo: make this work for osl
eval += sc->weight;
@@ -638,7 +638,7 @@ ccl_device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd)
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSDF_DIFFUSE(sc->type))
eval += sc->weight;
@@ -652,7 +652,7 @@ ccl_device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd)
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSDF_GLOSSY(sc->type))
eval += sc->weight;
@@ -666,7 +666,7 @@ ccl_device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd)
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSDF_TRANSMISSION(sc->type))
eval += sc->weight;
@@ -680,7 +680,7 @@ ccl_device float3 shader_bsdf_subsurface(KernelGlobals *kg, ShaderData *sd)
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSSRDF(sc->type) || CLOSURE_IS_BSDF_BSSRDF(sc->type))
eval += sc->weight;
@@ -695,7 +695,7 @@ ccl_device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_fac
float3 N = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSDF_DIFFUSE(sc->type)) {
eval += sc->weight*ao_factor;
@@ -723,7 +723,7 @@ ccl_device float3 shader_bssrdf_sum(ShaderData *sd, float3 *N_, float *texture_b
float texture_blur = 0.0f, weight_sum = 0.0f;
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BSSRDF(sc->type)) {
float avg_weight = fabsf(average(sc->weight));
@@ -757,7 +757,7 @@ ccl_device float3 shader_emissive_eval(KernelGlobals *kg, ShaderData *sd)
eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i < ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_EMISSION(sc->type))
eval += emissive_eval(kg, sd, sc)*sc->weight;
@@ -773,7 +773,7 @@ ccl_device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd)
float3 weight = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i < ccl_fetch(sd,num_closure); i++) {
- ShaderClosure *sc = sc_fetch(i);
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_HOLDOUT(sc->type))
weight += sc->weight;
@@ -799,11 +799,11 @@ ccl_device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd,
#ifdef __SVM__
svm_eval_nodes(kg, sd, SHADER_TYPE_SURFACE, path_flag);
#else
- sc_fetch(0)->weight = make_float3(0.8f, 0.8f, 0.8f);
- sc_fetch(0)->N = ccl_fetch(sd,N);
- sc_fetch(0)->data0 = 0.0f;
- sc_fetch(0)->data1 = 0.0f;
- ccl_fetch(sd,flag) |= bsdf_diffuse_setup(sc_fetch(0));
+ ccl_fetch_array(sd, closure, 0)->weight = make_float3(0.8f, 0.8f, 0.8f);
+ ccl_fetch_array(sd, closure, 0)->N = ccl_fetch(sd,N);
+ ccl_fetch_array(sd, closure, 0)->data0 = 0.0f;
+ ccl_fetch_array(sd, closure, 0)->data1 = 0.0f;
+ ccl_fetch(sd,flag) |= bsdf_diffuse_setup(ccl_fetch_array(sd, closure, 0));
#endif
}
}
@@ -829,7 +829,7 @@ ccl_device float3 shader_eval_background(KernelGlobals *kg, ShaderData *sd, int
float3 eval = make_float3(0.0f, 0.0f, 0.0f);
for(int i = 0; i< ccl_fetch(sd,num_closure); i++) {
- const ShaderClosure *sc = sc_fetch(i);
+ const ShaderClosure *sc = ccl_fetch_array(sd, closure, i);
if(CLOSURE_IS_BACKGROUND(sc->type))
eval += sc->weight;
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index bdfae8d..e3fa7fe 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -657,13 +657,13 @@ struct KernelGlobals;
#define SD_VAR(type, what) ccl_global type *what;
#define SD_CLOSURE_VAR(type, what, max_closure) type *what;
#define TIDX (get_global_id(1) * get_global_size(0) + get_global_id(0))
-#define ccl_fetch(s,t) (s->t[TIDX])
-#define sc_fetch(index) (&sd->closure[TIDX * MAX_CLOSURE + index])
+#define ccl_fetch(s, t) (s->t[TIDX])
+#define ccl_fetch_array(s, t, index) (&s->t[TIDX * MAX_CLOSURE + index])
#else
#define SD_VAR(type, what) type what;
#define SD_CLOSURE_VAR(type, what, max_closure) type what[max_closure];
-#define ccl_fetch(s,t) (s->t)
-#define sc_fetch(index) (&sd->closure[index])
+#define ccl_fetch(s, t) (s->t)
+#define ccl_fetch_array(s, t, index) (&s->t[index])
#endif
typedef ccl_addr_space struct ShaderData {
diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h
index a1bd55b..2185fd6 100644
--- a/intern/cycles/kernel/svm/svm_closure.h
+++ b/intern/cycles/kernel/svm/svm_closure.h
@@ -57,7 +57,7 @@ ccl_device void svm_node_glass_setup(ShaderData *sd, ShaderClosure *sc, int type
ccl_device_inline ShaderClosure *svm_node_closure_get_non_bsdf(ShaderData *sd, ClosureType type, float mix_weight)
{
- ShaderClosure *sc = sc_fetch(ccl_fetch(sd,num_closure));
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure));
if(ccl_fetch(sd,num_closure) < MAX_CLOSURE) {
sc->weight *= mix_weight;
@@ -74,7 +74,7 @@ ccl_device_inline ShaderClosure *svm_node_closure_get_non_bsdf(ShaderData *sd, C
ccl_device_inline ShaderClosure *svm_node_closure_get_bsdf(ShaderData *sd, float mix_weight)
{
- ShaderClosure *sc = sc_fetch(ccl_fetch(sd,num_closure));
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure));
float3 weight = sc->weight * mix_weight;
float sample_weight = fabsf(average(weight));
@@ -94,7 +94,7 @@ ccl_device_inline ShaderClosure *svm_node_closure_get_bsdf(ShaderData *sd, float
ccl_device_inline ShaderClosure *svm_node_closure_get_absorption(ShaderData *sd, float mix_weight)
{
- ShaderClosure *sc = sc_fetch(ccl_fetch(sd,num_closure));
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure));
float3 weight = (make_float3(1.0f, 1.0f, 1.0f) - sc->weight) * mix_weight;
float sample_weight = fabsf(average(weight));
@@ -264,7 +264,7 @@ ccl_device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *
float roughness = param1;
/* reflection */
- ShaderClosure *sc = sc_fetch(ccl_fetch(sd,num_closure));
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure));
float3 weight = sc->weight;
float sample_weight = sc->sample_weight;
@@ -285,7 +285,7 @@ ccl_device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *
#endif
/* refraction */
- sc = sc_fetch(ccl_fetch(sd,num_closure));
+ sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure));
sc->weight = weight;
sc->sample_weight = sample_weight;
@@ -395,7 +395,7 @@ ccl_device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *
}
}
else {
- ShaderClosure *sc = sc_fetch(ccl_fetch(sd,num_closure));
+ ShaderClosure *sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure));
sc = svm_node_closure_get_bsdf(sd, mix_weight);
if(sc) {
@@ -428,11 +428,11 @@ ccl_device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *
#ifndef __SPLIT_KERNEL__
# define sc_next(sc) sc++
# else
-# define sc_next(sc) sc = sc_fetch(ccl_fetch(sd,num_closure))
+# define sc_next(sc) sc = ccl_fetch_array(sd, closure, ccl_fetch(sd,num_closure))
# endif
case CLOS
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list