diff options
| author | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2013-11-16 03:17:10 +0400 |
|---|---|---|
| committer | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2013-11-18 11:48:15 +0400 |
| commit | c18712e86814d176433cea781fe1e68715c23bd4 (patch) | |
| tree | a9aa7a0b34940dceb00f290c720f4a2733819f44 /intern/cycles | |
| parent | 6f67f7c18cf2c6224f3e3d796bf0d83d2f098b2e (diff) | |
Cycles: change __device and similar qualifiers to ccl_device in kernel code.
This to avoids build conflicts with libc++ on FreeBSD, these __ prefixed values
are reserved for compilers. I apologize to anyone who has patches or branches
and has to go through the pain of merging this change, it may be easiest to do
these same replacements in your code and then apply/merge the patch.
Ref T37477.
Diffstat (limited to 'intern/cycles')
87 files changed, 804 insertions, 805 deletions
diff --git a/intern/cycles/kernel/closure/bsdf.h b/intern/cycles/kernel/closure/bsdf.h index 86fea48760f..b3141d1154f 100644 --- a/intern/cycles/kernel/closure/bsdf.h +++ b/intern/cycles/kernel/closure/bsdf.h @@ -35,7 +35,7 @@ CCL_NAMESPACE_BEGIN -__device int bsdf_sample(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc, float randu, float randv, float3 *eval, float3 *omega_in, differential3 *domega_in, float *pdf) +ccl_device int bsdf_sample(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc, float randu, float randv, float3 *eval, float3 *omega_in, differential3 *domega_in, float *pdf) { int label; @@ -132,7 +132,7 @@ __device int bsdf_sample(KernelGlobals *kg, const ShaderData *sd, const ShaderCl return label; } -__device float3 bsdf_eval(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_eval(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc, const float3 omega_in, float *pdf) { float3 eval; @@ -275,7 +275,7 @@ __device float3 bsdf_eval(KernelGlobals *kg, const ShaderData *sd, const ShaderC return eval; } -__device void bsdf_blur(KernelGlobals *kg, ShaderClosure *sc, float roughness) +ccl_device void bsdf_blur(KernelGlobals *kg, ShaderClosure *sc, float roughness) { #ifdef __OSL__ if(kg->osl && sc->prim) { diff --git a/intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h b/intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h index 94bc6eb0dc5..3631f90bf8c 100644 --- a/intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h +++ b/intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h @@ -35,7 +35,7 @@ CCL_NAMESPACE_BEGIN -__device int bsdf_ashikhmin_velvet_setup(ShaderClosure *sc) +ccl_device int bsdf_ashikhmin_velvet_setup(ShaderClosure *sc) { float sigma = fmaxf(sc->data0, 0.01f); sc->data0 = 1.0f/(sigma * sigma); /* m_invsigma2 */ @@ -45,11 +45,11 @@ __device int bsdf_ashikhmin_velvet_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL; } -__device void bsdf_ashikhmin_velvet_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_ashikhmin_velvet_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_ashikhmin_velvet_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_ashikhmin_velvet_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_invsigma2 = sc->data0; float3 N = sc->N; @@ -87,12 +87,12 @@ __device float3 bsdf_ashikhmin_velvet_eval_reflect(const ShaderClosure *sc, cons return make_float3(0, 0, 0); } -__device float3 bsdf_ashikhmin_velvet_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_ashikhmin_velvet_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_ashikhmin_velvet_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_ashikhmin_velvet_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_invsigma2 = sc->data0; float3 N = sc->N; diff --git a/intern/cycles/kernel/closure/bsdf_diffuse.h b/intern/cycles/kernel/closure/bsdf_diffuse.h index 46318ecd138..949fe869549 100644 --- a/intern/cycles/kernel/closure/bsdf_diffuse.h +++ b/intern/cycles/kernel/closure/bsdf_diffuse.h @@ -37,17 +37,17 @@ CCL_NAMESPACE_BEGIN /* DIFFUSE */ -__device int bsdf_diffuse_setup(ShaderClosure *sc) +ccl_device int bsdf_diffuse_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_DIFFUSE_ID; return SD_BSDF|SD_BSDF_HAS_EVAL; } -__device void bsdf_diffuse_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_diffuse_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_diffuse_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_diffuse_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float3 N = sc->N; @@ -56,12 +56,12 @@ __device float3 bsdf_diffuse_eval_reflect(const ShaderClosure *sc, const float3 return make_float3(cos_pi, cos_pi, cos_pi); } -__device float3 bsdf_diffuse_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_diffuse_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_diffuse_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_diffuse_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float3 N = sc->N; @@ -84,22 +84,22 @@ __device int bsdf_diffuse_sample(const ShaderClosure *sc, float3 Ng, float3 I, f /* TRANSLUCENT */ -__device int bsdf_translucent_setup(ShaderClosure *sc) +ccl_device int bsdf_translucent_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_TRANSLUCENT_ID; return SD_BSDF|SD_BSDF_HAS_EVAL; } -__device void bsdf_translucent_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_translucent_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_translucent_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_translucent_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_translucent_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_translucent_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float3 N = sc->N; @@ -108,12 +108,12 @@ __device float3 bsdf_translucent_eval_transmit(const ShaderClosure *sc, const fl return make_float3 (cos_pi, cos_pi, cos_pi); } -__device float bsdf_translucent_albedo(const ShaderClosure *sc, const float3 I) +ccl_device float bsdf_translucent_albedo(const ShaderClosure *sc, const float3 I) { return 1.0f; } -__device int bsdf_translucent_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_translucent_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float3 N = sc->N; diff --git a/intern/cycles/kernel/closure/bsdf_diffuse_ramp.h b/intern/cycles/kernel/closure/bsdf_diffuse_ramp.h index 2e43e16693f..b856774375f 100644 --- a/intern/cycles/kernel/closure/bsdf_diffuse_ramp.h +++ b/intern/cycles/kernel/closure/bsdf_diffuse_ramp.h @@ -35,7 +35,7 @@ CCL_NAMESPACE_BEGIN -__device float3 bsdf_diffuse_ramp_get_color(const ShaderClosure *sc, const float3 colors[8], float pos) +ccl_device float3 bsdf_diffuse_ramp_get_color(const ShaderClosure *sc, const float3 colors[8], float pos) { int MAXCOLORS = 8; @@ -49,17 +49,17 @@ __device float3 bsdf_diffuse_ramp_get_color(const ShaderClosure *sc, const float return colors[ipos] * (1.0f - offset) + colors[ipos+1] * offset; } -__device int bsdf_diffuse_ramp_setup(ShaderClosure *sc) +ccl_device int bsdf_diffuse_ramp_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_DIFFUSE_RAMP_ID; return SD_BSDF | SD_BSDF_HAS_EVAL; } -__device void bsdf_diffuse_ramp_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_diffuse_ramp_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_diffuse_ramp_eval_reflect(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_diffuse_ramp_eval_reflect(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) { float3 N = sc->N; @@ -68,12 +68,12 @@ __device float3 bsdf_diffuse_ramp_eval_reflect(const ShaderClosure *sc, const fl return bsdf_diffuse_ramp_get_color(sc, colors, cos_pi) * M_1_PI_F; } -__device float3 bsdf_diffuse_ramp_eval_transmit(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_diffuse_ramp_eval_transmit(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_diffuse_ramp_sample(const ShaderClosure *sc, const float3 colors[8], float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_diffuse_ramp_sample(const ShaderClosure *sc, const float3 colors[8], float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float3 N = sc->N; diff --git a/intern/cycles/kernel/closure/bsdf_hair.h b/intern/cycles/kernel/closure/bsdf_hair.h index 5825d2637ba..163e7cc5ee2 100644 --- a/intern/cycles/kernel/closure/bsdf_hair.h +++ b/intern/cycles/kernel/closure/bsdf_hair.h @@ -36,15 +36,15 @@ CCL_NAMESPACE_BEGIN -__device void bsdf_hair_reflection_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_hair_reflection_blur(ShaderClosure *sc, float roughness) { } -__device void bsdf_hair_transmission_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_hair_transmission_blur(ShaderClosure *sc, float roughness) { } -__device int bsdf_hair_reflection_setup(ShaderClosure *sc) +ccl_device int bsdf_hair_reflection_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_HAIR_REFLECTION_ID; sc->data0 = clamp(sc->data0, 0.001f, 1.0f); @@ -52,7 +52,7 @@ __device int bsdf_hair_reflection_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device int bsdf_hair_transmission_setup(ShaderClosure *sc) +ccl_device int bsdf_hair_transmission_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_HAIR_TRANSMISSION_ID; sc->data0 = clamp(sc->data0, 0.001f, 1.0f); @@ -60,7 +60,7 @@ __device int bsdf_hair_transmission_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device float3 bsdf_hair_reflection_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_hair_reflection_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { #ifdef __HAIR__ float offset = sc->offset; @@ -106,18 +106,18 @@ __device float3 bsdf_hair_reflection_eval_reflect(const ShaderClosure *sc, const return make_float3(*pdf, *pdf, *pdf); } -__device float3 bsdf_hair_transmission_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_hair_transmission_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_hair_reflection_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_hair_reflection_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_hair_transmission_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_hair_transmission_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { #ifdef __HAIR__ float offset = sc->offset; @@ -163,7 +163,7 @@ __device float3 bsdf_hair_transmission_eval_transmit(const ShaderClosure *sc, co return make_float3(*pdf, *pdf, *pdf); } -__device int bsdf_hair_reflection_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_hair_reflection_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { #ifdef __HAIR__ float offset = sc->offset; @@ -218,7 +218,7 @@ __device int bsdf_hair_reflection_sample(const ShaderClosure *sc, float3 Ng, flo return LABEL_REFLECT|LABEL_GLOSSY; } -__device int bsdf_hair_transmission_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_hair_transmission_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { #ifdef __HAIR__ float offset = sc->offset; diff --git a/intern/cycles/kernel/closure/bsdf_microfacet.h b/intern/cycles/kernel/closure/bsdf_microfacet.h index b159f585831..737cffb0f18 100644 --- a/intern/cycles/kernel/closure/bsdf_microfacet.h +++ b/intern/cycles/kernel/closure/bsdf_microfacet.h @@ -37,7 +37,7 @@ CCL_NAMESPACE_BEGIN /* GGX */ -__device int bsdf_microfacet_ggx_setup(ShaderClosure *sc) +ccl_device int bsdf_microfacet_ggx_setup(ShaderClosure *sc) { sc->data0 = clamp(sc->data0, 0.0f, 1.0f); /* m_ag */ @@ -46,7 +46,7 @@ __device int bsdf_microfacet_ggx_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device int bsdf_microfacet_ggx_refraction_setup(ShaderClosure *sc) +ccl_device int bsdf_microfacet_ggx_refraction_setup(ShaderClosure *sc) { sc->data0 = clamp(sc->data0, 0.0f, 1.0f); /* m_ag */ @@ -55,12 +55,12 @@ __device int bsdf_microfacet_ggx_refraction_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device void bsdf_microfacet_ggx_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_microfacet_ggx_blur(ShaderClosure *sc, float roughness) { sc->data0 = fmaxf(roughness, sc->data0); /* m_ag */ } -__device float3 bsdf_microfacet_ggx_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_microfacet_ggx_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_ag = max(sc->data0, 1e-4f); int m_refractive = sc->type == CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID; @@ -97,7 +97,7 @@ __device float3 bsdf_microfacet_ggx_eval_reflect(const ShaderClosure *sc, const return make_float3 (0, 0, 0); } -__device float3 bsdf_microfacet_ggx_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_microfacet_ggx_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_ag = max(sc->data0, 1e-4f); float m_eta = sc->data1; @@ -134,7 +134,7 @@ __device float3 bsdf_microfacet_ggx_eval_transmit(const ShaderClosure *sc, const return make_float3 (out, out, out); } -__device int bsdf_microfacet_ggx_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_microfacet_ggx_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_ag = sc->data0; int m_refractive = sc->type == CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID; @@ -255,7 +255,7 @@ __device int bsdf_microfacet_ggx_sample(const ShaderClosure *sc, float3 Ng, floa /* BECKMANN */ -__device int bsdf_microfacet_beckmann_setup(ShaderClosure *sc) +ccl_device int bsdf_microfacet_beckmann_setup(ShaderClosure *sc) { sc->data0 = clamp(sc->data0, 0.0f, 1.0f); /* m_ab */ @@ -263,7 +263,7 @@ __device int bsdf_microfacet_beckmann_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device int bsdf_microfacet_beckmann_refraction_setup(ShaderClosure *sc) +ccl_device int bsdf_microfacet_beckmann_refraction_setup(ShaderClosure *sc) { sc->data0 = clamp(sc->data0, 0.0f, 1.0f); /* m_ab */ @@ -271,12 +271,12 @@ __device int bsdf_microfacet_beckmann_refraction_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device void bsdf_microfacet_beckmann_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_microfacet_beckmann_blur(ShaderClosure *sc, float roughness) { sc->data0 = fmaxf(roughness, sc->data0); /* m_ab */ } -__device float3 bsdf_microfacet_beckmann_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_microfacet_beckmann_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_ab = max(sc->data0, 1e-4f); int m_refractive = sc->type == CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID; @@ -315,7 +315,7 @@ __device float3 bsdf_microfacet_beckmann_eval_reflect(const ShaderClosure *sc, c return make_float3 (0, 0, 0); } -__device float3 bsdf_microfacet_beckmann_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_microfacet_beckmann_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_ab = max(sc->data0, 1e-4f); float m_eta = sc->data1; @@ -354,7 +354,7 @@ __device float3 bsdf_microfacet_beckmann_eval_transmit(const ShaderClosure *sc, return make_float3 (out, out, out); } -__device int bsdf_microfacet_beckmann_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_microfacet_beckmann_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_ab = sc->data0; int m_refractive = sc->type == CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID; diff --git a/intern/cycles/kernel/closure/bsdf_oren_nayar.h b/intern/cycles/kernel/closure/bsdf_oren_nayar.h index c6c6811c007..6f685d5eeea 100644 --- a/intern/cycles/kernel/closure/bsdf_oren_nayar.h +++ b/intern/cycles/kernel/closure/bsdf_oren_nayar.h @@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN -__device float3 bsdf_oren_nayar_get_intensity(const ShaderClosure *sc, float3 n, float3 v, float3 l) +ccl_device float3 bsdf_oren_nayar_get_intensity(const ShaderClosure *sc, float3 n, float3 v, float3 l) { float nl = max(dot(n, l), 0.0f); float nv = max(dot(n, v), 0.0f); @@ -31,7 +31,7 @@ __device float3 bsdf_oren_nayar_get_intensity(const ShaderClosure *sc, float3 n, return make_float3(is, is, is); } -__device int bsdf_oren_nayar_setup(ShaderClosure *sc) +ccl_device int bsdf_oren_nayar_setup(ShaderClosure *sc) { float sigma = sc->data0; @@ -47,11 +47,11 @@ __device int bsdf_oren_nayar_setup(ShaderClosure *sc) return SD_BSDF | SD_BSDF_HAS_EVAL; } -__device void bsdf_oren_nayar_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_oren_nayar_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_oren_nayar_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_oren_nayar_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { if (dot(sc->N, omega_in) > 0.0f) { *pdf = 0.5f * M_1_PI_F; @@ -63,12 +63,12 @@ __device float3 bsdf_oren_nayar_eval_reflect(const ShaderClosure *sc, const floa } } -__device float3 bsdf_oren_nayar_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_oren_nayar_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_oren_nayar_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_oren_nayar_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { sample_uniform_hemisphere(sc->N, randu, randv, omega_in, pdf); diff --git a/intern/cycles/kernel/closure/bsdf_phong_ramp.h b/intern/cycles/kernel/closure/bsdf_phong_ramp.h index 1e332933287..219c5aea159 100644 --- a/intern/cycles/kernel/closure/bsdf_phong_ramp.h +++ b/intern/cycles/kernel/closure/bsdf_phong_ramp.h @@ -35,7 +35,7 @@ CCL_NAMESPACE_BEGIN -__device float3 bsdf_phong_ramp_get_color(const ShaderClosure *sc, const float3 colors[8], float pos) +ccl_device float3 bsdf_phong_ramp_get_color(const ShaderClosure *sc, const float3 colors[8], float pos) { int MAXCOLORS = 8; @@ -49,7 +49,7 @@ __device float3 bsdf_phong_ramp_get_color(const ShaderClosure *sc, const float3 return colors[ipos] * (1.0f - offset) + colors[ipos+1] * offset; } -__device int bsdf_phong_ramp_setup(ShaderClosure *sc) +ccl_device int bsdf_phong_ramp_setup(ShaderClosure *sc) { sc->data0 = max(sc->data0, 0.0f); @@ -57,11 +57,11 @@ __device int bsdf_phong_ramp_setup(ShaderClosure *sc) return SD_BSDF | SD_BSDF_HAS_EVAL | SD_BSDF_GLOSSY; } -__device void bsdf_phong_ramp_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_phong_ramp_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_phong_ramp_eval_reflect(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_phong_ramp_eval_reflect(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) { float m_exponent = sc->data0; float cosNI = dot(sc->N, omega_in); @@ -83,12 +83,12 @@ __device float3 bsdf_phong_ramp_eval_reflect(const ShaderClosure *sc, const floa return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_phong_ramp_eval_transmit(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_phong_ramp_eval_transmit(const ShaderClosure *sc, const float3 colors[8], const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_phong_ramp_sample(const ShaderClosure *sc, const float3 colors[8], float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_phong_ramp_sample(const ShaderClosure *sc, const float3 colors[8], float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float cosNO = dot(sc->N, I); float m_exponent = sc->data0; diff --git a/intern/cycles/kernel/closure/bsdf_reflection.h b/intern/cycles/kernel/closure/bsdf_reflection.h index 7715aac936f..0baccdf155c 100644 --- a/intern/cycles/kernel/closure/bsdf_reflection.h +++ b/intern/cycles/kernel/closure/bsdf_reflection.h @@ -37,27 +37,27 @@ CCL_NAMESPACE_BEGIN /* REFLECTION */ -__device int bsdf_reflection_setup(ShaderClosure *sc) +ccl_device int bsdf_reflection_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_REFLECTION_ID; return SD_BSDF; } -__device void bsdf_reflection_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_reflection_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_reflection_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_reflection_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_reflection_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_reflection_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_reflection_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_reflection_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float3 N = sc->N; diff --git a/intern/cycles/kernel/closure/bsdf_refraction.h b/intern/cycles/kernel/closure/bsdf_refraction.h index 8565c99d04e..c4698b42060 100644 --- a/intern/cycles/kernel/closure/bsdf_refraction.h +++ b/intern/cycles/kernel/closure/bsdf_refraction.h @@ -37,27 +37,27 @@ CCL_NAMESPACE_BEGIN /* REFRACTION */ -__device int bsdf_refraction_setup(ShaderClosure *sc) +ccl_device int bsdf_refraction_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_REFRACTION_ID; return SD_BSDF; } -__device void bsdf_refraction_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_refraction_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_refraction_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_refraction_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_refraction_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_refraction_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_refraction_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_refraction_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_eta = sc->data0; float3 N = sc->N; diff --git a/intern/cycles/kernel/closure/bsdf_toon.h b/intern/cycles/kernel/closure/bsdf_toon.h index e69981dba77..797fa4227ae 100644 --- a/intern/cycles/kernel/closure/bsdf_toon.h +++ b/intern/cycles/kernel/closure/bsdf_toon.h @@ -37,7 +37,7 @@ CCL_NAMESPACE_BEGIN /* DIFFUSE TOON */ -__device int bsdf_diffuse_toon_setup(ShaderClosure *sc) +ccl_device int bsdf_diffuse_toon_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_DIFFUSE_TOON_ID; sc->data0 = clamp(sc->data0, 0.0f, 1.0f); @@ -46,11 +46,11 @@ __device int bsdf_diffuse_toon_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL; } -__device void bsdf_diffuse_toon_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_diffuse_toon_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_toon_get_intensity(float max_angle, float smooth, float angle) +ccl_device float3 bsdf_toon_get_intensity(float max_angle, float smooth, float angle) { float is; @@ -64,12 +64,12 @@ __device float3 bsdf_toon_get_intensity(float max_angle, float smooth, float ang return make_float3(is, is, is); } -__device float bsdf_toon_get_sample_angle(float max_angle, float smooth) +ccl_device float bsdf_toon_get_sample_angle(float max_angle, float smooth) { return fminf(max_angle + smooth, M_PI_2_F); } -__device float3 bsdf_diffuse_toon_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_diffuse_toon_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float max_angle = sc->data0*M_PI_2_F; float smooth = sc->data1*M_PI_2_F; @@ -87,12 +87,12 @@ __device float3 bsdf_diffuse_toon_eval_reflect(const ShaderClosure *sc, const fl return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_diffuse_toon_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_diffuse_toon_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_diffuse_toon_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_diffuse_toon_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float max_angle = sc->data0*M_PI_2_F; float smooth = sc->data1*M_PI_2_F; @@ -121,7 +121,7 @@ __device int bsdf_diffuse_toon_sample(const ShaderClosure *sc, float3 Ng, float3 /* GLOSSY TOON */ -__device int bsdf_glossy_toon_setup(ShaderClosure *sc) +ccl_device int bsdf_glossy_toon_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_GLOSSY_TOON_ID; sc->data0 = clamp(sc->data0, 0.0f, 1.0f); @@ -130,11 +130,11 @@ __device int bsdf_glossy_toon_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL; } -__device void bsdf_glossy_toon_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_glossy_toon_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_glossy_toon_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_glossy_toon_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float max_angle = sc->data0*M_PI_2_F; float smooth = sc->data1*M_PI_2_F; @@ -158,12 +158,12 @@ __device float3 bsdf_glossy_toon_eval_reflect(const ShaderClosure *sc, const flo return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_glossy_toon_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_glossy_toon_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_glossy_toon_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_glossy_toon_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float max_angle = sc->data0*M_PI_2_F; float smooth = sc->data1*M_PI_2_F; diff --git a/intern/cycles/kernel/closure/bsdf_transparent.h b/intern/cycles/kernel/closure/bsdf_transparent.h index 81bc7690b50..e62aecf3da6 100644 --- a/intern/cycles/kernel/closure/bsdf_transparent.h +++ b/intern/cycles/kernel/closure/bsdf_transparent.h @@ -35,27 +35,27 @@ CCL_NAMESPACE_BEGIN -__device int bsdf_transparent_setup(ShaderClosure *sc) +ccl_device int bsdf_transparent_setup(ShaderClosure *sc) { sc->type = CLOSURE_BSDF_TRANSPARENT_ID; return SD_BSDF; } -__device void bsdf_transparent_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_transparent_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_transparent_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_transparent_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device float3 bsdf_transparent_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_transparent_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_transparent_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_transparent_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { // only one direction is possible *omega_in = -I; diff --git a/intern/cycles/kernel/closure/bsdf_util.h b/intern/cycles/kernel/closure/bsdf_util.h index 5ac26055e8d..f6dceb3ca82 100644 --- a/intern/cycles/kernel/closure/bsdf_util.h +++ b/intern/cycles/kernel/closure/bsdf_util.h @@ -35,7 +35,7 @@ CCL_NAMESPACE_BEGIN -__device float fresnel_dielectric(float eta, const float3 N, +ccl_device float fresnel_dielectric(float eta, const float3 N, const float3 I, float3 *R, float3 *T, #ifdef __RAY_DIFFERENTIALS__ const float3 dIdx, const float3 dIdy, @@ -95,7 +95,7 @@ __device float fresnel_dielectric(float eta, const float3 N, } } -__device float fresnel_dielectric_cos(float cosi, float eta) +ccl_device float fresnel_dielectric_cos(float cosi, float eta) { // compute fresnel reflectance without explicitly computing // the refracted direction @@ -110,7 +110,7 @@ __device float fresnel_dielectric_cos(float cosi, float eta) return 1.0f; // TIR(no refracted component) } -__device float fresnel_conductor(float cosi, float eta, float k) +ccl_device float fresnel_conductor(float cosi, float eta, float k) { float tmp_f = eta * eta + k * k; float tmp = tmp_f * cosi * cosi; @@ -121,7 +121,7 @@ __device float fresnel_conductor(float cosi, float eta, float k) return(Rparl2 + Rperp2) * 0.5f; } -__device float smooth_step(float edge0, float edge1, float x) +ccl_device float smooth_step(float edge0, float edge1, float x) { float result; if(x < edge0) result = 0.0f; diff --git a/intern/cycles/kernel/closure/bsdf_ward.h b/intern/cycles/kernel/closure/bsdf_ward.h index 0e5b0c544c7..c9de615a011 100644 --- a/intern/cycles/kernel/closure/bsdf_ward.h +++ b/intern/cycles/kernel/closure/bsdf_ward.h @@ -37,7 +37,7 @@ CCL_NAMESPACE_BEGIN /* WARD */ -__device int bsdf_ward_setup(ShaderClosure *sc) +ccl_device int bsdf_ward_setup(ShaderClosure *sc) { sc->data0 = clamp(sc->data0, 1e-4f, 1.0f); /* m_ax */ sc->data1 = clamp(sc->data1, 1e-4f, 1.0f); /* m_ay */ @@ -46,13 +46,13 @@ __device int bsdf_ward_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device void bsdf_ward_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_ward_blur(ShaderClosure *sc, float roughness) { sc->data0 = fmaxf(roughness, sc->data0); /* m_ax */ sc->data1 = fmaxf(roughness, sc->data1); /* m_ay */ } -__device float3 bsdf_ward_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_ward_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_ax = sc->data0; float m_ay = sc->data1; @@ -87,12 +87,12 @@ __device float3 bsdf_ward_eval_reflect(const ShaderClosure *sc, const float3 I, return make_float3 (0, 0, 0); } -__device float3 bsdf_ward_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_ward_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_ward_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_ward_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_ax = sc->data0; float m_ay = sc->data1; diff --git a/intern/cycles/kernel/closure/bsdf_westin.h b/intern/cycles/kernel/closure/bsdf_westin.h index e1a6b031d5e..ca4c05e91fe 100644 --- a/intern/cycles/kernel/closure/bsdf_westin.h +++ b/intern/cycles/kernel/closure/bsdf_westin.h @@ -37,7 +37,7 @@ CCL_NAMESPACE_BEGIN /* WESTIN BACKSCATTER */ -__device int bsdf_westin_backscatter_setup(ShaderClosure *sc) +ccl_device int bsdf_westin_backscatter_setup(ShaderClosure *sc) { float roughness = sc->data0; roughness = clamp(roughness, 1e-5f, 1.0f); @@ -49,14 +49,14 @@ __device int bsdf_westin_backscatter_setup(ShaderClosure *sc) return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device void bsdf_westin_backscatter_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_westin_backscatter_blur(ShaderClosure *sc, float roughness) { float m_invroughness = sc->data0; m_invroughness = min(1.0f/roughness, m_invroughness); sc->data0 = m_invroughness; } -__device float3 bsdf_westin_backscatter_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_westin_backscatter_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_invroughness = sc->data0; float3 N = sc->N; @@ -73,12 +73,12 @@ __device float3 bsdf_westin_backscatter_eval_reflect(const ShaderClosure *sc, co return make_float3 (0, 0, 0); } -__device float3 bsdf_westin_backscatter_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_westin_backscatter_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_westin_backscatter_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_westin_backscatter_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_invroughness = sc->data0; float3 N = sc->N; @@ -116,18 +116,18 @@ __device int bsdf_westin_backscatter_sample(const ShaderClosure *sc, float3 Ng, /* WESTIN SHEEN */ -__device int bsdf_westin_sheen_setup(ShaderClosure *sc) +ccl_device int bsdf_westin_sheen_setup(ShaderClosure *sc) { /* float edginess = sc->data0; */ sc->type = CLOSURE_BSDF_WESTIN_SHEEN_ID; return SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY; } -__device void bsdf_westin_sheen_blur(ShaderClosure *sc, float roughness) +ccl_device void bsdf_westin_sheen_blur(ShaderClosure *sc, float roughness) { } -__device float3 bsdf_westin_sheen_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_westin_sheen_eval_reflect(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { float m_edginess = sc->data0; float3 N = sc->N; @@ -144,12 +144,12 @@ __device float3 bsdf_westin_sheen_eval_reflect(const ShaderClosure *sc, const fl return make_float3 (0, 0, 0); } -__device float3 bsdf_westin_sheen_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) +ccl_device float3 bsdf_westin_sheen_eval_transmit(const ShaderClosure *sc, const float3 I, const float3 omega_in, float *pdf) { return make_float3(0.0f, 0.0f, 0.0f); } -__device int bsdf_westin_sheen_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) +ccl_device int bsdf_westin_sheen_sample(const ShaderClosure *sc, float3 Ng, float3 I, float3 dIdx, float3 dIdy, float randu, float randv, float3 *eval, float3 *omega_in, float3 *domega_in_dx, float3 *domega_in_dy, float *pdf) { float m_edginess = sc->data0; float3 N = sc->N; diff --git a/intern/cycles/kernel/closure/bssrdf.h b/intern/cycles/kernel/closure/bssrdf.h index 4ceff655dd5..3849dedc3b6 100644 --- a/intern/cycles/kernel/closure/bssrdf.h +++ b/intern/cycles/kernel/closure/bssrdf.h @@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN -__device int bssrdf_setup(ShaderClosure *sc, ClosureType type) +ccl_device int bssrdf_setup(ShaderClosure *sc, ClosureType type) { if(sc->data0 < BSSRDF_MIN_RADIUS) { /* revert to diffuse BSDF if radius too small */ @@ -47,7 +47,7 @@ __device int bssrdf_setup(ShaderClosure *sc, ClosureType type) /* paper suggests 1/12.46 which is much too small, suspect it's *12.46 */ #define GAUSS_TRUNCATE 12.46f -__device float bssrdf_gaussian_eval(ShaderClosure *sc, float r) +ccl_device float bssrdf_gaussian_eval(ShaderClosure *sc, float r) { /* integrate (2*pi*r * exp(-r*r/(2*v)))/(2*pi*v)) from 0 to Rm * = 1 - exp(-Rm*Rm/(2*v)) */ @@ -60,7 +60,7 @@ __device float bssrdf_gaussian_eval(ShaderClosure *sc, float r) return expf(-r*r/(2.0f*v))/(2.0f*M_PI_F*v); } -__device float bssrdf_gaussian_pdf(ShaderClosure *sc, float r) +ccl_device float bssrdf_gaussian_pdf(ShaderClosure *sc, float r) { /* 1.0 - expf(-Rm*Rm/(2*v)) simplified */ const float area_truncated = 1.0f - expf(-0.5f*GAUSS_TRUNCATE); @@ -68,7 +68,7 @@ __device float bssrdf_gaussian_pdf(ShaderClosure *sc, float r) return bssrdf_gaussian_eval(sc, r) * (1.0f/(area_truncated)); } -__device void bssrdf_gaussian_sample(ShaderClosure *sc, float xi, float *r, float *h) +ccl_device void bssrdf_gaussian_sample(ShaderClosure *sc, float xi, float *r, float *h) { /* xi = integrate (2*pi*r * exp(-r*r/(2*v)))/(2*pi*v)) = -exp(-r^2/(2*v)) * r = sqrt(-2*v*logf(xi)) */ @@ -94,7 +94,7 @@ __device void bssrdf_gaussian_sample(ShaderClosure *sc, float xi, float *r, floa * far as I can tell has no closed form solution. So we get an iterative solution * instead with newton-raphson. */ -__device float bssrdf_cubic_eval(ShaderClosure *sc, float r) +ccl_device float bssrdf_cubic_eval(ShaderClosure *sc, float r) { const float sharpness = sc->T.x; @@ -141,13 +141,13 @@ __device float bssrdf_cubic_eval(ShaderClosure *sc, float r) } } -__device float bssrdf_cubic_pdf(ShaderClosure *sc, float r) +ccl_device float bssrdf_cubic_pdf(ShaderClosure *sc, float r) { return bssrdf_cubic_eval(sc, r); } /* solve 10x^2 - 20x^3 + 15x^4 - 4x^5 - xi == 0 */ -__device float bssrdf_cubic_quintic_root_find(float xi) +ccl_device float bssrdf_cubic_quintic_root_find(float xi) { /* newton-raphson iteration, usually succeeds in 2-4 iterations, except * outside 0.02 ... 0.98 where it can go up to 10, so overall performance @@ -174,7 +174,7 @@ __device float bssrdf_cubic_quintic_root_find(float xi) return x; } -__device void bssrdf_cubic_sample(ShaderClosure *sc, float xi, float *r, float *h) +ccl_device void bssrdf_cubic_sample(ShaderClosure *sc, float xi, float *r, float *h) { float Rm = sc->data0; float r_ = bssrdf_cubic_quintic_root_find(xi); @@ -196,13 +196,13 @@ __device void bssrdf_cubic_sample(ShaderClosure *sc, float xi, float *r, float * * * Samples distributed over disk with no falloff, for reference. */ -__device float bssrdf_none_eval(ShaderClosure *sc, float r) +ccl_device float bssrdf_none_eval(ShaderClosure *sc, float r) { const float Rm = sc->data0; return (r < Rm)? 1.0f: 0.0f; } -__device float bssrdf_none_pdf(ShaderClosure *sc, float r) +ccl_device float bssrdf_none_pdf(ShaderClosure *sc, float r) { /* integrate (2*pi*r)/(pi*Rm*Rm) from 0 to Rm = 1 */ const float Rm = sc->data0; @@ -211,7 +211,7 @@ __device float bssrdf_none_pdf(ShaderClosure *sc, float r) return bssrdf_none_eval(sc, r) / area; } -__device void bssrdf_none_sample(ShaderClosure *sc, float xi, float *r, float *h) +ccl_device void bssrdf_none_sample(ShaderClosure *sc, float xi, float *r, float *h) { /* xi = integrate (2*pi*r)/(pi*Rm*Rm) = r^2/Rm^2 * r = sqrt(xi)*Rm */ @@ -226,7 +226,7 @@ __device void bssrdf_none_sample(ShaderClosure *sc, float xi, float *r, float *h /* Generic */ -__device void bssrdf_sample(ShaderClosure *sc, float xi, float *r, float *h) +ccl_device void bssrdf_sample(ShaderClosure *sc, float xi, float *r, float *h) { if(sc->type == CLOSURE_BSSRDF_CUBIC_ID) bssrdf_cubic_sample(sc, xi, r, h); @@ -234,7 +234,7 @@ __device void bssrdf_sample(ShaderClosure *sc, float xi, float *r, float *h) bssrdf_gaussian_sample(sc, xi, r, h); } -__device float bssrdf_pdf(ShaderClosure *sc, float r) +ccl_device float bssrdf_pdf(ShaderClosure *sc, float r) { if(sc->type == CLOSURE_BSSRDF_CUBIC_ID) return bssrdf_cubic_pdf(sc, r); diff --git a/intern/cycles/kernel/closure/emissive.h b/intern/cycles/kernel/closure/emissive.h index 33b1b695a9a..c534df373bd 100644 --- a/intern/cycles/kernel/closure/emissive.h +++ b/intern/cycles/kernel/closure/emissive.h @@ -37,19 +37,19 @@ CCL_NAMESPACE_BEGIN /* return the probability distribution function in the direction I, * given the parameters and the light's surface normal. This MUST match * the PDF computed by sample(). */ -__device float emissive_pdf(const float3 Ng, const float3 I) +ccl_device float emissive_pdf(const float3 Ng, const float3 I) { float cosNO = fabsf(dot(Ng, I)); return (cosNO > 0.0f)? 1.0f: 0.0f; } -__device void emissive_sample(const float3 Ng, float randu, float randv, +ccl_device void emissive_sample(const float3 Ng, float randu, float randv, float3 *omega_out, float *pdf) { /* todo: not implemented and used yet */ } -__device float3 emissive_simple_eval(const float3 Ng, const float3 I) +ccl_device float3 emissive_simple_eval(const float3 Ng, const float3 I) { float res = emissive_pdf(Ng, I); diff --git a/intern/cycles/kernel/closure/volume.h b/intern/cycles/kernel/closure/volume.h index ddaf939984e..f30b30c8c76 100644 --- a/intern/cycles/kernel/closure/volume.h +++ b/intern/cycles/kernel/closure/volume.h @@ -21,7 +21,7 @@ CCL_NAMESPACE_BEGIN /* ISOTROPIC VOLUME CLOSURE */ -__device int volume_isotropic_setup(ShaderClosure *sc, float density) +ccl_device int volume_isotropic_setup(ShaderClosure *sc, float density) { sc->type = CLOSURE_VOLUME_ISOTROPIC_ID; sc->data0 = density; @@ -29,14 +29,14 @@ __device int volume_isotropic_setup(ShaderClosure *sc, float density) return SD_VOLUME; } -__device float3 volume_isotropic_eval_phase(const ShaderClosure *sc, const float3 omega_in, const float3 omega_out) +ccl_device float3 volume_isotropic_eval_phase(const ShaderClosure *sc, const float3 omega_in, const float3 omega_out) { return make_float3(1.0f, 1.0f, 1.0f); } /* TRANSPARENT VOLUME CLOSURE */ -__device int volume_transparent_setup(ShaderClosure *sc, float density) +ccl_device int volume_transparent_setup(ShaderClosure *sc, float density) { sc->type = CLOSURE_VOLUME_TRANSPARENT_ID; sc->data0 = density; @@ -44,14 +44,14 @@ __device int volume_transparent_setup(ShaderClosure *sc, float density) return SD_VOLUME; } -__device float3 volume_transparent_eval_phase(const ShaderClosure *sc, const float3 omega_in, const float3 omega_out) +ccl_device float3 volume_transparent_eval_phase(const ShaderClosure *sc, const float3 omega_in, const float3 omega_out) { return make_float3(1.0f, 1.0f, 1.0f); } /* VOLUME CLOSURE */ -__device float3 volume_eval_phase(KernelGlobals *kg, const ShaderClosure *sc, const float3 omega_in, const float3 omega_out) +ccl_device float3 volume_eval_phase(KernelGlobals *kg, const ShaderClosure *sc, const float3 omega_in, const float3 omega_out) { #ifdef __OSL__ if(kg->osl && sc->prim) diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index 28e72d78731..6988ad6027f 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -26,12 +26,12 @@ #include "kernel_displace.h" __kernel void kernel_ocl_path_trace( - __constant KernelData *data, - __global float *buffer, - __global uint *rng_state, + ccl_constant KernelData *data, + ccl_global float *buffer, + ccl_global uint *rng_state, #define KERNEL_TEX(type, ttype, name) \ - __global type *name, + ccl_global type *name, #include "kernel_textures.h" int sample, @@ -53,12 +53,12 @@ __kernel void kernel_ocl_path_trace( } __kernel void kernel_ocl_convert_to_byte( - __constant KernelData *data, - __global uchar4 *rgba, - __global float *buffer, + ccl_constant KernelData *data, + ccl_global uchar4 *rgba, + ccl_global float *buffer, #define KERNEL_TEX(type, ttype, name) \ - __global type *name, + ccl_global type *name, #include "kernel_textures.h" float sample_scale, @@ -80,12 +80,12 @@ __kernel void kernel_ocl_convert_to_byte( } __kernel void kernel_ocl_convert_to_half_float( - __constant KernelData *data, - __global uchar4 *rgba, - __global float *buffer, + ccl_constant KernelData *data, + ccl_global uchar4 *rgba, + ccl_global float *buffer, #define KERNEL_TEX(type, ttype, name) \ - __global type *name, + ccl_global type *name, #include "kernel_textures.h" float sample_scale, @@ -107,12 +107,12 @@ __kernel void kernel_ocl_convert_to_half_float( } __kernel void kernel_ocl_shader( - __constant KernelData *data, - __global uint4 *input, - __global float4 *output, + ccl_constant KernelData *data, + ccl_global uint4 *input, + ccl_global float4 *output, #define KERNEL_TEX(type, ttype, name) \ - __global type *name, + ccl_global type *name, #include "kernel_textures.h" int type, int sx, int sw) diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h index d7531be0d8a..f4febd7cf2c 100644 --- a/intern/cycles/kernel/kernel_accumulate.h +++ b/intern/cycles/kernel/kernel_accumulate.h @@ -21,7 +21,7 @@ CCL_NAMESPACE_BEGIN * BSDF evaluation result, split per BSDF type. This is used to accumulate * render passes separately. */ -__device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass) +ccl_device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass) { #ifdef __PASSES__ eval->use_light_pass = use_light_pass; @@ -51,7 +51,7 @@ __device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 val #endif } -__device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value) +ccl_device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value) { #ifdef __PASSES__ if(eval->use_light_pass) { @@ -73,7 +73,7 @@ __device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 va #endif } -__device_inline bool bsdf_eval_is_zero(BsdfEval *eval) +ccl_device_inline bool bsdf_eval_is_zero(BsdfEval *eval) { #ifdef __PASSES__ if(eval->use_light_pass) { @@ -90,7 +90,7 @@ __device_inline bool bsdf_eval_is_zero(BsdfEval *eval) #endif } -__device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value) +ccl_device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value) { #ifdef __PASSES__ if(eval->use_light_pass) { @@ -115,7 +115,7 @@ __device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value) * visible as the first non-transparent hit, while indirectly visible are the * bounces after that. */ -__device_inline void path_radiance_init(PathRadiance *L, int use_light_pass) +ccl_device_inline void path_radiance_init(PathRadiance *L, int use_light_pass) { /* clear all */ #ifdef __PASSES__ @@ -159,7 +159,7 @@ __device_inline void path_radiance_init(PathRadiance *L, int use_light_pass) #endif } -__device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughput, +ccl_device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughput, BsdfEval *bsdf_eval, float bsdf_pdf, int bounce, int bsdf_label) { float inverse_pdf = 1.0f/bsdf_pdf; @@ -192,7 +192,7 @@ __device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughp #endif } -__device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce) +ccl_device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce) { #ifdef __PASSES__ if(L->use_light_pass) { @@ -210,7 +210,7 @@ __device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throug #endif } -__device_inline void path_radiance_accum_ao(PathRadiance *L, float3 throughput, float3 alpha, float3 bsdf, float3 ao, int bounce) +ccl_device_inline void path_radiance_accum_ao(PathRadiance *L, float3 throughput, float3 alpha, float3 bsdf, float3 ao, int bounce) { #ifdef __PASSES__ if(L->use_light_pass) { @@ -231,7 +231,7 @@ __device_inline void path_radiance_accum_ao(PathRadiance *L, float3 throughput, #endif } -__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, float3 shadow, float shadow_fac, int bounce, bool is_lamp) +ccl_device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, float3 shadow, float shadow_fac, int bounce, bool is_lamp) { #ifdef __PASSES__ if(L->use_light_pass) { @@ -261,7 +261,7 @@ __device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughpu #endif } -__device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce) +ccl_device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce) { #ifdef __PASSES__ if(L->use_light_pass) { @@ -279,7 +279,7 @@ __device_inline void path_radiance_accum_background(PathRadiance *L, float3 thro #endif } -__device_inline void path_radiance_sum_indirect(PathRadiance *L) +ccl_device_inline void path_radiance_sum_indirect(PathRadiance *L) { #ifdef __PASSES__ /* this division is a bit ugly, but means we only have to keep track of @@ -301,7 +301,7 @@ __device_inline void path_radiance_sum_indirect(PathRadiance *L) #endif } -__device_inline void path_radiance_reset_indirect(PathRadiance *L) +ccl_device_inline void path_radiance_reset_indirect(PathRadiance *L) { #ifdef __PASSES__ if(L->use_light_pass) { @@ -316,7 +316,7 @@ __device_inline void path_radiance_reset_indirect(PathRadiance *L) #endif } -__device_inline float3 path_radiance_sum(KernelGlobals *kg, PathRadiance *L) +ccl_device_inline float3 path_radiance_sum(KernelGlobals *kg, PathRadiance *L) { #ifdef __PASSES__ if(L->use_light_pass) { @@ -338,7 +338,7 @@ __device_inline float3 path_radiance_sum(KernelGlobals *kg, PathRadiance *L) #endif } -__device_inline void path_radiance_clamp(PathRadiance *L, float3 *L_sum, float clamp) +ccl_device_inline void path_radiance_clamp(PathRadiance *L, float3 *L_sum, float clamp) { float sum = fabsf((*L_sum).x) + fabsf((*L_sum).y) + fabsf((*L_sum).z); diff --git a/intern/cycles/kernel/kernel_bvh.h b/intern/cycles/kernel/kernel_bvh.h index 44a9822c103..5aae4111fca 100644 --- a/intern/cycles/kernel/kernel_bvh.h +++ b/intern/cycles/kernel/kernel_bvh.h @@ -42,7 +42,7 @@ CCL_NAMESPACE_BEGIN #define NO_EXTENDED_PRECISION volatile #endif -__device_inline float3 bvh_inverse_direction(float3 dir) +ccl_device_inline float3 bvh_inverse_direction(float3 dir) { /* avoid divide by zero (ooeps = exp2f(-80.0f)) */ float ooeps = 0.00000000000000000000000082718061255302767487140869206996285356581211090087890625f; @@ -55,7 +55,7 @@ __device_inline float3 bvh_inverse_direction(float3 dir) return idir; } -__device_inline void bvh_instance_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, const float tmax) +ccl_device_inline void bvh_instance_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, const float tmax) { Transform tfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM); @@ -72,7 +72,7 @@ __device_inline void bvh_instance_push(KernelGlobals *kg, int object, const Ray *t *= len; } -__device_inline void bvh_instance_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, const float tmax) +ccl_device_inline void bvh_instance_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, const float tmax) { if(*t != FLT_MAX) { Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM); @@ -84,7 +84,7 @@ __device_inline void bvh_instance_pop(KernelGlobals *kg, int object, const Ray * } #ifdef __OBJECT_MOTION__ -__device_inline void bvh_instance_motion_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, Transform *tfm, const float tmax) +ccl_device_inline void bvh_instance_motion_push(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, Transform *tfm, const float tmax) { Transform itfm; *tfm = object_fetch_transform_motion_test(kg, object, ray->time, &itfm); @@ -102,7 +102,7 @@ __device_inline void bvh_instance_motion_push(KernelGlobals *kg, int object, con *t *= len; } -__device_inline void bvh_instance_motion_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, Transform *tfm, const float tmax) +ccl_device_inline void bvh_instance_motion_pop(KernelGlobals *kg, int object, const Ray *ray, float3 *P, float3 *idir, float *t, Transform *tfm, const float tmax) { if(*t != FLT_MAX) *t *= len(transform_direction(tfm, 1.0f/(*idir))); @@ -113,7 +113,7 @@ __device_inline void bvh_instance_motion_pop(KernelGlobals *kg, int object, cons #endif /* Sven Woop's algorithm */ -__device_inline bool bvh_triangle_intersect(KernelGlobals *kg, Intersection *isect, +ccl_device_inline bool bvh_triangle_intersect(KernelGlobals *kg, Intersection *isect, float3 P, float3 idir, uint visibility, int object, int triAddr) { /* compute and check intersection t-value */ @@ -161,7 +161,7 @@ __device_inline bool bvh_triangle_intersect(KernelGlobals *kg, Intersection *ise } #ifdef __HAIR__ -__device_inline void curvebounds(float *lower, float *upper, float *extremta, float *extrema, float *extremtb, float *extremb, float p0, float p1, float p2, float p3) +ccl_device_inline void curvebounds(float *lower, float *upper, float *extremta, float *extrema, float *extremtb, float *extremb, float p0, float p1, float p2, float p3) { float halfdiscroot = (p2 * p2 - 3 * p3 * p1); float ta = -1.0f; @@ -211,7 +211,7 @@ __device_inline void curvebounds(float *lower, float *upper, float *extremta, fl } } -__device_inline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Intersection *isect, +ccl_device_inline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Intersection *isect, float3 P, float3 idir, uint visibility, int object, int curveAddr, int segment, uint *lcg_state, float difl, float extmax) { float epsilon = 0.0f; @@ -520,7 +520,7 @@ __device_inline bool bvh_cardinal_curve_intersect(KernelGlobals *kg, Intersectio return hit; } -__device_inline bool bvh_curve_intersect(KernelGlobals *kg, Intersection *isect, +ccl_device_inline bool bvh_curve_intersect(KernelGlobals *kg, Intersection *isect, float3 P, float3 idir, uint visibility, int object, int curveAddr, int segment, uint *lcg_state, float difl, float extmax) { /* curve Intersection check */ @@ -689,7 +689,7 @@ __device_inline bool bvh_curve_intersect(KernelGlobals *kg, Intersection *isect, * only want to intersect with primitives in the same object, and if case of * multiple hits we pick a single random primitive as the intersection point. */ -__device_inline void bvh_triangle_intersect_subsurface(KernelGlobals *kg, Intersection *isect_array, +ccl_device_inline void bvh_triangle_intersect_subsurface(KernelGlobals *kg, Intersection *isect_array, float3 P, float3 idir, int object, int triAddr, float tmax, uint *num_hits, uint *lcg_state, int max_hits) { /* compute and check intersection t-value */ @@ -811,9 +811,9 @@ __device_inline void bvh_triangle_intersect_subsurface(KernelGlobals *kg, Inters /* to work around titan bug when using arrays instead of textures */ #if !defined(__KERNEL_CUDA__) || defined(__KERNEL_CUDA_TEX_STORAGE__) -__device_inline +ccl_device_inline #else -__device_noinline +ccl_device_noinline #endif #ifdef __HAIR__ bool scene_intersect(KernelGlobals *kg, const Ray *ray, const uint visibility, Intersection *isect, uint *lcg_state, float difl, float extmax) @@ -859,9 +859,9 @@ bool scene_intersect(KernelGlobals *kg, const Ray *ray, const uint visibility, I /* to work around titan bug when using arrays instead of textures */ #ifdef __SUBSURFACE__ #if !defined(__KERNEL_CUDA__) || defined(__KERNEL_CUDA_TEX_STORAGE__) -__device_inline +ccl_device_inline #else -__device_noinline +ccl_device_noinline #endif uint scene_intersect_subsurface(KernelGlobals *kg, const Ray *ray, Intersection *isect, int subsurface_object, uint *lcg_state, int max_hits) { @@ -903,7 +903,7 @@ uint scene_intersect_subsurface(KernelGlobals *kg, const Ray *ray, Intersection /* Ray offset to avoid self intersection */ -__device_inline float3 ray_offset(float3 P, float3 Ng) +ccl_device_inline float3 ray_offset(float3 P, float3 Ng) { #ifdef __INTERSECTION_REFINE__ const float epsilon_f = 1e-5f; @@ -955,7 +955,7 @@ __device_inline float3 ray_offset(float3 P, float3 Ng) * far the precision is often not so good, this reintersects the primitive from * a closer distance. */ -__device_inline float3 bvh_triangle_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray) +ccl_device_inline float3 bvh_triangle_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray) { float3 P = ray->P; float3 D = ray->D; @@ -1000,7 +1000,7 @@ __device_inline float3 bvh_triangle_refine(KernelGlobals *kg, ShaderData *sd, co } /* same as above, except that isect->t is assumed to be in object space for instancing */ -__device_inline float3 bvh_triangle_refine_subsurface(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray) +ccl_device_inline float3 bvh_triangle_refine_subsurface(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray) { float3 P = ray->P; float3 D = ray->D; @@ -1046,7 +1046,7 @@ __device_inline float3 bvh_triangle_refine_subsurface(KernelGlobals *kg, ShaderD #ifdef __HAIR__ -__device_inline float3 curvetangent(float t, float3 p0, float3 p1, float3 p2, float3 p3) +ccl_device_inline float3 curvetangent(float t, float3 p0, float3 p1, float3 p2, float3 p3) { float fc = 0.71f; float data[4]; @@ -1058,7 +1058,7 @@ __device_inline float3 curvetangent(float t, float3 p0, float3 p1, float3 p2, fl return data[0] * p0 + data[1] * p1 + data[2] * p2 + data[3] * p3; } -__device_inline float3 curvepoint(float t, float3 p0, float3 p1, float3 p2, float3 p3) +ccl_device_inline float3 curvepoint(float t, float3 p0, float3 p1, float3 p2, float3 p3) { float data[4]; float fc = 0.71f; @@ -1071,7 +1071,7 @@ __device_inline float3 curvepoint(float t, float3 p0, float3 p1, float3 p2, floa return data[0] * p0 + data[1] * p1 + data[2] * p2 + data[3] * p3; } -__device_inline float3 bvh_curve_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, float t) +ccl_device_inline float3 bvh_curve_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, float t) { int flag = kernel_data.curve.curveflags; float3 P = ray->P; diff --git a/intern/cycles/kernel/kernel_bvh_subsurface.h b/intern/cycles/kernel/kernel_bvh_subsurface.h index 4446c1821d5..fb41bdcfa37 100644 --- a/intern/cycles/kernel/kernel_bvh_subsurface.h +++ b/intern/cycles/kernel/kernel_bvh_subsurface.h @@ -28,7 +28,7 @@ #define FEATURE(f) (((BVH_FUNCTION_FEATURES) & (f)) != 0) -__device uint BVH_FUNCTION_NAME(KernelGlobals *kg, const Ray *ray, Intersection *isect_array, +ccl_device uint BVH_FUNCTION_NAME(KernelGlobals *kg, const Ray *ray, Intersection *isect_array, int subsurface_object, uint *lcg_state, int max_hits) { /* todo: diff --git a/intern/cycles/kernel/kernel_bvh_traversal.h b/intern/cycles/kernel/kernel_bvh_traversal.h index a9264f318eb..8f69083575b 100644 --- a/intern/cycles/kernel/kernel_bvh_traversal.h +++ b/intern/cycles/kernel/kernel_bvh_traversal.h @@ -30,7 +30,7 @@ #define FEATURE(f) (((BVH_FUNCTION_FEATURES) & (f)) != 0) -__device bool BVH_FUNCTION_NAME +ccl_device bool BVH_FUNCTION_NAME (KernelGlobals *kg, const Ray *ray, Intersection *isect, const uint visibility #if FEATURE(BVH_HAIR_MINIMUM_WIDTH) , uint *lcg_state, float difl, float extmax diff --git a/intern/cycles/kernel/kernel_camera.h b/intern/cycles/kernel/kernel_camera.h index 966f28df05f..887b1afddd4 100644 --- a/intern/cycles/kernel/kernel_camera.h +++ b/intern/cycles/kernel/kernel_camera.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Perspective Camera */ -__device float2 camera_sample_aperture(KernelGlobals *kg, float u, float v) +ccl_device float2 camera_sample_aperture(KernelGlobals *kg, float u, float v) { float blades = kernel_data.cam.blades; @@ -33,7 +33,7 @@ __device float2 camera_sample_aperture(KernelGlobals *kg, float u, float v) } } -__device void camera_sample_perspective(KernelGlobals *kg, float raster_x, float raster_y, float lens_u, float lens_v, Ray *ray) +ccl_device void camera_sample_perspective(KernelGlobals *kg, float raster_x, float raster_y, float lens_u, float lens_v, Ray *ray) { /* create ray form raster position */ Transform rastertocamera = kernel_data.cam.rastertocamera; @@ -91,7 +91,7 @@ __device void camera_sample_perspective(KernelGlobals *kg, float raster_x, float /* Orthographic Camera */ -__device void camera_sample_orthographic(KernelGlobals *kg, float raster_x, float raster_y, float lens_u, float lens_v, Ray *ray) +ccl_device void camera_sample_orthographic(KernelGlobals *kg, float raster_x, float raster_y, float lens_u, float lens_v, Ray *ray) { /* create ray form raster position */ Transform rastertocamera = kernel_data.cam.rastertocamera; @@ -147,7 +147,7 @@ __device void camera_sample_orthographic(KernelGlobals *kg, float raster_x, floa /* Panorama Camera */ -__device void camera_sample_panorama(KernelGlobals *kg, float raster_x, float raster_y, float lens_u, float lens_v, Ray *ray) +ccl_device void camera_sample_panorama(KernelGlobals *kg, float raster_x, float raster_y, float lens_u, float lens_v, Ray *ray) { Transform rastertocamera = kernel_data.cam.rastertocamera; float3 Pcamera = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y, 0.0f)); @@ -216,7 +216,7 @@ __device void camera_sample_panorama(KernelGlobals *kg, float raster_x, float ra /* Common */ -__device void camera_sample(KernelGlobals *kg, int x, int y, float filter_u, float filter_v, +ccl_device void camera_sample(KernelGlobals *kg, int x, int y, float filter_u, float filter_v, float lens_u, float lens_v, float time, Ray *ray) { /* pixel filter */ @@ -243,13 +243,13 @@ __device void camera_sample(KernelGlobals *kg, int x, int y, float filter_u, flo /* Utilities */ -__device_inline float3 camera_position(KernelGlobals *kg) +ccl_device_inline float3 camera_position(KernelGlobals *kg) { Transform cameratoworld = kernel_data.cam.cameratoworld; return make_float3(cameratoworld.x.w, cameratoworld.y.w, cameratoworld.z.w); } -__device_inline float camera_distance(KernelGlobals *kg, float3 P) +ccl_device_inline float camera_distance(KernelGlobals *kg, float3 P) { Transform cameratoworld = kernel_data.cam.cameratoworld; float3 camP = make_float3(cameratoworld.x.w, cameratoworld.y.w, cameratoworld.z.w); @@ -262,7 +262,7 @@ __device_inline float camera_distance(KernelGlobals *kg, float3 P) return len(P - camP); } -__device_inline float3 camera_world_to_ndc(KernelGlobals *kg, ShaderData *sd, float3 P) +ccl_device_inline float3 camera_world_to_ndc(KernelGlobals *kg, ShaderData *sd, float3 P) { if(kernel_data.cam.type != CAMERA_PANORAMA) { /* perspective / ortho */ diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h index 44c2b9effe9..76f885aefe0 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -27,13 +27,12 @@ /* Qualifier wrappers for different names on different devices */ -#define __device __device__ __inline__ -#define __device_inline __device__ __inline__ -#define __device_noinline __device__ __noinline__ -#define __global -#define __shared __shared__ -#define __constant -#define __may_alias +#define ccl_device __device__ __inline__ +#define ccl_device_inline __device__ __inline__ +#define ccl_device_noinline __device__ __noinline__ +#define ccl_global +#define ccl_constant +#define ccl_may_alias /* No assert supported for CUDA */ diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h index e0102a01146..1ff3615e448 100644 --- a/intern/cycles/kernel/kernel_compat_opencl.h +++ b/intern/cycles/kernel/kernel_compat_opencl.h @@ -33,16 +33,18 @@ #endif #ifdef __CL_NOINLINE__ -#define __noinline __attribute__((noinline)) +#define ccl_noinline __attribute__((noinline)) #else -#define __noinline +#define ccl_noinline #endif /* in opencl all functions are device functions, so leave this empty */ -#define __device -#define __device_inline __device -#define __device_noinline __device __noinline -#define __may_alias +#define ccl_device +#define ccl_device_inline ccl_device +#define ccl_device_noinline ccl_device ccl_noinline +#define ccl_may_alias +#define ccl_constant __constant +#define ccl_global __global /* no assert in opencl */ #define kernel_assert(cond) diff --git a/intern/cycles/kernel/kernel_curve.h b/intern/cycles/kernel/kernel_curve.h index 9f7a1388a2b..821ac50eaa9 100644 --- a/intern/cycles/kernel/kernel_curve.h +++ b/intern/cycles/kernel/kernel_curve.h @@ -20,7 +20,7 @@ CCL_NAMESPACE_BEGIN /* curve attributes */ -__device float curve_attribute_float(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float *dx, float *dy) +ccl_device float curve_attribute_float(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float *dx, float *dy) { if(elem == ATTR_ELEMENT_CURVE) { #ifdef __RAY_DIFFERENTIALS__ @@ -55,7 +55,7 @@ __device float curve_attribute_float(KernelGlobals *kg, const ShaderData *sd, At } } -__device float3 curve_attribute_float3(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float3 *dx, float3 *dy) +ccl_device float3 curve_attribute_float3(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float3 *dx, float3 *dy) { if(elem == ATTR_ELEMENT_CURVE) { /* idea: we can't derive any useful differentials here, but for tiled @@ -96,7 +96,7 @@ __device float3 curve_attribute_float3(KernelGlobals *kg, const ShaderData *sd, /* hair info node functions */ -__device float curve_thickness(KernelGlobals *kg, ShaderData *sd) +ccl_device float curve_thickness(KernelGlobals *kg, ShaderData *sd) { float r = 0.0f; @@ -113,7 +113,7 @@ __device float curve_thickness(KernelGlobals *kg, ShaderData *sd) return r*2.0f; } -__device float3 curve_tangent_normal(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 curve_tangent_normal(KernelGlobals *kg, ShaderData *sd) { float3 tgN = make_float3(0.0f,0.0f,0.0f); diff --git a/intern/cycles/kernel/kernel_differential.h b/intern/cycles/kernel/kernel_differential.h index 71d6e87a4d9..daba2d927b7 100644 --- a/intern/cycles/kernel/kernel_differential.h +++ b/intern/cycles/kernel/kernel_differential.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* See "Tracing Ray Differentials", Homan Igehy, 1999. */ -__device void differential_transfer(differential3 *dP_, const differential3 dP, float3 D, const differential3 dD, float3 Ng, float t) +ccl_device void differential_transfer(differential3 *dP_, const differential3 dP, float3 D, const differential3 dD, float3 Ng, float t) { /* ray differential transfer through homogeneous medium, to * compute dPdx/dy at a shading point from the incoming ray */ @@ -31,7 +31,7 @@ __device void differential_transfer(differential3 *dP_, const differential3 dP, dP_->dy = tmpy - dot(tmpy, Ng)*tmp; } -__device void differential_incoming(differential3 *dI, const differential3 dD) +ccl_device void differential_incoming(differential3 *dI, const differential3 dD) { /* compute dIdx/dy at a shading point, we just need to negate the * differential of the ray direction */ @@ -40,7 +40,7 @@ __device void differential_incoming(differential3 *dI, const differential3 dD) dI->dy = -dD.dy; } -__device void differential_dudv(differential *du, differential *dv, float3 dPdu, float3 dPdv, differential3 dP, float3 Ng) +ccl_device void differential_dudv(differential *du, differential *dv, float3 dPdu, float3 dPdv, differential3 dP, float3 Ng) { /* now we have dPdx/dy from the ray differential transfer, and dPdu/dv * from the primitive, we can compute dudx/dy and dvdx/dy. these are @@ -84,7 +84,7 @@ __device void differential_dudv(differential *du, differential *dv, float3 dPdu, dv->dy = (dP.dy.y*dPdu.x - dP.dy.x*dPdu.y)*det; } -__device differential differential_zero() +ccl_device differential differential_zero() { differential d; d.dx = 0.0f; @@ -93,7 +93,7 @@ __device differential differential_zero() return d; } -__device differential3 differential3_zero() +ccl_device differential3 differential3_zero() { differential3 d; d.dx = make_float3(0.0f, 0.0f, 0.0f); diff --git a/intern/cycles/kernel/kernel_displace.h b/intern/cycles/kernel/kernel_displace.h index 38152b5571e..c50e2166660 100644 --- a/intern/cycles/kernel/kernel_displace.h +++ b/intern/cycles/kernel/kernel_displace.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void kernel_shader_evaluate(KernelGlobals *kg, __global uint4 *input, __global float4 *output, ShaderEvalType type, int i) +ccl_device void kernel_shader_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_global float4 *output, ShaderEvalType type, int i) { ShaderData sd; uint4 in = input[i]; diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index 0c8d69fb594..2ce0b758972 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Direction Emission */ -__device_noinline float3 direct_emissive_eval(KernelGlobals *kg, float rando, +ccl_device_noinline float3 direct_emissive_eval(KernelGlobals *kg, float rando, LightSample *ls, float u, float v, float3 I, differential3 dI, float t, float time, int bounce) { /* setup shading at emitter */ @@ -70,7 +70,7 @@ __device_noinline float3 direct_emissive_eval(KernelGlobals *kg, float rando, return eval; } -__device_noinline bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, +ccl_device_noinline bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, float randt, float rando, float randu, float randv, Ray *ray, BsdfEval *eval, bool *is_lamp, int bounce) { @@ -160,7 +160,7 @@ __device_noinline bool direct_emission(KernelGlobals *kg, ShaderData *sd, int li /* Indirect Primitive Emission */ -__device_noinline float3 indirect_primitive_emission(KernelGlobals *kg, ShaderData *sd, float t, int path_flag, float bsdf_pdf) +ccl_device_noinline float3 indirect_primitive_emission(KernelGlobals *kg, ShaderData *sd, float t, int path_flag, float bsdf_pdf) { /* evaluate emissive closure */ float3 L = shader_emissive_eval(kg, sd); @@ -183,7 +183,7 @@ __device_noinline float3 indirect_primitive_emission(KernelGlobals *kg, ShaderDa /* Indirect Lamp Emission */ -__device_noinline bool indirect_lamp_emission(KernelGlobals *kg, Ray *ray, int path_flag, float bsdf_pdf, float randt, float3 *emission, int bounce) +ccl_device_noinline bool indirect_lamp_emission(KernelGlobals *kg, Ray *ray, int path_flag, float bsdf_pdf, float randt, float3 *emission, int bounce) { LightSample ls; int lamp = lamp_light_eval_sample(kg, randt); @@ -222,7 +222,7 @@ __device_noinline bool indirect_lamp_emission(KernelGlobals *kg, Ray *ray, int p /* Indirect Background */ -__device_noinline float3 indirect_background(KernelGlobals *kg, Ray *ray, int path_flag, float bsdf_pdf, int bounce) +ccl_device_noinline float3 indirect_background(KernelGlobals *kg, Ray *ray, int path_flag, float bsdf_pdf, int bounce) { #ifdef __BACKGROUND__ int shader = kernel_data.background.shader; diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h index 370c550a515..b4118666491 100644 --- a/intern/cycles/kernel/kernel_film.h +++ b/intern/cycles/kernel/kernel_film.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device float4 film_map(KernelGlobals *kg, float4 irradiance, float scale) +ccl_device float4 film_map(KernelGlobals *kg, float4 irradiance, float scale) { float exposure = kernel_data.film.exposure; float4 result = irradiance*scale; @@ -32,7 +32,7 @@ __device float4 film_map(KernelGlobals *kg, float4 irradiance, float scale) return result; } -__device uchar4 film_float_to_byte(float4 color) +ccl_device uchar4 film_float_to_byte(float4 color) { uchar4 result; @@ -45,8 +45,8 @@ __device uchar4 film_float_to_byte(float4 color) return result; } -__device void kernel_film_convert_to_byte(KernelGlobals *kg, - __global uchar4 *rgba, __global float *buffer, +ccl_device void kernel_film_convert_to_byte(KernelGlobals *kg, + ccl_global uchar4 *rgba, ccl_global float *buffer, float sample_scale, int x, int y, int offset, int stride) { /* buffer offset */ @@ -56,22 +56,22 @@ __device void kernel_film_convert_to_byte(KernelGlobals *kg, buffer += index*kernel_data.film.pass_stride; /* map colors */ - float4 irradiance = *((__global float4*)buffer); + float4 irradiance = *((ccl_global float4*)buffer); float4 float_result = film_map(kg, irradiance, sample_scale); uchar4 byte_result = film_float_to_byte(float_result); *rgba = byte_result; } -__device void kernel_film_convert_to_half_float(KernelGlobals *kg, - __global uchar4 *rgba, __global float *buffer, +ccl_device void kernel_film_convert_to_half_float(KernelGlobals *kg, + ccl_global uchar4 *rgba, ccl_global float *buffer, float sample_scale, int x, int y, int offset, int stride) { /* buffer offset */ int index = offset + x + y*stride; - __global float4 *in = (__global float4*)(buffer + index*kernel_data.film.pass_stride); - __global half *out = (__global half*)rgba + index*4; + ccl_global float4 *in = (ccl_global float4*)(buffer + index*kernel_data.film.pass_stride); + ccl_global half *out = (ccl_global half*)rgba + index*4; float exposure = kernel_data.film.exposure; diff --git a/intern/cycles/kernel/kernel_globals.h b/intern/cycles/kernel/kernel_globals.h index b5e691eb615..e60bd6c0067 100644 --- a/intern/cycles/kernel/kernel_globals.h +++ b/intern/cycles/kernel/kernel_globals.h @@ -81,10 +81,10 @@ typedef struct KernelGlobals {} KernelGlobals; #ifdef __KERNEL_OPENCL__ typedef struct KernelGlobals { - __constant KernelData *data; + ccl_constant KernelData *data; #define KERNEL_TEX(type, ttype, name) \ - __global type *name; + ccl_global type *name; #include "kernel_textures.h" } KernelGlobals; @@ -92,7 +92,7 @@ typedef struct KernelGlobals { /* Interpolated lookup table access */ -__device float lookup_table_read(KernelGlobals *kg, float x, int offset, int size) +ccl_device float lookup_table_read(KernelGlobals *kg, float x, int offset, int size) { x = clamp(x, 0.0f, 1.0f)*(size-1); @@ -108,7 +108,7 @@ __device float lookup_table_read(KernelGlobals *kg, float x, int offset, int siz return (1.0f - t)*data0 + t*data1; } -__device float lookup_table_read_2D(KernelGlobals *kg, float x, float y, int offset, int xsize, int ysize) +ccl_device float lookup_table_read_2D(KernelGlobals *kg, float x, float y, int offset, int xsize, int ysize) { y = clamp(y, 0.0f, 1.0f)*(ysize-1); diff --git a/intern/cycles/kernel/kernel_jitter.h b/intern/cycles/kernel/kernel_jitter.h index 18666b51c0c..7a850844bf2 100644 --- a/intern/cycles/kernel/kernel_jitter.h +++ b/intern/cycles/kernel/kernel_jitter.h @@ -22,18 +22,18 @@ CCL_NAMESPACE_BEGIN /* todo: find good value, suggested 64 gives pattern on cornell box ceiling */ #define CMJ_RANDOM_OFFSET_LIMIT 4096 -__device_inline bool cmj_is_pow2(int i) +ccl_device_inline bool cmj_is_pow2(int i) { return (i & (i - 1)) == 0; } -__device_inline int cmj_fast_mod_pow2(int a, int b) +ccl_device_inline int cmj_fast_mod_pow2(int a, int b) { return (a & (b - 1)); } /* a must be > 0 and b must be > 1 */ -__device_inline int cmj_fast_div_pow2(int a, int b) +ccl_device_inline int cmj_fast_div_pow2(int a, int b) { #if defined(__KERNEL_SSE2__) && !defined(_MSC_VER) return a >> __builtin_ctz(b); @@ -42,7 +42,7 @@ __device_inline int cmj_fast_div_pow2(int a, int b) #endif } -__device_inline uint cmj_w_mask(uint w) +ccl_device_inline uint cmj_w_mask(uint w) { #if defined(__KERNEL_SSE2__) && !defined(_MSC_VER) return ((1 << (32 - __builtin_clz(w))) - 1); @@ -57,7 +57,7 @@ __device_inline uint cmj_w_mask(uint w) #endif } -__device_inline uint cmj_permute(uint i, uint l, uint p) +ccl_device_inline uint cmj_permute(uint i, uint l, uint p) { uint w = l - 1; @@ -113,7 +113,7 @@ __device_inline uint cmj_permute(uint i, uint l, uint p) } } -__device_inline uint cmj_hash(uint i, uint p) +ccl_device_inline uint cmj_hash(uint i, uint p) { i ^= p; i ^= i >> 17; @@ -129,13 +129,13 @@ __device_inline uint cmj_hash(uint i, uint p) return i; } -__device_inline float cmj_randfloat(uint i, uint p) +ccl_device_inline float cmj_randfloat(uint i, uint p) { return cmj_hash(i, p) * (1.0f / 4294967808.0f); } #ifdef __CMJ__ -__device float cmj_sample_1D(int s, int N, int p) +ccl_device float cmj_sample_1D(int s, int N, int p) { kernel_assert(s < N); @@ -146,7 +146,7 @@ __device float cmj_sample_1D(int s, int N, int p) return (x + jx)*invN; } -__device void cmj_sample_2D(int s, int N, int p, float *fx, float *fy) +ccl_device void cmj_sample_2D(int s, int N, int p, float *fx, float *fy) { kernel_assert(s < N); diff --git a/intern/cycles/kernel/kernel_light.h b/intern/cycles/kernel/kernel_light.h index 6b00bd2ab01..9915cd2495f 100644 --- a/intern/cycles/kernel/kernel_light.h +++ b/intern/cycles/kernel/kernel_light.h @@ -36,7 +36,7 @@ typedef struct LightSample { #ifdef __BACKGROUND_MIS__ -__device float3 background_light_sample(KernelGlobals *kg, float randu, float randv, float *pdf) +ccl_device float3 background_light_sample(KernelGlobals *kg, float randu, float randv, float *pdf) { /* for the following, the CDF values are actually a pair of floats, with the * function value as X and the actual CDF as Y. The last entry's function @@ -112,7 +112,7 @@ __device float3 background_light_sample(KernelGlobals *kg, float randu, float ra return -equirectangular_to_direction(u, v); } -__device float background_light_pdf(KernelGlobals *kg, float3 direction) +ccl_device float background_light_pdf(KernelGlobals *kg, float3 direction) { float2 uv = direction_to_equirectangular(direction); int res = kernel_data.integrator.pdf_background_res; @@ -146,7 +146,7 @@ __device float background_light_pdf(KernelGlobals *kg, float3 direction) /* Regular Light */ -__device float3 disk_light_sample(float3 v, float randu, float randv) +ccl_device float3 disk_light_sample(float3 v, float randu, float randv) { float3 ru, rv; @@ -156,17 +156,17 @@ __device float3 disk_light_sample(float3 v, float randu, float randv) return ru*randu + rv*randv; } -__device float3 distant_light_sample(float3 D, float radius, float randu, float randv) +ccl_device float3 distant_light_sample(float3 D, float radius, float randu, float randv) { return normalize(D + disk_light_sample(D, randu, randv)*radius); } -__device float3 sphere_light_sample(float3 P, float3 center, float radius, float randu, float randv) +ccl_device float3 sphere_light_sample(float3 P, float3 center, float radius, float randu, float randv) { return disk_light_sample(normalize(P - center), randu, randv)*radius; } -__device float3 area_light_sample(float3 axisu, float3 axisv, float randu, float randv) +ccl_device float3 area_light_sample(float3 axisu, float3 axisv, float randu, float randv) { randu = randu - 0.5f; randv = randv - 0.5f; @@ -174,7 +174,7 @@ __device float3 area_light_sample(float3 axisu, float3 axisv, float randu, float return axisu*randu + axisv*randv; } -__device float spot_light_attenuation(float4 data1, float4 data2, LightSample *ls) +ccl_device float spot_light_attenuation(float4 data1, float4 data2, LightSample *ls) { float3 dir = make_float3(data2.y, data2.z, data2.w); float3 I = ls->Ng; @@ -197,7 +197,7 @@ __device float spot_light_attenuation(float4 data1, float4 data2, LightSample *l return attenuation; } -__device float lamp_light_pdf(KernelGlobals *kg, const float3 Ng, const float3 I, float t) +ccl_device float lamp_light_pdf(KernelGlobals *kg, const float3 Ng, const float3 I, float t) { float cos_pi = dot(Ng, I); @@ -207,7 +207,7 @@ __device float lamp_light_pdf(KernelGlobals *kg, const float3 Ng, const float3 I return t*t/cos_pi; } -__device void lamp_light_sample(KernelGlobals *kg, int lamp, +ccl_device void lamp_light_sample(KernelGlobals *kg, int lamp, float randu, float randv, float3 P, LightSample *ls) { float4 data0 = kernel_tex_fetch(__light_data, lamp*LIGHT_SIZE + 0); @@ -298,7 +298,7 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp, } } -__device bool lamp_light_eval(KernelGlobals *kg, int lamp, float3 P, float3 D, float t, LightSample *ls) +ccl_device bool lamp_light_eval(KernelGlobals *kg, int lamp, float3 P, float3 D, float t, LightSample *ls) { float4 data0 = kernel_tex_fetch(__light_data, lamp*LIGHT_SIZE + 0); float4 data1 = kernel_tex_fetch(__light_data, lamp*LIGHT_SIZE + 1); @@ -422,7 +422,7 @@ __device bool lamp_light_eval(KernelGlobals *kg, int lamp, float3 P, float3 D, f /* Triangle Light */ -__device void object_transform_light_sample(KernelGlobals *kg, LightSample *ls, int object, float time) +ccl_device void object_transform_light_sample(KernelGlobals *kg, LightSample *ls, int object, float time) { #ifdef __INSTANCING__ /* instance transform */ @@ -440,7 +440,7 @@ __device void object_transform_light_sample(KernelGlobals *kg, LightSample *ls, #endif } -__device void triangle_light_sample(KernelGlobals *kg, int prim, int object, +ccl_device void triangle_light_sample(KernelGlobals *kg, int prim, int object, float randu, float randv, float time, LightSample *ls) { /* triangle, so get position, normal, shader */ @@ -457,7 +457,7 @@ __device void triangle_light_sample(KernelGlobals *kg, int prim, int object, object_transform_light_sample(kg, ls, object, time); } -__device float triangle_light_pdf(KernelGlobals *kg, +ccl_device float triangle_light_pdf(KernelGlobals *kg, const float3 Ng, const float3 I, float t) { float pdf = kernel_data.integrator.pdf_triangles; @@ -473,7 +473,7 @@ __device float triangle_light_pdf(KernelGlobals *kg, #ifdef __HAIR__ -__device void curve_segment_light_sample(KernelGlobals *kg, int prim, int object, +ccl_device void curve_segment_light_sample(KernelGlobals *kg, int prim, int object, int segment, float randu, float randv, float time, LightSample *ls) { /* this strand code needs completion */ @@ -515,7 +515,7 @@ __device void curve_segment_light_sample(KernelGlobals *kg, int prim, int object /* Light Distribution */ -__device int light_distribution_sample(KernelGlobals *kg, float randt) +ccl_device int light_distribution_sample(KernelGlobals *kg, float randt) { /* this is basically std::upper_bound as used by pbrt, to find a point light or * triangle to emit from, proportional to area. a good improvement would be to @@ -544,7 +544,7 @@ __device int light_distribution_sample(KernelGlobals *kg, float randt) /* Generic Light */ -__device void light_sample(KernelGlobals *kg, float randt, float randu, float randv, float time, float3 P, LightSample *ls) +ccl_device void light_sample(KernelGlobals *kg, float randt, float randu, float randv, float time, float3 P, LightSample *ls) { /* sample index */ int index = light_distribution_sample(kg, randt); @@ -577,18 +577,18 @@ __device void light_sample(KernelGlobals *kg, float randt, float randu, float ra } } -__device int light_select_num_samples(KernelGlobals *kg, int index) +ccl_device int light_select_num_samples(KernelGlobals *kg, int index) { float4 data3 = kernel_tex_fetch(__light_data, index*LIGHT_SIZE + 3); return __float_as_int(data3.x); } -__device void light_select(KernelGlobals *kg, int index, float randu, float randv, float3 P, LightSample *ls) +ccl_device void light_select(KernelGlobals *kg, int index, float randu, float randv, float3 P, LightSample *ls) { lamp_light_sample(kg, index, randu, randv, P, ls); } -__device int lamp_light_eval_sample(KernelGlobals *kg, float randt) +ccl_device int lamp_light_eval_sample(KernelGlobals *kg, float randt) { /* sample index */ int index = light_distribution_sample(kg, randt); diff --git a/intern/cycles/kernel/kernel_montecarlo.h b/intern/cycles/kernel/kernel_montecarlo.h index b3d53e00be7..92f3420a218 100644 --- a/intern/cycles/kernel/kernel_montecarlo.h +++ b/intern/cycles/kernel/kernel_montecarlo.h @@ -36,7 +36,7 @@ CCL_NAMESPACE_BEGIN /* distribute uniform xy on [0,1] over unit disk [-1,1] */ -__device void to_unit_disk(float *x, float *y) +ccl_device void to_unit_disk(float *x, float *y) { float phi = M_2PI_F * (*x); float r = sqrtf(*y); @@ -47,14 +47,14 @@ __device void to_unit_disk(float *x, float *y) /* return an orthogonal tangent and bitangent given a normal and tangent that * may not be exactly orthogonal */ -__device void make_orthonormals_tangent(const float3 N, const float3 T, float3 *a, float3 *b) +ccl_device void make_orthonormals_tangent(const float3 N, const float3 T, float3 *a, float3 *b) { *b = normalize(cross(N, T)); *a = cross(*b, N); } /* sample direction with cosine weighted distributed in hemisphere */ -__device_inline void sample_cos_hemisphere(const float3 N, +ccl_device_inline void sample_cos_hemisphere(const float3 N, float randu, float randv, float3 *omega_in, float *pdf) { to_unit_disk(&randu, &randv); @@ -66,7 +66,7 @@ __device_inline void sample_cos_hemisphere(const float3 N, } /* sample direction uniformly distributed in hemisphere */ -__device_inline void sample_uniform_hemisphere(const float3 N, +ccl_device_inline void sample_uniform_hemisphere(const float3 N, float randu, float randv, float3 *omega_in, float *pdf) { @@ -83,7 +83,7 @@ __device_inline void sample_uniform_hemisphere(const float3 N, } /* sample direction uniformly distributed in cone */ -__device_inline void sample_uniform_cone(const float3 N, float angle, +ccl_device_inline void sample_uniform_cone(const float3 N, float angle, float randu, float randv, float3 *omega_in, float *pdf) { @@ -100,7 +100,7 @@ __device_inline void sample_uniform_cone(const float3 N, float angle, } /* sample uniform point on the surface of a sphere */ -__device float3 sample_uniform_sphere(float u1, float u2) +ccl_device float3 sample_uniform_sphere(float u1, float u2) { float z = 1.0f - 2.0f*u1; float r = sqrtf(fmaxf(0.0f, 1.0f - z*z)); @@ -111,29 +111,29 @@ __device float3 sample_uniform_sphere(float u1, float u2) return make_float3(x, y, z); } -__device float balance_heuristic(float a, float b) +ccl_device float balance_heuristic(float a, float b) { return (a)/(a + b); } -__device float balance_heuristic_3(float a, float b, float c) +ccl_device float balance_heuristic_3(float a, float b, float c) { return (a)/(a + b + c); } -__device float power_heuristic(float a, float b) +ccl_device float power_heuristic(float a, float b) { return (a*a)/(a*a + b*b); } -__device float power_heuristic_3(float a, float b, float c) +ccl_device float power_heuristic_3(float a, float b, float c) { return (a*a)/(a*a + b*b + c*c); } /* distribute uniform xy on [0,1] over unit disk [-1,1], with concentric mapping * to better preserve stratification for some RNG sequences */ -__device float2 concentric_sample_disk(float u1, float u2) +ccl_device float2 concentric_sample_disk(float u1, float u2) { float phi, r; float a = 2.0f*u1 - 1.0f; @@ -155,7 +155,7 @@ __device float2 concentric_sample_disk(float u1, float u2) } /* sample point in unit polygon with given number of corners and rotation */ -__device float2 regular_polygon_sample(float corners, float rotation, float u, float v) +ccl_device float2 regular_polygon_sample(float corners, float rotation, float u, float v) { /* sample corner number and reuse u */ float corner = floorf(u*corners); diff --git a/intern/cycles/kernel/kernel_object.h b/intern/cycles/kernel/kernel_object.h index d0aae119476..a66277e10cd 100644 --- a/intern/cycles/kernel/kernel_object.h +++ b/intern/cycles/kernel/kernel_object.h @@ -30,7 +30,7 @@ enum ObjectVectorTransform { OBJECT_VECTOR_MOTION_POST = 3 }; -__device_inline Transform object_fetch_transform(KernelGlobals *kg, int object, enum ObjectTransform type) +ccl_device_inline Transform object_fetch_transform(KernelGlobals *kg, int object, enum ObjectTransform type) { int offset = object*OBJECT_SIZE + (int)type; @@ -43,7 +43,7 @@ __device_inline Transform object_fetch_transform(KernelGlobals *kg, int object, return tfm; } -__device_inline Transform object_fetch_vector_transform(KernelGlobals *kg, int object, enum ObjectVectorTransform type) +ccl_device_inline Transform object_fetch_vector_transform(KernelGlobals *kg, int object, enum ObjectVectorTransform type) { int offset = object*OBJECT_VECTOR_SIZE + (int)type; @@ -57,7 +57,7 @@ __device_inline Transform object_fetch_vector_transform(KernelGlobals *kg, int o } #ifdef __OBJECT_MOTION__ -__device_inline Transform object_fetch_transform_motion(KernelGlobals *kg, int object, float time) +ccl_device_inline Transform object_fetch_transform_motion(KernelGlobals *kg, int object, float time) { DecompMotionTransform motion; @@ -79,7 +79,7 @@ __device_inline Transform object_fetch_transform_motion(KernelGlobals *kg, int o return tfm; } -__device_inline Transform object_fetch_transform_motion_test(KernelGlobals *kg, int object, float time, Transform *itfm) +ccl_device_inline Transform object_fetch_transform_motion_test(KernelGlobals *kg, int object, float time, Transform *itfm) { int object_flag = kernel_tex_fetch(__object_flag, object); @@ -102,7 +102,7 @@ __device_inline Transform object_fetch_transform_motion_test(KernelGlobals *kg, } #endif -__device_inline void object_position_transform(KernelGlobals *kg, ShaderData *sd, float3 *P) +ccl_device_inline void object_position_transform(KernelGlobals *kg, ShaderData *sd, float3 *P) { #ifdef __OBJECT_MOTION__ *P = transform_point(&sd->ob_tfm, *P); @@ -112,7 +112,7 @@ __device_inline void object_position_transform(KernelGlobals *kg, ShaderData *sd #endif } -__device_inline void object_inverse_position_transform(KernelGlobals *kg, ShaderData *sd, float3 *P) +ccl_device_inline void object_inverse_position_transform(KernelGlobals *kg, ShaderData *sd, float3 *P) { #ifdef __OBJECT_MOTION__ *P = transform_point(&sd->ob_itfm, *P); @@ -122,7 +122,7 @@ __device_inline void object_inverse_position_transform(KernelGlobals *kg, Shader #endif } -__device_inline void object_inverse_normal_transform(KernelGlobals *kg, ShaderData *sd, float3 *N) +ccl_device_inline void object_inverse_normal_transform(KernelGlobals *kg, ShaderData *sd, float3 *N) { #ifdef __OBJECT_MOTION__ *N = normalize(transform_direction_transposed(&sd->ob_tfm, *N)); @@ -132,7 +132,7 @@ __device_inline void object_inverse_normal_transform(KernelGlobals *kg, ShaderDa #endif } -__device_inline void object_normal_transform(KernelGlobals *kg, ShaderData *sd, float3 *N) +ccl_device_inline void object_normal_transform(KernelGlobals *kg, ShaderData *sd, float3 *N) { #ifdef __OBJECT_MOTION__ *N = normalize(transform_direction_transposed(&sd->ob_itfm, *N)); @@ -142,7 +142,7 @@ __device_inline void object_normal_transform(KernelGlobals *kg, ShaderData *sd, #endif } -__device_inline void object_dir_transform(KernelGlobals *kg, ShaderData *sd, float3 *D) +ccl_device_inline void object_dir_transform(KernelGlobals *kg, ShaderData *sd, float3 *D) { #ifdef __OBJECT_MOTION__ *D = transform_direction(&sd->ob_tfm, *D); @@ -152,7 +152,7 @@ __device_inline void object_dir_transform(KernelGlobals *kg, ShaderData *sd, flo #endif } -__device_inline void object_inverse_dir_transform(KernelGlobals *kg, ShaderData *sd, float3 *D) +ccl_device_inline void object_inverse_dir_transform(KernelGlobals *kg, ShaderData *sd, float3 *D) { #ifdef __OBJECT_MOTION__ *D = transform_direction(&sd->ob_itfm, *D); @@ -162,7 +162,7 @@ __device_inline void object_inverse_dir_transform(KernelGlobals *kg, ShaderData #endif } -__device_inline float3 object_location(KernelGlobals *kg, ShaderData *sd) +ccl_device_inline float3 object_location(KernelGlobals *kg, ShaderData *sd) { if(sd->object == ~0) return make_float3(0.0f, 0.0f, 0.0f); @@ -175,14 +175,14 @@ __device_inline float3 object_location(KernelGlobals *kg, ShaderData *sd) #endif } -__device_inline float object_surface_area(KernelGlobals *kg, int object) +ccl_device_inline float object_surface_area(KernelGlobals *kg, int object) { int offset = object*OBJECT_SIZE + OBJECT_PROPERTIES; float4 f = kernel_tex_fetch(__objects, offset); return f.x; } -__device_inline float object_pass_id(KernelGlobals *kg, int object) +ccl_device_inline float object_pass_id(KernelGlobals *kg, int object) { if(object == ~0) return 0.0f; @@ -192,7 +192,7 @@ __device_inline float object_pass_id(KernelGlobals *kg, int object) return f.y; } -__device_inline float object_random_number(KernelGlobals *kg, int object) +ccl_device_inline float object_random_number(KernelGlobals *kg, int object) { if(object == ~0) return 0.0f; @@ -202,7 +202,7 @@ __device_inline float object_random_number(KernelGlobals *kg, int object) return f.z; } -__device_inline uint object_particle_id(KernelGlobals *kg, int object) +ccl_device_inline uint object_particle_id(KernelGlobals *kg, int object) { if(object == ~0) return 0.0f; @@ -212,7 +212,7 @@ __device_inline uint object_particle_id(KernelGlobals *kg, int object) return __float_as_uint(f.w); } -__device_inline float3 object_dupli_generated(KernelGlobals *kg, int object) +ccl_device_inline float3 object_dupli_generated(KernelGlobals *kg, int object) { if(object == ~0) return make_float3(0.0f, 0.0f, 0.0f); @@ -222,7 +222,7 @@ __device_inline float3 object_dupli_generated(KernelGlobals *kg, int object) return make_float3(f.x, f.y, f.z); } -__device_inline float3 object_dupli_uv(KernelGlobals *kg, int object) +ccl_device_inline float3 object_dupli_uv(KernelGlobals *kg, int object) { if(object == ~0) return make_float3(0.0f, 0.0f, 0.0f); @@ -233,54 +233,54 @@ __device_inline float3 object_dupli_uv(KernelGlobals *kg, int object) } -__device int shader_pass_id(KernelGlobals *kg, ShaderData *sd) +ccl_device int shader_pass_id(KernelGlobals *kg, ShaderData *sd) { return kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2 + 1); } -__device_inline float particle_index(KernelGlobals *kg, int particle) +ccl_device_inline float particle_index(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f = kernel_tex_fetch(__particles, offset + 0); return f.x; } -__device float particle_age(KernelGlobals *kg, int particle) +ccl_device float particle_age(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f = kernel_tex_fetch(__particles, offset + 0); return f.y; } -__device float particle_lifetime(KernelGlobals *kg, int particle) +ccl_device float particle_lifetime(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f = kernel_tex_fetch(__particles, offset + 0); return f.z; } -__device float particle_size(KernelGlobals *kg, int particle) +ccl_device float particle_size(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f = kernel_tex_fetch(__particles, offset + 0); return f.w; } -__device float4 particle_rotation(KernelGlobals *kg, int particle) +ccl_device float4 particle_rotation(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f = kernel_tex_fetch(__particles, offset + 1); return f; } -__device float3 particle_location(KernelGlobals *kg, int particle) +ccl_device float3 particle_location(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f = kernel_tex_fetch(__particles, offset + 2); return make_float3(f.x, f.y, f.z); } -__device float3 particle_velocity(KernelGlobals *kg, int particle) +ccl_device float3 particle_velocity(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f2 = kernel_tex_fetch(__particles, offset + 2); @@ -288,7 +288,7 @@ __device float3 particle_velocity(KernelGlobals *kg, int particle) return make_float3(f2.w, f3.x, f3.y); } -__device float3 particle_angular_velocity(KernelGlobals *kg, int particle) +ccl_device float3 particle_angular_velocity(KernelGlobals *kg, int particle) { int offset = particle*PARTICLE_SIZE; float4 f3 = kernel_tex_fetch(__particles, offset + 3); diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h index 5e91b13f90c..512db9ec392 100644 --- a/intern/cycles/kernel/kernel_passes.h +++ b/intern/cycles/kernel/kernel_passes.h @@ -16,25 +16,25 @@ CCL_NAMESPACE_BEGIN -__device_inline void kernel_write_pass_float(__global float *buffer, int sample, float value) +ccl_device_inline void kernel_write_pass_float(ccl_global float *buffer, int sample, float value) { - __global float *buf = buffer; + ccl_global float *buf = buffer; *buf = (sample == 0)? value: *buf + value; } -__device_inline void kernel_write_pass_float3(__global float *buffer, int sample, float3 value) +ccl_device_inline void kernel_write_pass_float3(ccl_global float *buffer, int sample, float3 value) { - __global float3 *buf = (__global float3*)buffer; + ccl_global float3 *buf = (ccl_global float3*)buffer; *buf = (sample == 0)? value: *buf + value; } -__device_inline void kernel_write_pass_float4(__global float *buffer, int sample, float4 value) +ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, int sample, float4 value) { - __global float4 *buf = (__global float4*)buffer; + ccl_global float4 *buf = (ccl_global float4*)buffer; *buf = (sample == 0)? value: *buf + value; } -__device_inline void kernel_write_data_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, +ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L, ShaderData *sd, int sample, int path_flag, float3 throughput) { #ifdef __PASSES__ @@ -114,7 +114,7 @@ __device_inline void kernel_write_data_passes(KernelGlobals *kg, __global float #endif } -__device_inline void kernel_write_light_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, int sample) +ccl_device_inline void kernel_write_light_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L, int sample) { #ifdef __PASSES__ int flag = kernel_data.film.pass_flag; diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 5354738d378..4f3957b66ef 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -42,7 +42,7 @@ CCL_NAMESPACE_BEGIN -__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, float3 *shadow) +ccl_device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, float3 *shadow) { *shadow = make_float3(1.0f, 1.0f, 1.0f); @@ -122,7 +122,7 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra #if defined(__BRANCHED_PATH__) || defined(__SUBSURFACE__) -__device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer, +ccl_device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray ray, ccl_global float *buffer, float3 throughput, int num_samples, int num_total_samples, float min_ray_pdf, float ray_pdf, PathState state, int rng_offset, PathRadiance *L) { @@ -359,7 +359,7 @@ __device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray #ifdef __SUBSURFACE__ -__device_inline bool kernel_path_integrate_lighting(KernelGlobals *kg, RNG *rng, +ccl_device_inline bool kernel_path_integrate_lighting(KernelGlobals *kg, RNG *rng, int sample, int num_samples, ShaderData *sd, float3 *throughput, float *min_ray_pdf, float *ray_pdf, PathState *state, @@ -452,7 +452,7 @@ __device_inline bool kernel_path_integrate_lighting(KernelGlobals *kg, RNG *rng, #endif -__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer) +ccl_device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, ccl_global float *buffer) { /* initialize */ PathRadiance L; @@ -790,11 +790,11 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R #ifdef __BRANCHED_PATH__ -__device_noinline void kernel_branched_path_integrate_lighting(KernelGlobals *kg, RNG *rng, +ccl_device_noinline void kernel_branched_path_integrate_lighting(KernelGlobals *kg, RNG *rng, int sample, int aa_samples, ShaderData *sd, float3 throughput, float num_samples_adjust, float min_ray_pdf, float ray_pdf, PathState state, - int rng_offset, PathRadiance *L, __global float *buffer) + int rng_offset, PathRadiance *L, ccl_global float *buffer) { #ifdef __EMISSION__ /* sample illumination from lights to find path contribution */ @@ -941,7 +941,7 @@ __device_noinline void kernel_branched_path_integrate_lighting(KernelGlobals *kg } } -__device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer) +ccl_device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, ccl_global float *buffer) { /* initialize */ PathRadiance L; @@ -1166,7 +1166,7 @@ __device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, int #endif -__device_inline void kernel_path_trace_setup(KernelGlobals *kg, __global uint *rng_state, int sample, int x, int y, RNG *rng, Ray *ray) +ccl_device_inline void kernel_path_trace_setup(KernelGlobals *kg, ccl_global uint *rng_state, int sample, int x, int y, RNG *rng, Ray *ray) { float filter_u; float filter_v; @@ -1195,8 +1195,8 @@ __device_inline void kernel_path_trace_setup(KernelGlobals *kg, __global uint *r camera_sample(kg, x, y, filter_u, filter_v, lens_u, lens_v, time, ray); } -__device void kernel_path_trace(KernelGlobals *kg, - __global float *buffer, __global uint *rng_state, +ccl_device void kernel_path_trace(KernelGlobals *kg, + ccl_global float *buffer, ccl_global uint *rng_state, int sample, int x, int y, int offset, int stride) { /* buffer offset */ @@ -1227,8 +1227,8 @@ __device void kernel_path_trace(KernelGlobals *kg, } #ifdef __BRANCHED_PATH__ -__device void kernel_branched_path_trace(KernelGlobals *kg, - __global float *buffer, __global uint *rng_state, +ccl_device void kernel_branched_path_trace(KernelGlobals *kg, + ccl_global float *buffer, ccl_global uint *rng_state, int sample, int x, int y, int offset, int stride) { /* buffer offset */ diff --git a/intern/cycles/kernel/kernel_path_state.h b/intern/cycles/kernel/kernel_path_state.h index e0e0f43fd26..0ded332b3b9 100644 --- a/intern/cycles/kernel/kernel_path_state.h +++ b/intern/cycles/kernel/kernel_path_state.h @@ -26,7 +26,7 @@ typedef struct PathState { int transparent_bounce; } PathState; -__device_inline void path_state_init(PathState *state) +ccl_device_inline void path_state_init(PathState *state) { state->flag = PATH_RAY_CAMERA|PATH_RAY_SINGULAR|PATH_RAY_MIS_SKIP; state->bounce = 0; @@ -36,7 +36,7 @@ __device_inline void path_state_init(PathState *state) state->transparent_bounce = 0; } -__device_inline void path_state_next(KernelGlobals *kg, PathState *state, int label) +ccl_device_inline void path_state_next(KernelGlobals *kg, PathState *state, int label) { /* ray through transparent keeps same flags from previous ray and is * not counted as a regular bounce, transparent has separate max */ @@ -88,7 +88,7 @@ __device_inline void path_state_next(KernelGlobals *kg, PathState *state, int la } } -__device_inline uint path_state_ray_visibility(KernelGlobals *kg, PathState *state) +ccl_device_inline uint path_state_ray_visibility(KernelGlobals *kg, PathState *state) { uint flag = state->flag & PATH_RAY_ALL_VISIBILITY; @@ -102,7 +102,7 @@ __device_inline uint path_state_ray_visibility(KernelGlobals *kg, PathState *sta return flag; } -__device_inline float path_state_terminate_probability(KernelGlobals *kg, PathState *state, const float3 throughput) +ccl_device_inline float path_state_terminate_probability(KernelGlobals *kg, PathState *state, const float3 throughput) { if(state->flag & PATH_RAY_TRANSPARENT) { /* transparent rays treated separately */ diff --git a/intern/cycles/kernel/kernel_primitive.h b/intern/cycles/kernel/kernel_primitive.h index 636cfd06532..ababad28f35 100644 --- a/intern/cycles/kernel/kernel_primitive.h +++ b/intern/cycles/kernel/kernel_primitive.h @@ -21,7 +21,7 @@ CCL_NAMESPACE_BEGIN /* attribute lookup */ -__device_inline int find_attribute(KernelGlobals *kg, ShaderData *sd, uint id, AttributeElement *elem) +ccl_device_inline int find_attribute(KernelGlobals *kg, ShaderData *sd, uint id, AttributeElement *elem) { if(sd->object == ~0) return (int)ATTR_STD_NOT_FOUND; @@ -52,7 +52,7 @@ __device_inline int find_attribute(KernelGlobals *kg, ShaderData *sd, uint id, A } } -__device float primitive_attribute_float(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float *dx, float *dy) +ccl_device float primitive_attribute_float(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float *dx, float *dy) { #ifdef __HAIR__ if(sd->segment == ~0) @@ -64,7 +64,7 @@ __device float primitive_attribute_float(KernelGlobals *kg, const ShaderData *sd #endif } -__device float3 primitive_attribute_float3(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float3 *dx, float3 *dy) +ccl_device float3 primitive_attribute_float3(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float3 *dx, float3 *dy) { #ifdef __HAIR__ if(sd->segment == ~0) @@ -76,7 +76,7 @@ __device float3 primitive_attribute_float3(KernelGlobals *kg, const ShaderData * #endif } -__device float3 primitive_uv(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 primitive_uv(KernelGlobals *kg, ShaderData *sd) { AttributeElement elem_uv; int offset_uv = find_attribute(kg, sd, ATTR_STD_UV, &elem_uv); @@ -89,7 +89,7 @@ __device float3 primitive_uv(KernelGlobals *kg, ShaderData *sd) return uv; } -__device float3 primitive_tangent(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 primitive_tangent(KernelGlobals *kg, ShaderData *sd) { #ifdef __HAIR__ if(sd->segment != ~0) @@ -122,7 +122,7 @@ __device float3 primitive_tangent(KernelGlobals *kg, ShaderData *sd) /* motion */ -__device float4 primitive_motion_vector(KernelGlobals *kg, ShaderData *sd) +ccl_device float4 primitive_motion_vector(KernelGlobals *kg, ShaderData *sd) { float3 motion_pre = sd->P, motion_post = sd->P; diff --git a/intern/cycles/kernel/kernel_projection.h b/intern/cycles/kernel/kernel_projection.h index d9520de7956..e2108604bc8 100644 --- a/intern/cycles/kernel/kernel_projection.h +++ b/intern/cycles/kernel/kernel_projection.h @@ -37,7 +37,7 @@ CCL_NAMESPACE_BEGIN /* Spherical coordinates <-> Cartesian direction */ -__device float2 direction_to_spherical(float3 dir) +ccl_device float2 direction_to_spherical(float3 dir) { float theta = acosf(dir.z); float phi = atan2f(dir.x, dir.y); @@ -45,7 +45,7 @@ __device float2 direction_to_spherical(float3 dir) return make_float2(theta, phi); } -__device float3 spherical_to_direction(float theta, float phi) +ccl_device float3 spherical_to_direction(float theta, float phi) { return make_float3( sinf(theta)*cosf(phi), @@ -55,7 +55,7 @@ __device float3 spherical_to_direction(float theta, float phi) /* Equirectangular coordinates <-> Cartesian direction */ -__device float2 direction_to_equirectangular(float3 dir) +ccl_device float2 direction_to_equirectangular(float3 dir) { float u = -atan2f(dir.y, dir.x)/(M_2PI_F) + 0.5f; float v = atan2f(dir.z, hypotf(dir.x, dir.y))/M_PI_F + 0.5f; @@ -63,7 +63,7 @@ __device float2 direction_to_equirectangular(float3 dir) return make_float2(u, v); } -__device float3 equirectangular_to_direction(float u, float v) +ccl_device float3 equirectangular_to_direction(float u, float v) { float phi = M_PI_F*(1.0f - 2.0f*u); float theta = M_PI_F*(1.0f - v); @@ -76,7 +76,7 @@ __device float3 equirectangular_to_direction(float u, float v) /* Fisheye <-> Cartesian direction */ -__device float2 direction_to_fisheye(float3 dir, float fov) +ccl_device float2 direction_to_fisheye(float3 dir, float fov) { float r = atan2f(sqrtf(dir.y*dir.y + dir.z*dir.z), dir.x) / fov; float phi = atan2f(dir.z, dir.y); @@ -87,7 +87,7 @@ __device float2 direction_to_fisheye(float3 dir, float fov) return make_float2(u, v); } -__device float3 fisheye_to_direction(float u, float v, float fov) +ccl_device float3 fisheye_to_direction(float u, float v, float fov) { u = (u - 0.5f) * 2.0f; v = (v - 0.5f) * 2.0f; @@ -109,7 +109,7 @@ __device float3 fisheye_to_direction(float u, float v, float fov) ); } -__device float2 direction_to_fisheye_equisolid(float3 dir, float lens, float width, float height) +ccl_device float2 direction_to_fisheye_equisolid(float3 dir, float lens, float width, float height) { float theta = acosf(dir.x); float r = 2.0f * lens * sinf(theta * 0.5f); @@ -121,7 +121,7 @@ __device float2 direction_to_fisheye_equisolid(float3 dir, float lens, float wid return make_float2(u, v); } -__device float3 fisheye_equisolid_to_direction(float u, float v, float lens, float fov, float width, float height) +ccl_device float3 fisheye_equisolid_to_direction(float u, float v, float lens, float fov, float width, float height) { u = (u - 0.5f) * width; v = (v - 0.5f) * height; @@ -146,7 +146,7 @@ __device float3 fisheye_equisolid_to_direction(float u, float v, float lens, flo /* Mirror Ball <-> Cartesion direction */ -__device float3 mirrorball_to_direction(float u, float v) +ccl_device float3 mirrorball_to_direction(float u, float v) { /* point on sphere */ float3 dir; @@ -161,7 +161,7 @@ __device float3 mirrorball_to_direction(float u, float v) return 2.0f*dot(dir, I)*dir - I; } -__device float2 direction_to_mirrorball(float3 dir) +ccl_device float2 direction_to_mirrorball(float3 dir) { /* inverse of mirrorball_to_direction */ dir.y -= 1.0f; @@ -176,7 +176,7 @@ __device float2 direction_to_mirrorball(float3 dir) return make_float2(u, v); } -__device float3 panorama_to_direction(KernelGlobals *kg, float u, float v) +ccl_device float3 panorama_to_direction(KernelGlobals *kg, float u, float v) { switch(kernel_data.cam.panorama_type) { case PANORAMA_EQUIRECTANGULAR: @@ -190,7 +190,7 @@ __device float3 panorama_to_direction(KernelGlobals *kg, float u, float v) } } -__device float2 direction_to_panorama(KernelGlobals *kg, float3 dir) +ccl_device float2 direction_to_panorama(KernelGlobals *kg, float3 dir) { switch(kernel_data.cam.panorama_type) { case PANORAMA_EQUIRECTANGULAR: diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h index dc977a8780f..69e7b439e1c 100644 --- a/intern/cycles/kernel/kernel_random.h +++ b/intern/cycles/kernel/kernel_random.h @@ -30,7 +30,7 @@ typedef uint RNG; /* High Dimensional Sobol */ /* van der corput radical inverse */ -__device uint van_der_corput(uint bits) +ccl_device uint van_der_corput(uint bits) { bits = (bits << 16) | (bits >> 16); bits = ((bits & 0x00ff00ff) << 8) | ((bits & 0xff00ff00) >> 8); @@ -41,7 +41,7 @@ __device uint van_der_corput(uint bits) } /* sobol radical inverse */ -__device uint sobol(uint i) +ccl_device uint sobol(uint i) { uint r = 0; @@ -53,7 +53,7 @@ __device uint sobol(uint i) } /* inverse of sobol radical inverse */ -__device uint sobol_inverse(uint i) +ccl_device uint sobol_inverse(uint i) { const uint msb = 1U << 31; uint r = 0; @@ -67,7 +67,7 @@ __device uint sobol_inverse(uint i) /* multidimensional sobol with generator matrices * dimension 0 and 1 are equal to van_der_corput() and sobol() respectively */ -__device uint sobol_dimension(KernelGlobals *kg, int index, int dimension) +ccl_device uint sobol_dimension(KernelGlobals *kg, int index, int dimension) { uint result = 0; uint i = index; @@ -80,7 +80,7 @@ __device uint sobol_dimension(KernelGlobals *kg, int index, int dimension) } /* lookup index and x/y coordinate, assumes m is a power of two */ -__device uint sobol_lookup(const uint m, const uint frame, const uint ex, const uint ey, uint *x, uint *y) +ccl_device uint sobol_lookup(const uint m, const uint frame, const uint ex, const uint ey, uint *x, uint *y) { /* shift is constant per frame */ const uint shift = frame << (m << 1); @@ -100,7 +100,7 @@ __device uint sobol_lookup(const uint m, const uint frame, const uint ex, const return index; } -__device_inline float path_rng_1D(KernelGlobals *kg, RNG *rng, int sample, int num_samples, int dimension) +ccl_device_inline float path_rng_1D(KernelGlobals *kg, RNG *rng, int sample, int num_samples, int dimension) { #ifdef __CMJ__ if(kernel_data.integrator.sampling_pattern == SAMPLING_PATTERN_CMJ) { @@ -131,7 +131,7 @@ __device_inline float path_rng_1D(KernelGlobals *kg, RNG *rng, int sample, int n #endif } -__device_inline void path_rng_2D(KernelGlobals *kg, RNG *rng, int sample, int num_samples, int dimension, float *fx, float *fy) +ccl_device_inline void path_rng_2D(KernelGlobals *kg, RNG *rng, int sample, int num_samples, int dimension, float *fx, float *fy) { #ifdef __CMJ__ if(kernel_data.integrator.sampling_pattern == SAMPLING_PATTERN_CMJ) { @@ -148,7 +148,7 @@ __device_inline void path_rng_2D(KernelGlobals *kg, RNG *rng, int sample, int nu } } -__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, int num_samples, RNG *rng, int x, int y, float *fx, float *fy) +ccl_device_inline void path_rng_init(KernelGlobals *kg, ccl_global uint *rng_state, int sample, int num_samples, RNG *rng, int x, int y, float *fx, float *fy) { #ifdef __SOBOL_FULL_SCREEN__ uint px, py; @@ -183,7 +183,7 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, #endif } -__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) +ccl_device void path_rng_end(KernelGlobals *kg, ccl_global uint *rng_state, RNG rng) { /* nothing to do */ } @@ -192,24 +192,24 @@ __device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) /* Linear Congruential Generator */ -__device float path_rng(KernelGlobals *kg, RNG& rng, int sample, int dimension) +ccl_device float path_rng(KernelGlobals *kg, RNG& rng, int sample, int dimension) { } -__device_inline float path_rng_1D(KernelGlobals *kg, RNG& rng, int sample, int num_samples, int dimension) +ccl_device_inline float path_rng_1D(KernelGlobals *kg, RNG& rng, int sample, int num_samples, int dimension) { /* implicit mod 2^32 */ rng = (1103515245*(rng) + 12345); return (float)rng * (1.0f/(float)0xFFFFFFFF); } -__device_inline void path_rng_2D(KernelGlobals *kg, RNG& rng, int sample, int num_samples, int dimension, float *fx, float *fy) +ccl_device_inline void path_rng_2D(KernelGlobals *kg, RNG& rng, int sample, int num_samples, int dimension, float *fx, float *fy) { *fx = path_rng_1D(kg, rng, sample, num_samples, dimension); *fy = path_rng_1D(kg, rng, sample, num_samples, dimension + 1); } -__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, int num_samples, RNG *rng, int x, int y, float *fx, float *fy) +ccl_device void path_rng_init(KernelGlobals *kg, ccl_global uint *rng_state, int sample, int num_samples, RNG *rng, int x, int y, float *fx, float *fy) { /* load state */ *rng = *rng_state; @@ -225,7 +225,7 @@ __device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sam } } -__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) +ccl_device void path_rng_end(KernelGlobals *kg, ccl_global uint *rng_state, RNG rng) { /* store state for next sample */ *rng_state = rng; @@ -233,21 +233,21 @@ __device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) #endif -__device uint lcg_step_uint(uint *rng) +ccl_device uint lcg_step_uint(uint *rng) { /* implicit mod 2^32 */ *rng = (1103515245*(*rng) + 12345); return *rng; } -__device float lcg_step_float(uint *rng) +ccl_device float lcg_step_float(uint *rng) { /* implicit mod 2^32 */ *rng = (1103515245*(*rng) + 12345); return (float)*rng * (1.0f/(float)0xFFFFFFFF); } -__device uint lcg_init(uint seed) +ccl_device uint lcg_init(uint seed) { uint rng = seed; lcg_step_uint(&rng); diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index 81630caed9a..77154ce3aef 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -36,7 +36,7 @@ CCL_NAMESPACE_BEGIN /* ShaderData setup from incoming ray */ #ifdef __OBJECT_MOTION__ -__device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, float time) +ccl_device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, float time) { if(sd->flag & SD_OBJECT_MOTION) { sd->ob_tfm = object_fetch_transform_motion(kg, sd->object, time); @@ -49,7 +49,7 @@ __device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, } #endif -__device void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, +ccl_device void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, int bounce) { #ifdef __INSTANCING__ @@ -161,7 +161,7 @@ __device void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, /* ShaderData setup from BSSRDF scatter */ #ifdef __SUBSURFACE__ -__device_inline void shader_setup_from_subsurface(KernelGlobals *kg, ShaderData *sd, +ccl_device_inline void shader_setup_from_subsurface(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray) { bool backfacing = sd->flag & SD_BACKFACING; @@ -237,7 +237,7 @@ __device_inline void shader_setup_from_subsurface(KernelGlobals *kg, ShaderData /* ShaderData setup from position sampled on mesh */ -__device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, +ccl_device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, const float3 P, const float3 Ng, const float3 I, int shader, int object, int prim, float u, float v, float t, float time, int bounce, int segment) { @@ -357,7 +357,7 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, /* ShaderData setup for displacement */ -__device void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd, +ccl_device void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd, int object, int prim, float u, float v) { float3 P, Ng, I = make_float3(0.0f, 0.0f, 0.0f); @@ -376,7 +376,7 @@ __device void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd, /* ShaderData setup from ray into background */ -__device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray, int bounce) +ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray, int bounce) { /* vectors */ sd->P = ray->D; @@ -426,7 +426,7 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData #ifdef __MULTI_CLOSURE__ -__device_inline void _shader_bsdf_multi_eval(KernelGlobals *kg, const ShaderData *sd, const float3 omega_in, float *pdf, +ccl_device_inline void _shader_bsdf_multi_eval(KernelGlobals *kg, const ShaderData *sd, const float3 omega_in, float *pdf, int skip_bsdf, BsdfEval *result_eval, float sum_pdf, float sum_sample_weight) { /* this is the veach one-sample model with balance heuristic, some pdf @@ -455,7 +455,7 @@ __device_inline void _shader_bsdf_multi_eval(KernelGlobals *kg, const ShaderData #endif -__device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, +ccl_device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, const float3 omega_in, BsdfEval *eval, float *pdf) { #ifdef __MULTI_CLOSURE__ @@ -470,7 +470,7 @@ __device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, #endif } -__device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, +ccl_device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, float randu, float randv, BsdfEval *bsdf_eval, float3 *omega_in, differential3 *domega_in, float *pdf) { @@ -534,7 +534,7 @@ __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, #endif } -__device int shader_bsdf_sample_closure(KernelGlobals *kg, const ShaderData *sd, +ccl_device int shader_bsdf_sample_closure(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc, float randu, float randv, BsdfEval *bsdf_eval, float3 *omega_in, differential3 *domega_in, float *pdf) { @@ -550,7 +550,7 @@ __device int shader_bsdf_sample_closure(KernelGlobals *kg, const ShaderData *sd, return label; } -__device void shader_bsdf_blur(KernelGlobals *kg, ShaderData *sd, float roughness) +ccl_device void shader_bsdf_blur(KernelGlobals *kg, ShaderData *sd, float roughness) { #ifdef __MULTI_CLOSURE__ for(int i = 0; i< sd->num_closure; i++) { @@ -564,7 +564,7 @@ __device void shader_bsdf_blur(KernelGlobals *kg, ShaderData *sd, float roughnes #endif } -__device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -585,7 +585,7 @@ __device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd) #endif } -__device float3 shader_bsdf_alpha(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_bsdf_alpha(KernelGlobals *kg, ShaderData *sd) { float3 alpha = make_float3(1.0f, 1.0f, 1.0f) - shader_bsdf_transparency(kg, sd); @@ -595,7 +595,7 @@ __device float3 shader_bsdf_alpha(KernelGlobals *kg, ShaderData *sd) return alpha; } -__device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -616,7 +616,7 @@ __device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd) #endif } -__device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -637,7 +637,7 @@ __device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd) #endif } -__device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -658,7 +658,7 @@ __device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd) #endif } -__device float3 shader_bsdf_subsurface(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_bsdf_subsurface(KernelGlobals *kg, ShaderData *sd) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -679,7 +679,7 @@ __device float3 shader_bsdf_subsurface(KernelGlobals *kg, ShaderData *sd) #endif } -__device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_factor, float3 *N_) +ccl_device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_factor, float3 *N_) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -717,7 +717,7 @@ __device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_facto #endif } -__device float3 shader_bssrdf_sum(ShaderData *sd, float3 *N_, float *texture_blur_) +ccl_device float3 shader_bssrdf_sum(ShaderData *sd, float3 *N_, float *texture_blur_) { #ifdef __MULTI_CLOSURE__ float3 eval = make_float3(0.0f, 0.0f, 0.0f); @@ -762,7 +762,7 @@ __device float3 shader_bssrdf_sum(ShaderData *sd, float3 *N_, float *texture_blu /* Emission */ -__device float3 emissive_eval(KernelGlobals *kg, ShaderData *sd, ShaderClosure *sc) +ccl_device float3 emissive_eval(KernelGlobals *kg, ShaderData *sd, ShaderClosure *sc) { #ifdef __OSL__ if(kg->osl && sc->prim) @@ -772,7 +772,7 @@ __device float3 emissive_eval(KernelGlobals *kg, ShaderData *sd, ShaderClosure * return emissive_simple_eval(sd->Ng, sd->I); } -__device float3 shader_emissive_eval(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_emissive_eval(KernelGlobals *kg, ShaderData *sd) { float3 eval; #ifdef __MULTI_CLOSURE__ @@ -793,7 +793,7 @@ __device float3 shader_emissive_eval(KernelGlobals *kg, ShaderData *sd) /* Holdout */ -__device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd) +ccl_device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd) { #ifdef __MULTI_CLOSURE__ float3 weight = make_float3(0.0f, 0.0f, 0.0f); @@ -816,7 +816,7 @@ __device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd) /* Surface Evaluation */ -__device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd, +ccl_device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd, float randb, int path_flag, ShaderContext ctx) { #ifdef __OSL__ @@ -837,7 +837,7 @@ __device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd, /* Background Evaluation */ -__device float3 shader_eval_background(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx) +ccl_device float3 shader_eval_background(KernelGlobals *kg, ShaderData *sd, int path_flag, ShaderContext ctx) { #ifdef __OSL__ if (kg->osl) @@ -875,7 +875,7 @@ __device float3 shader_eval_background(KernelGlobals *kg, ShaderData *sd, int pa /* Volume */ -__device float3 shader_volume_eval_phase(KernelGlobals *kg, ShaderData *sd, +ccl_device float3 shader_volume_eval_phase(KernelGlobals *kg, ShaderData *sd, float3 omega_in, float3 omega_out) { #ifdef __MULTI_CLOSURE__ @@ -896,7 +896,7 @@ __device float3 shader_volume_eval_phase(KernelGlobals *kg, ShaderData *sd, /* Volume Evaluation */ -__device void shader_eval_volume(KernelGlobals *kg, ShaderData *sd, +ccl_device void shader_eval_volume(KernelGlobals *kg, ShaderData *sd, float randb, int path_flag, ShaderContext ctx) { #ifdef __SVM__ @@ -911,7 +911,7 @@ __device void shader_eval_volume(KernelGlobals *kg, ShaderData *sd, /* Displacement Evaluation */ -__device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, ShaderContext ctx) +ccl_device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, ShaderContext ctx) { /* this will modify sd->P */ #ifdef __SVM__ @@ -927,7 +927,7 @@ __device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, Shader /* Transparent Shadows */ #ifdef __TRANSPARENT_SHADOWS__ -__device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect) +ccl_device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect) { int prim = kernel_tex_fetch(__prim_index, isect->prim); int shader = 0; @@ -953,7 +953,7 @@ __device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect) /* Merging */ #ifdef __BRANCHED_PATH__ -__device void shader_merge_closures(KernelGlobals *kg, ShaderData *sd) +ccl_device void shader_merge_closures(KernelGlobals *kg, ShaderData *sd) { /* merge identical closures, better when we sample a single closure at a time */ for(int i = 0; i < sd->num_closure; i++) { diff --git a/intern/cycles/kernel/kernel_subsurface.h b/intern/cycles/kernel/kernel_subsurface.h index d16b9328bf2..2326ca18b55 100644 --- a/intern/cycles/kernel/kernel_subsurface.h +++ b/intern/cycles/kernel/kernel_subsurface.h @@ -25,7 +25,7 @@ CCL_NAMESPACE_BEGIN #define BSSRDF_MULTI_EVAL -__device ShaderClosure *subsurface_scatter_pick_closure(KernelGlobals *kg, ShaderData *sd, float *probability) +ccl_device ShaderClosure *subsurface_scatter_pick_closure(KernelGlobals *kg, ShaderData *sd, float *probability) { /* sum sample weights of bssrdf and bsdf */ float bsdf_sum = 0.0f; @@ -80,7 +80,7 @@ __device ShaderClosure *subsurface_scatter_pick_closure(KernelGlobals *kg, Shade return NULL; } -__device float3 subsurface_scatter_eval(ShaderData *sd, ShaderClosure *sc, float disk_r, float r, bool all) +ccl_device float3 subsurface_scatter_eval(ShaderData *sd, ShaderClosure *sc, float disk_r, float r, bool all) { #ifdef BSSRDF_MULTI_EVAL /* this is the veach one-sample model with balance heuristic, some pdf @@ -133,7 +133,7 @@ __device float3 subsurface_scatter_eval(ShaderData *sd, ShaderClosure *sc, float } /* replace closures with a single diffuse bsdf closure after scatter step */ -__device void subsurface_scatter_setup_diffuse_bsdf(ShaderData *sd, float3 weight, bool hit, float3 N) +ccl_device void subsurface_scatter_setup_diffuse_bsdf(ShaderData *sd, float3 weight, bool hit, float3 N) { sd->flag &= ~SD_CLOSURE_FLAGS; sd->randb_closure = 0.0f; @@ -158,7 +158,7 @@ __device void subsurface_scatter_setup_diffuse_bsdf(ShaderData *sd, float3 weigh } /* optionally do blurring of color and/or bump mapping, at the cost of a shader evaluation */ -__device float3 subsurface_color_pow(float3 color, float exponent) +ccl_device float3 subsurface_color_pow(float3 color, float exponent) { color = max(color, make_float3(0.0f, 0.0f, 0.0f)); @@ -179,7 +179,7 @@ __device float3 subsurface_color_pow(float3 color, float exponent) return color; } -__device void subsurface_color_bump_blur(KernelGlobals *kg, ShaderData *out_sd, ShaderData *in_sd, int state_flag, float3 *eval, float3 *N) +ccl_device void subsurface_color_bump_blur(KernelGlobals *kg, ShaderData *out_sd, ShaderData *in_sd, int state_flag, float3 *eval, float3 *N) { /* average color and texture blur at outgoing point */ float texture_blur; @@ -207,7 +207,7 @@ __device void subsurface_color_bump_blur(KernelGlobals *kg, ShaderData *out_sd, } /* subsurface scattering step, from a point on the surface to other nearby points on the same object */ -__device int subsurface_scatter_multi_step(KernelGlobals *kg, ShaderData *sd, ShaderData bssrdf_sd[BSSRDF_MAX_HITS], +ccl_device int subsurface_scatter_multi_step(KernelGlobals *kg, ShaderData *sd, ShaderData bssrdf_sd[BSSRDF_MAX_HITS], int state_flag, ShaderClosure *sc, uint *lcg_state, float disk_u, float disk_v, bool all) { /* pick random axis in local frame and point on disk */ @@ -313,7 +313,7 @@ __device int subsurface_scatter_multi_step(KernelGlobals *kg, ShaderData *sd, Sh } /* subsurface scattering step, from a point on the surface to another nearby point on the same object */ -__device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd, +ccl_device void subsurface_scatter_step(KernelGlobals *kg, ShaderData *sd, int state_flag, ShaderClosure *sc, uint *lcg_state, float disk_u, float disk_v, bool all) { float3 eval = make_float3(0.0f, 0.0f, 0.0f); diff --git a/intern/cycles/kernel/kernel_triangle.h b/intern/cycles/kernel/kernel_triangle.h index 71389e0ec32..d457b67e77e 100644 --- a/intern/cycles/kernel/kernel_triangle.h +++ b/intern/cycles/kernel/kernel_triangle.h @@ -17,7 +17,7 @@ CCL_NAMESPACE_BEGIN /* Point on triangle for Moller-Trumbore triangles */ -__device_inline float3 triangle_point_MT(KernelGlobals *kg, int tri_index, float u, float v) +ccl_device_inline float3 triangle_point_MT(KernelGlobals *kg, int tri_index, float u, float v) { /* load triangle vertices */ float3 tri_vindex = float4_to_float3(kernel_tex_fetch(__tri_vindex, tri_index)); @@ -32,7 +32,7 @@ __device_inline float3 triangle_point_MT(KernelGlobals *kg, int tri_index, float } /* Sample point on triangle */ -__device_inline float3 triangle_sample_MT(KernelGlobals *kg, int tri_index, float randu, float randv) +ccl_device_inline float3 triangle_sample_MT(KernelGlobals *kg, int tri_index, float randu, float randv) { /* compute point */ randu = sqrtf(randu); @@ -44,7 +44,7 @@ __device_inline float3 triangle_sample_MT(KernelGlobals *kg, int tri_index, floa } /* Normal for Moller-Trumbore triangles */ -__device_inline float3 triangle_normal_MT(KernelGlobals *kg, int tri_index, int *shader) +ccl_device_inline float3 triangle_normal_MT(KernelGlobals *kg, int tri_index, int *shader) { #if 0 /* load triangle vertices */ @@ -64,7 +64,7 @@ __device_inline float3 triangle_normal_MT(KernelGlobals *kg, int tri_index, int } /* Return 3 triangle vertex locations */ -__device_inline void triangle_vertices(KernelGlobals *kg, int tri_index, float3 P[3]) +ccl_device_inline void triangle_vertices(KernelGlobals *kg, int tri_index, float3 P[3]) { /* load triangle vertices */ float3 tri_vindex = float4_to_float3(kernel_tex_fetch(__tri_vindex, tri_index)); @@ -74,7 +74,7 @@ __device_inline void triangle_vertices(KernelGlobals *kg, int tri_index, float3 P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, __float_as_int(tri_vindex.z))); } -__device_inline float3 triangle_smooth_normal(KernelGlobals *kg, int tri_index, float u, float v) +ccl_device_inline float3 triangle_smooth_normal(KernelGlobals *kg, int tri_index, float u, float v) { /* load triangle vertices */ float3 tri_vindex = float4_to_float3(kernel_tex_fetch(__tri_vindex, tri_index)); @@ -86,7 +86,7 @@ __device_inline float3 triangle_smooth_normal(KernelGlobals *kg, int tri_index, return normalize((1.0f - u - v)*n2 + u*n0 + v*n1); } -__device_inline void triangle_dPdudv(KernelGlobals *kg, float3 *dPdu, float3 *dPdv, int tri) +ccl_device_inline void triangle_dPdudv(KernelGlobals *kg, float3 *dPdu, float3 *dPdv, int tri) { /* fetch triangle vertex coordinates */ float3 tri_vindex = float4_to_float3(kernel_tex_fetch(__tri_vindex, tri)); @@ -102,7 +102,7 @@ __device_inline void triangle_dPdudv(KernelGlobals *kg, float3 *dPdu, float3 *dP /* attributes */ -__device float triangle_attribute_float(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float *dx, float *dy) +ccl_device float triangle_attribute_float(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float *dx, float *dy) { if(elem == ATTR_ELEMENT_FACE) { if(dx) *dx = 0.0f; @@ -145,7 +145,7 @@ __device float triangle_attribute_float(KernelGlobals *kg, const ShaderData *sd, } } -__device float3 triangle_attribute_float3(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float3 *dx, float3 *dy) +ccl_device float3 triangle_attribute_float3(KernelGlobals *kg, const ShaderData *sd, AttributeElement elem, int offset, float3 *dx, float3 *dy) { if(elem == ATTR_ELEMENT_FACE) { if(dx) *dx = make_float3(0.0f, 0.0f, 0.0f); diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h index 168127c620c..b2be9deb938 100644 --- a/intern/cycles/kernel/svm/svm.h +++ b/intern/cycles/kernel/svm/svm.h @@ -45,14 +45,14 @@ CCL_NAMESPACE_BEGIN /* Stack */ -__device_inline float3 stack_load_float3(float *stack, uint a) +ccl_device_inline float3 stack_load_float3(float *stack, uint a) { kernel_assert(a+2 < SVM_STACK_SIZE); return make_float3(stack[a+0], stack[a+1], stack[a+2]); } -__device_inline void stack_store_float3(float *stack, uint a, float3 f) +ccl_device_inline void stack_store_float3(float *stack, uint a, float3 f) { kernel_assert(a+2 < SVM_STACK_SIZE); @@ -61,59 +61,59 @@ __device_inline void stack_store_float3(float *stack, uint a, float3 f) stack[a+2] = f.z; } -__device_inline float stack_load_float(float *stack, uint a) +ccl_device_inline float stack_load_float(float *stack, uint a) { kernel_assert(a < SVM_STACK_SIZE); return stack[a]; } -__device_inline float stack_load_float_default(float *stack, uint a, uint value) +ccl_device_inline float stack_load_float_default(float *stack, uint a, uint value) { return (a == (uint)SVM_STACK_INVALID)? __uint_as_float(value): stack_load_float(stack, a); } -__device_inline void stack_store_float(float *stack, uint a, float f) +ccl_device_inline void stack_store_float(float *stack, uint a, float f) { kernel_assert(a < SVM_STACK_SIZE); stack[a] = f; } -__device_inline int stack_load_int(float *stack, uint a) +ccl_device_inline int stack_load_int(float *stack, uint a) { kernel_assert(a < SVM_STACK_SIZE); return __float_as_int(stack[a]); } -__device_inline float stack_load_int_default(float *stack, uint a, uint value) +ccl_device_inline float stack_load_int_default(float *stack, uint a, uint value) { return (a == (uint)SVM_STACK_INVALID)? (int)value: stack_load_int(stack, a); } -__device_inline void stack_store_int(float *stack, uint a, int i) +ccl_device_inline void stack_store_int(float *stack, uint a, int i) { kernel_assert(a < SVM_STACK_SIZE); stack[a] = __int_as_float(i); } -__device_inline bool stack_valid(uint a) +ccl_device_inline bool stack_valid(uint a) { return a != (uint)SVM_STACK_INVALID; } /* Reading Nodes */ -__device_inline uint4 read_node(KernelGlobals *kg, int *offset) +ccl_device_inline uint4 read_node(KernelGlobals *kg, int *offset) { uint4 node = kernel_tex_fetch(__svm_nodes, *offset); (*offset)++; return node; } -__device_inline float4 read_node_float(KernelGlobals *kg, int *offset) +ccl_device_inline float4 read_node_float(KernelGlobals *kg, int *offset) { uint4 node = kernel_tex_fetch(__svm_nodes, *offset); float4 f = make_float4(__uint_as_float(node.x), __uint_as_float(node.y), __uint_as_float(node.z), __uint_as_float(node.w)); @@ -121,13 +121,13 @@ __device_inline float4 read_node_float(KernelGlobals *kg, int *offset) return f; } -__device_inline float4 fetch_node_float(KernelGlobals *kg, int offset) +ccl_device_inline float4 fetch_node_float(KernelGlobals *kg, int offset) { uint4 node = kernel_tex_fetch(__svm_nodes, offset); return make_float4(__uint_as_float(node.x), __uint_as_float(node.y), __uint_as_float(node.z), __uint_as_float(node.w)); } -__device_inline void decode_node_uchar4(uint i, uint *x, uint *y, uint *z, uint *w) +ccl_device_inline void decode_node_uchar4(uint i, uint *x, uint *y, uint *z, uint *w) { if(x) *x = (i & 0xFF); if(y) *y = ((i >> 8) & 0xFF); @@ -182,7 +182,7 @@ CCL_NAMESPACE_BEGIN /* Main Interpreter Loop */ -__device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag) +ccl_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; diff --git a/intern/cycles/kernel/svm/svm_attribute.h b/intern/cycles/kernel/svm/svm_attribute.h index 8e71e7cdd56..90409e16477 100644 --- a/intern/cycles/kernel/svm/svm_attribute.h +++ b/intern/cycles/kernel/svm/svm_attribute.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Attribute Node */ -__device void svm_node_attr_init(KernelGlobals *kg, ShaderData *sd, +ccl_device void svm_node_attr_init(KernelGlobals *kg, ShaderData *sd, uint4 node, NodeAttributeType *type, NodeAttributeType *mesh_type, AttributeElement *elem, int *offset, uint *out_offset) { @@ -52,7 +52,7 @@ __device void svm_node_attr_init(KernelGlobals *kg, ShaderData *sd, *type = (NodeAttributeType)node.w; } -__device void svm_node_attr(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_attr(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { NodeAttributeType type, mesh_type; AttributeElement elem; @@ -84,7 +84,7 @@ __device void svm_node_attr(KernelGlobals *kg, ShaderData *sd, float *stack, uin } } -__device void svm_node_attr_bump_dx(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_attr_bump_dx(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { NodeAttributeType type, mesh_type; AttributeElement elem; @@ -120,7 +120,7 @@ __device void svm_node_attr_bump_dx(KernelGlobals *kg, ShaderData *sd, float *st } } -__device void svm_node_attr_bump_dy(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_attr_bump_dy(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { NodeAttributeType type, mesh_type; AttributeElement elem; diff --git a/intern/cycles/kernel/svm/svm_blackbody.h b/intern/cycles/kernel/svm/svm_blackbody.h index 2fc2c770a83..63dbf27d35e 100644 --- a/intern/cycles/kernel/svm/svm_blackbody.h +++ b/intern/cycles/kernel/svm/svm_blackbody.h @@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN /* Blackbody Node */ -__device void svm_node_blackbody(KernelGlobals *kg, ShaderData *sd, float *stack, uint temperature_offset, uint col_offset) +ccl_device void svm_node_blackbody(KernelGlobals *kg, ShaderData *sd, float *stack, uint temperature_offset, uint col_offset) { /* Output */ float3 color_rgb = make_float3(0.0f, 0.0f, 0.0f); diff --git a/intern/cycles/kernel/svm/svm_brick.h b/intern/cycles/kernel/svm/svm_brick.h index 19b4b5e779f..7cac922d8a6 100644 --- a/intern/cycles/kernel/svm/svm_brick.h +++ b/intern/cycles/kernel/svm/svm_brick.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Brick */ -__device_noinline float brick_noise(int n) /* fast integer noise */ +ccl_device_noinline float brick_noise(int n) /* fast integer noise */ { int nn; n = (n >> 13) ^ n; @@ -26,7 +26,7 @@ __device_noinline float brick_noise(int n) /* fast integer noise */ return 0.5f * ((float)nn / 1073741824.0f); } -__device_noinline float2 svm_brick(float3 p, float scale, float mortar_size, float bias, +ccl_device_noinline float2 svm_brick(float3 p, float scale, float mortar_size, float bias, float brick_width, float row_height, float offset_amount, int offset_frequency, float squash_amount, int squash_frequency) { @@ -56,7 +56,7 @@ __device_noinline float2 svm_brick(float3 p, float scale, float mortar_size, flo y > (row_height - mortar_size)) ? 1.0f : 0.0f); } -__device void svm_node_tex_brick(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_brick(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint4 node2 = read_node(kg, offset); uint4 node3 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_brightness.h b/intern/cycles/kernel/svm/svm_brightness.h index 3e977dcbe1b..9b330b3213f 100644 --- a/intern/cycles/kernel/svm/svm_brightness.h +++ b/intern/cycles/kernel/svm/svm_brightness.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_brightness(ShaderData *sd, float *stack, uint in_color, uint out_color, uint node) +ccl_device void svm_node_brightness(ShaderData *sd, float *stack, uint in_color, uint out_color, uint node) { uint bright_offset, contrast_offset; float3 color = stack_load_float3(stack, in_color); diff --git a/intern/cycles/kernel/svm/svm_camera.h b/intern/cycles/kernel/svm/svm_camera.h index 76f50e196eb..bfe9289fa02 100644 --- a/intern/cycles/kernel/svm/svm_camera.h +++ b/intern/cycles/kernel/svm/svm_camera.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_camera(KernelGlobals *kg, ShaderData *sd, float *stack, uint out_vector, uint out_zdepth, uint out_distance) +ccl_device void svm_node_camera(KernelGlobals *kg, ShaderData *sd, float *stack, uint out_vector, uint out_zdepth, uint out_distance) { float distance; float zdepth; diff --git a/intern/cycles/kernel/svm/svm_checker.h b/intern/cycles/kernel/svm/svm_checker.h index 70fe2ac5a92..ebc48e16d68 100644 --- a/intern/cycles/kernel/svm/svm_checker.h +++ b/intern/cycles/kernel/svm/svm_checker.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Checker */ -__device_noinline float svm_checker(float3 p, float scale) +ccl_device_noinline float svm_checker(float3 p, float scale) { p *= scale; @@ -34,7 +34,7 @@ __device_noinline float svm_checker(float3 p, float scale) return ((xi % 2 == yi % 2) == (zi % 2))? 1.0f: 0.0f; } -__device void svm_node_tex_checker(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_tex_checker(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint co_offset, color1_offset, color2_offset, scale_offset; uint color_offset, fac_offset; diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h index 0d4716ab078..2c6fb5deca4 100644 --- a/intern/cycles/kernel/svm/svm_closure.h +++ b/intern/cycles/kernel/svm/svm_closure.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Closure Nodes */ -__device void svm_node_glass_setup(ShaderData *sd, ShaderClosure *sc, int type, float eta, float roughness, bool refract) +ccl_device void svm_node_glass_setup(ShaderData *sd, ShaderClosure *sc, int type, float eta, float roughness, bool refract) { if(type == CLOSURE_BSDF_SHARP_GLASS_ID) { if(refract) { @@ -49,7 +49,7 @@ __device void svm_node_glass_setup(ShaderData *sd, ShaderClosure *sc, int type, } } -__device_inline ShaderClosure *svm_node_closure_get_non_bsdf(ShaderData *sd, ClosureType type, float mix_weight) +ccl_device_inline ShaderClosure *svm_node_closure_get_non_bsdf(ShaderData *sd, ClosureType type, float mix_weight) { #ifdef __MULTI_CLOSURE__ ShaderClosure *sc = &sd->closure[sd->num_closure]; @@ -70,7 +70,7 @@ __device_inline ShaderClosure *svm_node_closure_get_non_bsdf(ShaderData *sd, Clo #endif } -__device_inline ShaderClosure *svm_node_closure_get_bsdf(ShaderData *sd, float mix_weight) +ccl_device_inline ShaderClosure *svm_node_closure_get_bsdf(ShaderData *sd, float mix_weight) { #ifdef __MULTI_CLOSURE__ ShaderClosure *sc = &sd->closure[sd->num_closure]; @@ -93,7 +93,7 @@ __device_inline ShaderClosure *svm_node_closure_get_bsdf(ShaderData *sd, float m #endif } -__device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, float randb, int path_flag, int *offset) +ccl_device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, float randb, int path_flag, int *offset) { uint type, param1_offset, param2_offset; @@ -456,7 +456,7 @@ __device void svm_node_closure_bsdf(KernelGlobals *kg, ShaderData *sd, float *st } } -__device void svm_node_closure_volume(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int path_flag) +ccl_device void svm_node_closure_volume(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int path_flag) { uint type, param1_offset, param2_offset; @@ -499,7 +499,7 @@ __device void svm_node_closure_volume(KernelGlobals *kg, ShaderData *sd, float * } } -__device void svm_node_closure_emission(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_closure_emission(ShaderData *sd, float *stack, uint4 node) { #ifdef __MULTI_CLOSURE__ uint mix_weight_offset = node.y; @@ -522,7 +522,7 @@ __device void svm_node_closure_emission(ShaderData *sd, float *stack, uint4 node sd->flag |= SD_EMISSION; } -__device void svm_node_closure_background(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_closure_background(ShaderData *sd, float *stack, uint4 node) { #ifdef __MULTI_CLOSURE__ uint mix_weight_offset = node.y; @@ -543,7 +543,7 @@ __device void svm_node_closure_background(ShaderData *sd, float *stack, uint4 no #endif } -__device void svm_node_closure_holdout(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_closure_holdout(ShaderData *sd, float *stack, uint4 node) { #ifdef __MULTI_CLOSURE__ uint mix_weight_offset = node.y; @@ -566,7 +566,7 @@ __device void svm_node_closure_holdout(ShaderData *sd, float *stack, uint4 node) sd->flag |= SD_HOLDOUT; } -__device void svm_node_closure_ambient_occlusion(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_closure_ambient_occlusion(ShaderData *sd, float *stack, uint4 node) { #ifdef __MULTI_CLOSURE__ uint mix_weight_offset = node.y; @@ -591,7 +591,7 @@ __device void svm_node_closure_ambient_occlusion(ShaderData *sd, float *stack, u /* Closure Nodes */ -__device_inline void svm_node_closure_store_weight(ShaderData *sd, float3 weight) +ccl_device_inline void svm_node_closure_store_weight(ShaderData *sd, float3 weight) { #ifdef __MULTI_CLOSURE__ if(sd->num_closure < MAX_CLOSURE) @@ -601,13 +601,13 @@ __device_inline void svm_node_closure_store_weight(ShaderData *sd, float3 weight #endif } -__device void svm_node_closure_set_weight(ShaderData *sd, uint r, uint g, uint b) +ccl_device void svm_node_closure_set_weight(ShaderData *sd, uint r, uint g, uint b) { float3 weight = make_float3(__uint_as_float(r), __uint_as_float(g), __uint_as_float(b)); svm_node_closure_store_weight(sd, weight); } -__device void svm_node_emission_set_weight_total(KernelGlobals *kg, ShaderData *sd, uint r, uint g, uint b) +ccl_device void svm_node_emission_set_weight_total(KernelGlobals *kg, ShaderData *sd, uint r, uint g, uint b) { float3 weight = make_float3(__uint_as_float(r), __uint_as_float(g), __uint_as_float(b)); @@ -617,14 +617,14 @@ __device void svm_node_emission_set_weight_total(KernelGlobals *kg, ShaderData * svm_node_closure_store_weight(sd, weight); } -__device void svm_node_closure_weight(ShaderData *sd, float *stack, uint weight_offset) +ccl_device void svm_node_closure_weight(ShaderData *sd, float *stack, uint weight_offset) { float3 weight = stack_load_float3(stack, weight_offset); svm_node_closure_store_weight(sd, weight); } -__device void svm_node_emission_weight(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_emission_weight(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint color_offset = node.y; uint strength_offset = node.z; @@ -639,7 +639,7 @@ __device void svm_node_emission_weight(KernelGlobals *kg, ShaderData *sd, float svm_node_closure_store_weight(sd, weight); } -__device void svm_node_mix_closure(ShaderData *sd, float *stack, +ccl_device void svm_node_mix_closure(ShaderData *sd, float *stack, uint4 node, int *offset, float *randb) { #ifdef __MULTI_CLOSURE__ @@ -675,7 +675,7 @@ __device void svm_node_mix_closure(ShaderData *sd, float *stack, #endif } -__device void svm_node_add_closure(ShaderData *sd, float *stack, uint unused, +ccl_device void svm_node_add_closure(ShaderData *sd, float *stack, uint unused, uint node_jump, int *offset, float *randb, float *closure_weight) { #ifdef __MULTI_CLOSURE__ @@ -699,7 +699,7 @@ __device void svm_node_add_closure(ShaderData *sd, float *stack, uint unused, /* (Bump) normal */ -__device void svm_node_set_normal(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_direction, uint out_normal) +ccl_device void svm_node_set_normal(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_direction, uint out_normal) { float3 normal = stack_load_float3(stack, in_direction); sd->N = normal; diff --git a/intern/cycles/kernel/svm/svm_convert.h b/intern/cycles/kernel/svm/svm_convert.h index 22f4651689d..2503912c5c6 100644 --- a/intern/cycles/kernel/svm/svm_convert.h +++ b/intern/cycles/kernel/svm/svm_convert.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Conversion Nodes */ -__device void svm_node_convert(ShaderData *sd, float *stack, uint type, uint from, uint to) +ccl_device void svm_node_convert(ShaderData *sd, float *stack, uint type, uint from, uint to) { switch(type) { case NODE_CONVERT_FI: { diff --git a/intern/cycles/kernel/svm/svm_displace.h b/intern/cycles/kernel/svm/svm_displace.h index d0bac647a7c..6cd5ee4b375 100644 --- a/intern/cycles/kernel/svm/svm_displace.h +++ b/intern/cycles/kernel/svm/svm_displace.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Bump Node */ -__device void svm_node_set_bump(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_set_bump(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { #ifdef __RAY_DIFFERENTIALS__ /* get normal input */ @@ -62,7 +62,7 @@ __device void svm_node_set_bump(KernelGlobals *kg, ShaderData *sd, float *stack, /* Displacement Node */ -__device void svm_node_set_displacement(ShaderData *sd, float *stack, uint fac_offset) +ccl_device void svm_node_set_displacement(ShaderData *sd, float *stack, uint fac_offset) { float d = stack_load_float(stack, fac_offset); sd->P += sd->N*d*0.1f; /* todo: get rid of this factor */ diff --git a/intern/cycles/kernel/svm/svm_fresnel.h b/intern/cycles/kernel/svm/svm_fresnel.h index d97d6a3738f..bb70a3faa2a 100644 --- a/intern/cycles/kernel/svm/svm_fresnel.h +++ b/intern/cycles/kernel/svm/svm_fresnel.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Fresnel Node */ -__device void svm_node_fresnel(ShaderData *sd, float *stack, uint ior_offset, uint ior_value, uint node) +ccl_device void svm_node_fresnel(ShaderData *sd, float *stack, uint ior_offset, uint ior_value, uint node) { uint normal_offset, out_offset; decode_node_uchar4(node, &normal_offset, &out_offset, NULL, NULL); @@ -35,7 +35,7 @@ __device void svm_node_fresnel(ShaderData *sd, float *stack, uint ior_offset, ui /* Layer Weight Node */ -__device void svm_node_layer_weight(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_layer_weight(ShaderData *sd, float *stack, uint4 node) { uint blend_offset = node.y; uint blend_value = node.z; diff --git a/intern/cycles/kernel/svm/svm_gamma.h b/intern/cycles/kernel/svm/svm_gamma.h index ef1581fba8d..c4749e7b936 100644 --- a/intern/cycles/kernel/svm/svm_gamma.h +++ b/intern/cycles/kernel/svm/svm_gamma.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_gamma(ShaderData *sd, float *stack, uint in_gamma, uint in_color, uint out_color) +ccl_device void svm_node_gamma(ShaderData *sd, float *stack, uint in_gamma, uint in_color, uint out_color) { float3 color = stack_load_float3(stack, in_color); float gamma = stack_load_float(stack, in_gamma); diff --git a/intern/cycles/kernel/svm/svm_geometry.h b/intern/cycles/kernel/svm/svm_geometry.h index 818d8694453..ad0cacb027a 100644 --- a/intern/cycles/kernel/svm/svm_geometry.h +++ b/intern/cycles/kernel/svm/svm_geometry.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Geometry Node */ -__device void svm_node_geometry(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) +ccl_device void svm_node_geometry(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) { float3 data; @@ -38,7 +38,7 @@ __device void svm_node_geometry(KernelGlobals *kg, ShaderData *sd, float *stack, stack_store_float3(stack, out_offset, data); } -__device void svm_node_geometry_bump_dx(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) +ccl_device void svm_node_geometry_bump_dx(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) { #ifdef __RAY_DIFFERENTIALS__ float3 data; @@ -55,7 +55,7 @@ __device void svm_node_geometry_bump_dx(KernelGlobals *kg, ShaderData *sd, float #endif } -__device void svm_node_geometry_bump_dy(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) +ccl_device void svm_node_geometry_bump_dy(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) { #ifdef __RAY_DIFFERENTIALS__ float3 data; @@ -74,7 +74,7 @@ __device void svm_node_geometry_bump_dy(KernelGlobals *kg, ShaderData *sd, float /* Object Info */ -__device void svm_node_object_info(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) +ccl_device void svm_node_object_info(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) { float data; @@ -94,7 +94,7 @@ __device void svm_node_object_info(KernelGlobals *kg, ShaderData *sd, float *sta /* Particle Info */ -__device void svm_node_particle_info(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) +ccl_device void svm_node_particle_info(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) { switch(type) { case NODE_INFO_PAR_INDEX: { @@ -146,7 +146,7 @@ __device void svm_node_particle_info(KernelGlobals *kg, ShaderData *sd, float *s /* Hair Info */ -__device void svm_node_hair_info(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) +ccl_device void svm_node_hair_info(KernelGlobals *kg, ShaderData *sd, float *stack, uint type, uint out_offset) { float data; float3 data3; diff --git a/intern/cycles/kernel/svm/svm_gradient.h b/intern/cycles/kernel/svm/svm_gradient.h index 1c0fe511f9b..a4b3c0583f7 100644 --- a/intern/cycles/kernel/svm/svm_gradient.h +++ b/intern/cycles/kernel/svm/svm_gradient.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Gradient */ -__device float svm_gradient(float3 p, NodeGradientType type) +ccl_device float svm_gradient(float3 p, NodeGradientType type) { float x, y, z; @@ -57,7 +57,7 @@ __device float svm_gradient(float3 p, NodeGradientType type) return 0.0f; } -__device void svm_node_tex_gradient(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_tex_gradient(ShaderData *sd, float *stack, uint4 node) { uint type, co_offset, color_offset, fac_offset; diff --git a/intern/cycles/kernel/svm/svm_hsv.h b/intern/cycles/kernel/svm/svm_hsv.h index e16fb7582c1..11dfc4f096b 100644 --- a/intern/cycles/kernel/svm/svm_hsv.h +++ b/intern/cycles/kernel/svm/svm_hsv.h @@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_hsv(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_color_offset, uint fac_offset, uint out_color_offset, int *offset) +ccl_device void svm_node_hsv(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_color_offset, uint fac_offset, uint out_color_offset, int *offset) { /* read extra data */ uint4 node1 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_image.h b/intern/cycles/kernel/svm/svm_image.h index e18fe7c53a7..58e5775265a 100644 --- a/intern/cycles/kernel/svm/svm_image.h +++ b/intern/cycles/kernel/svm/svm_image.h @@ -21,14 +21,14 @@ CCL_NAMESPACE_BEGIN /* For OpenCL all images are packed in a single array, and we do manual lookup * and interpolation. */ -__device_inline float4 svm_image_texture_read(KernelGlobals *kg, int offset) +ccl_device_inline float4 svm_image_texture_read(KernelGlobals *kg, int offset) { uchar4 r = kernel_tex_fetch(__tex_image_packed, offset); float f = 1.0f/255.0f; return make_float4(r.x*f, r.y*f, r.z*f, r.w*f); } -__device_inline int svm_image_texture_wrap_periodic(int x, int width) +ccl_device_inline int svm_image_texture_wrap_periodic(int x, int width) { x %= width; if(x < 0) @@ -36,19 +36,19 @@ __device_inline int svm_image_texture_wrap_periodic(int x, int width) return x; } -__device_inline int svm_image_texture_wrap_clamp(int x, int width) +ccl_device_inline int svm_image_texture_wrap_clamp(int x, int width) { return clamp(x, 0, width-1); } -__device_inline float svm_image_texture_frac(float x, int *ix) +ccl_device_inline float svm_image_texture_frac(float x, int *ix) { int i = float_to_int(x) - ((x < 0.0f)? 1: 0); *ix = i; return x - (float)i; } -__device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, uint srgb, uint use_alpha) +ccl_device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, uint srgb, uint use_alpha) { /* first slots are used by float textures, which are not supported here */ if(id < TEX_NUM_FLOAT_IMAGES) @@ -110,7 +110,7 @@ __device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, u #else -__device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, uint srgb, uint use_alpha) +ccl_device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, uint srgb, uint use_alpha) { float4 r; @@ -257,7 +257,7 @@ __device float4 svm_image_texture(KernelGlobals *kg, int id, float x, float y, u #endif -__device void svm_node_tex_image(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_tex_image(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint id = node.y; uint co_offset, out_offset, alpha_offset, srgb; @@ -274,7 +274,7 @@ __device void svm_node_tex_image(KernelGlobals *kg, ShaderData *sd, float *stack stack_store_float(stack, alpha_offset, f.w); } -__device void svm_node_tex_image_box(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_tex_image_box(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { /* get object space normal */ float3 N = sd->N; @@ -363,7 +363,7 @@ __device void svm_node_tex_image_box(KernelGlobals *kg, ShaderData *sd, float *s } -__device void svm_node_tex_environment(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_tex_environment(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint id = node.y; uint co_offset, out_offset, alpha_offset, srgb; diff --git a/intern/cycles/kernel/svm/svm_invert.h b/intern/cycles/kernel/svm/svm_invert.h index 4c40afeadd9..eb47e9ad4ab 100644 --- a/intern/cycles/kernel/svm/svm_invert.h +++ b/intern/cycles/kernel/svm/svm_invert.h @@ -16,12 +16,12 @@ CCL_NAMESPACE_BEGIN -__device float invert(float color, float factor) +ccl_device float invert(float color, float factor) { return factor*(1.0f - color) + (1.0f - factor) * color; } -__device void svm_node_invert(ShaderData *sd, float *stack, uint in_fac, uint in_color, uint out_color) +ccl_device void svm_node_invert(ShaderData *sd, float *stack, uint in_fac, uint in_color, uint out_color) { float factor = stack_load_float(stack, in_fac); float3 color = stack_load_float3(stack, in_color); diff --git a/intern/cycles/kernel/svm/svm_light_path.h b/intern/cycles/kernel/svm/svm_light_path.h index ff6776d751e..e7afa2d2200 100644 --- a/intern/cycles/kernel/svm/svm_light_path.h +++ b/intern/cycles/kernel/svm/svm_light_path.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Light Path Node */ -__device void svm_node_light_path(ShaderData *sd, float *stack, uint type, uint out_offset, int path_flag) +ccl_device void svm_node_light_path(ShaderData *sd, float *stack, uint type, uint out_offset, int path_flag) { float info = 0.0f; @@ -40,7 +40,7 @@ __device void svm_node_light_path(ShaderData *sd, float *stack, uint type, uint /* Light Falloff Node */ -__device void svm_node_light_falloff(ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_light_falloff(ShaderData *sd, float *stack, uint4 node) { uint strength_offset, out_offset, smooth_offset; diff --git a/intern/cycles/kernel/svm/svm_magic.h b/intern/cycles/kernel/svm/svm_magic.h index 7a5eba3f564..b661f5cacf8 100644 --- a/intern/cycles/kernel/svm/svm_magic.h +++ b/intern/cycles/kernel/svm/svm_magic.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Magic */ -__device_noinline float3 svm_magic(float3 p, int n, float distortion) +ccl_device_noinline float3 svm_magic(float3 p, int n, float distortion) { float x = sinf((p.x + p.y + p.z)*5.0f); float y = cosf((-p.x + p.y - p.z)*5.0f); @@ -87,7 +87,7 @@ __device_noinline float3 svm_magic(float3 p, int n, float distortion) return make_float3(0.5f - x, 0.5f - y, 0.5f - z); } -__device void svm_node_tex_magic(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_magic(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint depth; uint scale_offset, distortion_offset, co_offset, fac_offset, color_offset; diff --git a/intern/cycles/kernel/svm/svm_mapping.h b/intern/cycles/kernel/svm/svm_mapping.h index fcdd92dd575..c9fa8502dd1 100644 --- a/intern/cycles/kernel/svm/svm_mapping.h +++ b/intern/cycles/kernel/svm/svm_mapping.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Mapping Node */ -__device void svm_node_mapping(KernelGlobals *kg, ShaderData *sd, float *stack, uint vec_offset, uint out_offset, int *offset) +ccl_device void svm_node_mapping(KernelGlobals *kg, ShaderData *sd, float *stack, uint vec_offset, uint out_offset, int *offset) { float3 v = stack_load_float3(stack, vec_offset); @@ -32,7 +32,7 @@ __device void svm_node_mapping(KernelGlobals *kg, ShaderData *sd, float *stack, stack_store_float3(stack, out_offset, r); } -__device void svm_node_min_max(KernelGlobals *kg, ShaderData *sd, float *stack, uint vec_offset, uint out_offset, int *offset) +ccl_device void svm_node_min_max(KernelGlobals *kg, ShaderData *sd, float *stack, uint vec_offset, uint out_offset, int *offset) { float3 v = stack_load_float3(stack, vec_offset); diff --git a/intern/cycles/kernel/svm/svm_math.h b/intern/cycles/kernel/svm/svm_math.h index d4863dd6216..bb46d443a6b 100644 --- a/intern/cycles/kernel/svm/svm_math.h +++ b/intern/cycles/kernel/svm/svm_math.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device float svm_math(NodeMath type, float Fac1, float Fac2) +ccl_device float svm_math(NodeMath type, float Fac1, float Fac2) { float Fac; @@ -64,12 +64,12 @@ __device float svm_math(NodeMath type, float Fac1, float Fac2) return Fac; } -__device float average_fac(float3 v) +ccl_device float average_fac(float3 v) { return (fabsf(v.x) + fabsf(v.y) + fabsf(v.z))/3.0f; } -__device void svm_vector_math(float *Fac, float3 *Vector, NodeVectorMath type, float3 Vector1, float3 Vector2) +ccl_device void svm_vector_math(float *Fac, float3 *Vector, NodeVectorMath type, float3 Vector1, float3 Vector2) { if(type == NODE_VECTOR_MATH_ADD) { *Vector = Vector1 + Vector2; @@ -104,7 +104,7 @@ __device void svm_vector_math(float *Fac, float3 *Vector, NodeVectorMath type, f /* Nodes */ -__device void svm_node_math(KernelGlobals *kg, ShaderData *sd, float *stack, uint itype, uint f1_offset, uint f2_offset, int *offset) +ccl_device void svm_node_math(KernelGlobals *kg, ShaderData *sd, float *stack, uint itype, uint f1_offset, uint f2_offset, int *offset) { NodeMath type = (NodeMath)itype; float f1 = stack_load_float(stack, f1_offset); @@ -116,7 +116,7 @@ __device void svm_node_math(KernelGlobals *kg, ShaderData *sd, float *stack, uin stack_store_float(stack, node1.y, f); } -__device void svm_node_vector_math(KernelGlobals *kg, ShaderData *sd, float *stack, uint itype, uint v1_offset, uint v2_offset, int *offset) +ccl_device void svm_node_vector_math(KernelGlobals *kg, ShaderData *sd, float *stack, uint itype, uint v1_offset, uint v2_offset, int *offset) { NodeVectorMath type = (NodeVectorMath)itype; float3 v1 = stack_load_float3(stack, v1_offset); diff --git a/intern/cycles/kernel/svm/svm_mix.h b/intern/cycles/kernel/svm/svm_mix.h index 506f772dba5..0eeb4cf9b05 100644 --- a/intern/cycles/kernel/svm/svm_mix.h +++ b/intern/cycles/kernel/svm/svm_mix.h @@ -16,22 +16,22 @@ CCL_NAMESPACE_BEGIN -__device float3 svm_mix_blend(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_blend(float t, float3 col1, float3 col2) { return interp(col1, col2, t); } -__device float3 svm_mix_add(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_add(float t, float3 col1, float3 col2) { return interp(col1, col1 + col2, t); } -__device float3 svm_mix_mul(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_mul(float t, float3 col1, float3 col2) { return interp(col1, col1 * col2, t); } -__device float3 svm_mix_screen(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_screen(float t, float3 col1, float3 col2) { float tm = 1.0f - t; float3 one = make_float3(1.0f, 1.0f, 1.0f); @@ -40,7 +40,7 @@ __device float3 svm_mix_screen(float t, float3 col1, float3 col2) return one - (tm3 + t*(one - col2))*(one - col1); } -__device float3 svm_mix_overlay(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_overlay(float t, float3 col1, float3 col2) { float tm = 1.0f - t; @@ -64,12 +64,12 @@ __device float3 svm_mix_overlay(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_sub(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_sub(float t, float3 col1, float3 col2) { return interp(col1, col1 - col2, t); } -__device float3 svm_mix_div(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_div(float t, float3 col1, float3 col2) { float tm = 1.0f - t; @@ -82,22 +82,22 @@ __device float3 svm_mix_div(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_diff(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_diff(float t, float3 col1, float3 col2) { return interp(col1, fabs(col1 - col2), t); } -__device float3 svm_mix_dark(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_dark(float t, float3 col1, float3 col2) { return min(col1, col2*t); } -__device float3 svm_mix_light(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_light(float t, float3 col1, float3 col2) { return max(col1, col2*t); } -__device float3 svm_mix_dodge(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_dodge(float t, float3 col1, float3 col2) { float3 outcol = col1; @@ -132,7 +132,7 @@ __device float3 svm_mix_dodge(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_burn(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_burn(float t, float3 col1, float3 col2) { float tmp, tm = 1.0f - t; @@ -171,7 +171,7 @@ __device float3 svm_mix_burn(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_hue(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_hue(float t, float3 col1, float3 col2) { float3 outcol = col1; @@ -188,7 +188,7 @@ __device float3 svm_mix_hue(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_sat(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_sat(float t, float3 col1, float3 col2) { float tm = 1.0f - t; @@ -206,7 +206,7 @@ __device float3 svm_mix_sat(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_val(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_val(float t, float3 col1, float3 col2) { float tm = 1.0f - t; @@ -218,7 +218,7 @@ __device float3 svm_mix_val(float t, float3 col1, float3 col2) return hsv_to_rgb(hsv); } -__device float3 svm_mix_color(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_color(float t, float3 col1, float3 col2) { float3 outcol = col1; float3 hsv2 = rgb_to_hsv(col2); @@ -235,7 +235,7 @@ __device float3 svm_mix_color(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_soft(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_soft(float t, float3 col1, float3 col2) { float tm = 1.0f - t; @@ -245,7 +245,7 @@ __device float3 svm_mix_soft(float t, float3 col1, float3 col2) return tm*col1 + t*((one - col1)*col2*col1 + col1*scr); } -__device float3 svm_mix_linear(float t, float3 col1, float3 col2) +ccl_device float3 svm_mix_linear(float t, float3 col1, float3 col2) { float3 outcol = col1; @@ -267,7 +267,7 @@ __device float3 svm_mix_linear(float t, float3 col1, float3 col2) return outcol; } -__device float3 svm_mix_clamp(float3 col) +ccl_device float3 svm_mix_clamp(float3 col) { float3 outcol = col; @@ -278,7 +278,7 @@ __device float3 svm_mix_clamp(float3 col) return outcol; } -__device float3 svm_mix(NodeMix type, float fac, float3 c1, float3 c2) +ccl_device float3 svm_mix(NodeMix type, float fac, float3 c1, float3 c2) { float t = clamp(fac, 0.0f, 1.0f); @@ -309,7 +309,7 @@ __device float3 svm_mix(NodeMix type, float fac, float3 c1, float3 c2) /* Node */ -__device void svm_node_mix(KernelGlobals *kg, ShaderData *sd, float *stack, uint fac_offset, uint c1_offset, uint c2_offset, int *offset) +ccl_device void svm_node_mix(KernelGlobals *kg, ShaderData *sd, float *stack, uint fac_offset, uint c1_offset, uint c2_offset, int *offset) { /* read extra data */ uint4 node1 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_musgrave.h b/intern/cycles/kernel/svm/svm_musgrave.h index 65dcf1a9f83..c67dc8297e4 100644 --- a/intern/cycles/kernel/svm/svm_musgrave.h +++ b/intern/cycles/kernel/svm/svm_musgrave.h @@ -25,7 +25,7 @@ CCL_NAMESPACE_BEGIN * from "Texturing and Modelling: A procedural approach" */ -__device_noinline float noise_musgrave_fBm(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves) +ccl_device_noinline float noise_musgrave_fBm(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves) { float rmd; float value = 0.0f; @@ -53,7 +53,7 @@ __device_noinline float noise_musgrave_fBm(float3 p, NodeNoiseBasis basis, float * octaves: number of frequencies in the fBm */ -__device_noinline float noise_musgrave_multi_fractal(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves) +ccl_device_noinline float noise_musgrave_multi_fractal(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves) { float rmd; float value = 1.0f; @@ -82,7 +82,7 @@ __device_noinline float noise_musgrave_multi_fractal(float3 p, NodeNoiseBasis ba * offset: raises the terrain from `sea level' */ -__device_noinline float noise_musgrave_hetero_terrain(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves, float offset) +ccl_device_noinline float noise_musgrave_hetero_terrain(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves, float offset) { float value, increment, rmd; float pwHL = powf(lacunarity, -H); @@ -117,7 +117,7 @@ __device_noinline float noise_musgrave_hetero_terrain(float3 p, NodeNoiseBasis b * offset: raises the terrain from `sea level' */ -__device_noinline float noise_musgrave_hybrid_multi_fractal(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves, float offset, float gain) +ccl_device_noinline float noise_musgrave_hybrid_multi_fractal(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves, float offset, float gain) { float result, signal, weight, rmd; float pwHL = powf(lacunarity, -H); @@ -154,7 +154,7 @@ __device_noinline float noise_musgrave_hybrid_multi_fractal(float3 p, NodeNoiseB * offset: raises the terrain from `sea level' */ -__device_noinline float noise_musgrave_ridged_multi_fractal(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves, float offset, float gain) +ccl_device_noinline float noise_musgrave_ridged_multi_fractal(float3 p, NodeNoiseBasis basis, float H, float lacunarity, float octaves, float offset, float gain) { float result, signal, weight; float pwHL = powf(lacunarity, -H); @@ -181,7 +181,7 @@ __device_noinline float noise_musgrave_ridged_multi_fractal(float3 p, NodeNoiseB /* Shader */ -__device float svm_musgrave(NodeMusgraveType type, float dimension, float lacunarity, float octaves, float offset, float intensity, float gain, float scale, float3 p) +ccl_device float svm_musgrave(NodeMusgraveType type, float dimension, float lacunarity, float octaves, float offset, float intensity, float gain, float scale, float3 p) { NodeNoiseBasis basis = NODE_NOISE_PERLIN; p *= scale; @@ -200,7 +200,7 @@ __device float svm_musgrave(NodeMusgraveType type, float dimension, float lacuna return 0.0f; } -__device void svm_node_tex_musgrave(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_musgrave(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint4 node2 = read_node(kg, offset); uint4 node3 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_noise.h b/intern/cycles/kernel/svm/svm_noise.h index a55c635b679..2055b0e3ec7 100644 --- a/intern/cycles/kernel/svm/svm_noise.h +++ b/intern/cycles/kernel/svm/svm_noise.h @@ -32,17 +32,17 @@ CCL_NAMESPACE_BEGIN -__device int quick_floor(float x) +ccl_device int quick_floor(float x) { return float_to_int(x) - ((x < 0) ? 1 : 0); } -__device float bits_to_01(uint bits) +ccl_device float bits_to_01(uint bits) { return bits * (1.0f/(float)0xFFFFFFFF); } -__device uint hash(uint kx, uint ky, uint kz) +ccl_device uint hash(uint kx, uint ky, uint kz) { // define some handy macros #define rot(x,k) (((x)<<(k)) | ((x)>>(32-(k)))) @@ -71,34 +71,34 @@ __device uint hash(uint kx, uint ky, uint kz) #undef final } -__device int imod(int a, int b) +ccl_device int imod(int a, int b) { a %= b; return a < 0 ? a + b : a; } -__device uint phash(int kx, int ky, int kz, int3 p) +ccl_device uint phash(int kx, int ky, int kz, int3 p) { return hash(imod(kx, p.x), imod(ky, p.y), imod(kz, p.z)); } -__device float floorfrac(float x, int* i) +ccl_device float floorfrac(float x, int* i) { *i = quick_floor(x); return x - *i; } -__device float fade(float t) +ccl_device float fade(float t) { return t * t * t * (t * (t * 6.0f - 15.0f) + 10.0f); } -__device float nerp(float t, float a, float b) +ccl_device float nerp(float t, float a, float b) { return (1.0f - t) * a + t * b; } -__device float grad(int hash, float x, float y, float z) +ccl_device float grad(int hash, float x, float y, float z) { // use vectors pointing to the edges of the cube int h = hash & 15; @@ -107,12 +107,12 @@ __device float grad(int hash, float x, float y, float z) return ((h&1) ? -u : u) + ((h&2) ? -v : v); } -__device float scale3(float result) +ccl_device float scale3(float result) { return 0.9820f * result; } -__device_noinline float perlin(float x, float y, float z) +ccl_device_noinline float perlin(float x, float y, float z) { int X; float fx = floorfrac(x, &X); int Y; float fy = floorfrac(y, &Y); @@ -138,7 +138,7 @@ __device_noinline float perlin(float x, float y, float z) return (isfinite(r))? r: 0.0f; } -__device_noinline float perlin_periodic(float x, float y, float z, float3 pperiod) +ccl_device_noinline float perlin_periodic(float x, float y, float z, float3 pperiod) { int X; float fx = floorfrac(x, &X); int Y; float fy = floorfrac(y, &Y); @@ -171,20 +171,20 @@ __device_noinline float perlin_periodic(float x, float y, float z, float3 pperio } /* perlin noise in range 0..1 */ -__device float noise(float3 p) +ccl_device float noise(float3 p) { float r = perlin(p.x, p.y, p.z); return 0.5f*r + 0.5f; } /* perlin noise in range -1..1 */ -__device float snoise(float3 p) +ccl_device float snoise(float3 p) { return perlin(p.x, p.y, p.z); } /* cell noise */ -__device_noinline float cellnoise(float3 p) +ccl_device_noinline float cellnoise(float3 p) { uint ix = quick_floor(p.x); uint iy = quick_floor(p.y); @@ -193,7 +193,7 @@ __device_noinline float cellnoise(float3 p) return bits_to_01(hash(ix, iy, iz)); } -__device float3 cellnoise_color(float3 p) +ccl_device float3 cellnoise_color(float3 p) { float r = cellnoise(p); float g = cellnoise(make_float3(p.y, p.x, p.z)); @@ -203,14 +203,14 @@ __device float3 cellnoise_color(float3 p) } /* periodic perlin noise in range 0..1 */ -__device float pnoise(float3 p, float3 pperiod) +ccl_device float pnoise(float3 p, float3 pperiod) { float r = perlin_periodic(p.x, p.y, p.z, pperiod); return 0.5f*r + 0.5f; } /* periodic perlin noise in range -1..1 */ -__device float psnoise(float3 p, float3 pperiod) +ccl_device float psnoise(float3 p, float3 pperiod) { return perlin_periodic(p.x, p.y, p.z, pperiod); } diff --git a/intern/cycles/kernel/svm/svm_noisetex.h b/intern/cycles/kernel/svm/svm_noisetex.h index acb3f20246e..02583131704 100644 --- a/intern/cycles/kernel/svm/svm_noisetex.h +++ b/intern/cycles/kernel/svm/svm_noisetex.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Noise */ -__device_inline void svm_noise(float3 p, float scale, float detail, float distortion, float *fac, float3 *color) +ccl_device_inline void svm_noise(float3 p, float scale, float detail, float distortion, float *fac, float3 *color) { NodeNoiseBasis basis = NODE_NOISE_PERLIN; int hard = 0; @@ -41,7 +41,7 @@ __device_inline void svm_noise(float3 p, float scale, float detail, float distor noise_turbulence(make_float3(p.y, p.z, p.x), basis, detail, hard)); } -__device void svm_node_tex_noise(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_noise(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint co_offset, scale_offset, detail_offset, distortion_offset, fac_offset, color_offset; diff --git a/intern/cycles/kernel/svm/svm_normal.h b/intern/cycles/kernel/svm/svm_normal.h index dd7506bb5fc..8695031b8b9 100644 --- a/intern/cycles/kernel/svm/svm_normal.h +++ b/intern/cycles/kernel/svm/svm_normal.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_normal(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_normal_offset, uint out_normal_offset, uint out_dot_offset, int *offset) +ccl_device void svm_node_normal(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_normal_offset, uint out_normal_offset, uint out_dot_offset, int *offset) { /* read extra data */ uint4 node1 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_ramp.h b/intern/cycles/kernel/svm/svm_ramp.h index 3cb23a2b2dd..55eee3d24c3 100644 --- a/intern/cycles/kernel/svm/svm_ramp.h +++ b/intern/cycles/kernel/svm/svm_ramp.h @@ -19,7 +19,7 @@ CCL_NAMESPACE_BEGIN -__device float4 rgb_ramp_lookup(KernelGlobals *kg, int offset, float f, bool interpolate) +ccl_device float4 rgb_ramp_lookup(KernelGlobals *kg, int offset, float f, bool interpolate) { f = clamp(f, 0.0f, 1.0f)*(RAMP_TABLE_SIZE-1); @@ -35,7 +35,7 @@ __device float4 rgb_ramp_lookup(KernelGlobals *kg, int offset, float f, bool int return a; } -__device void svm_node_rgb_ramp(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_rgb_ramp(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint fac_offset, color_offset, alpha_offset; uint interpolate = node.z; @@ -53,7 +53,7 @@ __device void svm_node_rgb_ramp(KernelGlobals *kg, ShaderData *sd, float *stack, *offset += RAMP_TABLE_SIZE; } -__device void svm_node_rgb_curves(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_rgb_curves(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint fac_offset = node.y; uint color_offset = node.z; @@ -72,7 +72,7 @@ __device void svm_node_rgb_curves(KernelGlobals *kg, ShaderData *sd, float *stac *offset += RAMP_TABLE_SIZE; } -__device void svm_node_vector_curves(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_vector_curves(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint fac_offset = node.y; uint color_offset = node.z; diff --git a/intern/cycles/kernel/svm/svm_sepcomb_hsv.h b/intern/cycles/kernel/svm/svm_sepcomb_hsv.h index 130890fdc8e..0f68ecbea03 100644 --- a/intern/cycles/kernel/svm/svm_sepcomb_hsv.h +++ b/intern/cycles/kernel/svm/svm_sepcomb_hsv.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_combine_hsv(KernelGlobals *kg, ShaderData *sd, float *stack, uint hue_in, uint saturation_in, uint value_in, int *offset) +ccl_device void svm_node_combine_hsv(KernelGlobals *kg, ShaderData *sd, float *stack, uint hue_in, uint saturation_in, uint value_in, int *offset) { uint4 node1 = read_node(kg, offset); uint color_out = node1.y; @@ -32,7 +32,7 @@ __device void svm_node_combine_hsv(KernelGlobals *kg, ShaderData *sd, float *sta stack_store_float3(stack, color_out, color); } -__device void svm_node_separate_hsv(KernelGlobals *kg, ShaderData *sd, float *stack, uint color_in, uint hue_out, uint saturation_out, int *offset) +ccl_device void svm_node_separate_hsv(KernelGlobals *kg, ShaderData *sd, float *stack, uint color_in, uint hue_out, uint saturation_out, int *offset) { uint4 node1 = read_node(kg, offset); uint value_out = node1.y; diff --git a/intern/cycles/kernel/svm/svm_sepcomb_rgb.h b/intern/cycles/kernel/svm/svm_sepcomb_rgb.h index 5c3d95435f2..34c4449ecdb 100644 --- a/intern/cycles/kernel/svm/svm_sepcomb_rgb.h +++ b/intern/cycles/kernel/svm/svm_sepcomb_rgb.h @@ -16,7 +16,7 @@ CCL_NAMESPACE_BEGIN -__device void svm_node_combine_rgb(ShaderData *sd, float *stack, uint in_offset, uint color_index, uint out_offset) +ccl_device void svm_node_combine_rgb(ShaderData *sd, float *stack, uint in_offset, uint color_index, uint out_offset) { float color = stack_load_float(stack, in_offset); @@ -24,7 +24,7 @@ __device void svm_node_combine_rgb(ShaderData *sd, float *stack, uint in_offset, stack_store_float(stack, out_offset+color_index, color); } -__device void svm_node_separate_rgb(ShaderData *sd, float *stack, uint icolor_offset, uint color_index, uint out_offset) +ccl_device void svm_node_separate_rgb(ShaderData *sd, float *stack, uint icolor_offset, uint color_index, uint out_offset) { float3 color = stack_load_float3(stack, icolor_offset); diff --git a/intern/cycles/kernel/svm/svm_sky.h b/intern/cycles/kernel/svm/svm_sky.h index 81b5f1a201f..1e3552647bd 100644 --- a/intern/cycles/kernel/svm/svm_sky.h +++ b/intern/cycles/kernel/svm/svm_sky.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Sky texture */ -__device float sky_angle_between(float thetav, float phiv, float theta, float phi) +ccl_device float sky_angle_between(float thetav, float phiv, float theta, float phi) { float cospsi = sinf(thetav)*sinf(theta)*cosf(phi - phiv) + cosf(thetav)*cosf(theta); return safe_acosf(cospsi); @@ -28,7 +28,7 @@ __device float sky_angle_between(float thetav, float phiv, float theta, float ph * "A Practical Analytic Model for Daylight" * A. J. Preetham, Peter Shirley, Brian Smits */ -__device float sky_perez_function(float *lam, float theta, float gamma) +ccl_device float sky_perez_function(float *lam, float theta, float gamma) { float ctheta = cosf(theta); float cgamma = cosf(gamma); @@ -36,7 +36,7 @@ __device float sky_perez_function(float *lam, float theta, float gamma) return (1.0f + lam[0]*expf(lam[1]/ctheta)) * (1.0f + lam[2]*expf(lam[3]*gamma) + lam[4]*cgamma*cgamma); } -__device float3 sky_radiance_old(KernelGlobals *kg, float3 dir, +ccl_device float3 sky_radiance_old(KernelGlobals *kg, float3 dir, float sunphi, float suntheta, float radiance_x, float radiance_y, float radiance_z, float *config_x, float *config_y, float *config_z) @@ -66,7 +66,7 @@ __device float3 sky_radiance_old(KernelGlobals *kg, float3 dir, * "An Analytic Model for Full Spectral Sky-Dome Radiance" * Lukas Hosek, Alexander Wilkie */ -__device float sky_radiance_internal(float *configuration, float theta, float gamma) +ccl_device float sky_radiance_internal(float *configuration, float theta, float gamma) { float ctheta = cosf(theta); float cgamma = cosf(gamma); @@ -80,7 +80,7 @@ __device float sky_radiance_internal(float *configuration, float theta, float ga (configuration[2] + configuration[3] * expM + configuration[5] * rayM + configuration[6] * mieM + configuration[7] * zenith); } -__device float3 sky_radiance_new(KernelGlobals *kg, float3 dir, +ccl_device float3 sky_radiance_new(KernelGlobals *kg, float3 dir, float sunphi, float suntheta, float radiance_x, float radiance_y, float radiance_z, float *config_x, float *config_y, float *config_z) @@ -105,7 +105,7 @@ __device float3 sky_radiance_new(KernelGlobals *kg, float3 dir, return xyz_to_rgb(x, y, z) * (M_2PI_F/683); } -__device void svm_node_tex_sky(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_sky(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { /* Define variables */ float sunphi, suntheta, radiance_x, radiance_y, radiance_z; diff --git a/intern/cycles/kernel/svm/svm_tex_coord.h b/intern/cycles/kernel/svm/svm_tex_coord.h index 9f88389fcb1..3044cbf81e0 100644 --- a/intern/cycles/kernel/svm/svm_tex_coord.h +++ b/intern/cycles/kernel/svm/svm_tex_coord.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Texture Coordinate Node */ -__device void svm_node_tex_coord(KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint type, uint out_offset) +ccl_device void svm_node_tex_coord(KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint type, uint out_offset) { float3 data; @@ -78,7 +78,7 @@ __device void svm_node_tex_coord(KernelGlobals *kg, ShaderData *sd, int path_fla stack_store_float3(stack, out_offset, data); } -__device void svm_node_tex_coord_bump_dx(KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint type, uint out_offset) +ccl_device void svm_node_tex_coord_bump_dx(KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint type, uint out_offset) { #ifdef __RAY_DIFFERENTIALS__ float3 data; @@ -142,7 +142,7 @@ __device void svm_node_tex_coord_bump_dx(KernelGlobals *kg, ShaderData *sd, int #endif } -__device void svm_node_tex_coord_bump_dy(KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint type, uint out_offset) +ccl_device void svm_node_tex_coord_bump_dy(KernelGlobals *kg, ShaderData *sd, int path_flag, float *stack, uint type, uint out_offset) { #ifdef __RAY_DIFFERENTIALS__ float3 data; @@ -206,7 +206,7 @@ __device void svm_node_tex_coord_bump_dy(KernelGlobals *kg, ShaderData *sd, int #endif } -__device void svm_node_normal_map(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_normal_map(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint color_offset, strength_offset, normal_offset, space; decode_node_uchar4(node.y, &color_offset, &strength_offset, &normal_offset, &space); @@ -280,7 +280,7 @@ __device void svm_node_normal_map(KernelGlobals *kg, ShaderData *sd, float *stac stack_store_float3(stack, normal_offset, N); } -__device void svm_node_tangent(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_tangent(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint tangent_offset, direction_type, axis; decode_node_uchar4(node.y, &tangent_offset, &direction_type, &axis, NULL); diff --git a/intern/cycles/kernel/svm/svm_texture.h b/intern/cycles/kernel/svm/svm_texture.h index 7f3e09a481d..8ced8390b0b 100644 --- a/intern/cycles/kernel/svm/svm_texture.h +++ b/intern/cycles/kernel/svm/svm_texture.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Voronoi Distances */ -__device float voronoi_distance(NodeDistanceMetric distance_metric, float3 d, float e) +ccl_device float voronoi_distance(NodeDistanceMetric distance_metric, float3 d, float e) { #if 0 if(distance_metric == NODE_VORONOI_DISTANCE_SQUARED) @@ -44,7 +44,7 @@ __device float voronoi_distance(NodeDistanceMetric distance_metric, float3 d, fl /* Voronoi / Worley like */ -__device_noinline float4 voronoi_Fn(float3 p, float e, int n1, int n2) +ccl_device_noinline float4 voronoi_Fn(float3 p, float e, int n1, int n2) { float da[4]; float3 pa[4]; @@ -120,29 +120,29 @@ __device_noinline float4 voronoi_Fn(float3 p, float e, int n1, int n2) return result; } -__device float voronoi_F1(float3 p) { return voronoi_Fn(p, 0.0f, 0, -1).w; } -__device float voronoi_F2(float3 p) { return voronoi_Fn(p, 0.0f, 1, -1).w; } -__device float voronoi_F3(float3 p) { return voronoi_Fn(p, 0.0f, 2, -1).w; } -__device float voronoi_F4(float3 p) { return voronoi_Fn(p, 0.0f, 3, -1).w; } -__device float voronoi_F1F2(float3 p) { return voronoi_Fn(p, 0.0f, 0, 1).w; } +ccl_device float voronoi_F1(float3 p) { return voronoi_Fn(p, 0.0f, 0, -1).w; } +ccl_device float voronoi_F2(float3 p) { return voronoi_Fn(p, 0.0f, 1, -1).w; } +ccl_device float voronoi_F3(float3 p) { return voronoi_Fn(p, 0.0f, 2, -1).w; } +ccl_device float voronoi_F4(float3 p) { return voronoi_Fn(p, 0.0f, 3, -1).w; } +ccl_device float voronoi_F1F2(float3 p) { return voronoi_Fn(p, 0.0f, 0, 1).w; } -__device float voronoi_Cr(float3 p) +ccl_device float voronoi_Cr(float3 p) { /* crackle type pattern, just a scale/clamp of F2-F1 */ float t = 10.0f*voronoi_F1F2(p); return (t > 1.0f)? 1.0f: t; } -__device float voronoi_F1S(float3 p) { return 2.0f*voronoi_F1(p) - 1.0f; } -__device float voronoi_F2S(float3 p) { return 2.0f*voronoi_F2(p) - 1.0f; } -__device float voronoi_F3S(float3 p) { return 2.0f*voronoi_F3(p) - 1.0f; } -__device float voronoi_F4S(float3 p) { return 2.0f*voronoi_F4(p) - 1.0f; } -__device float voronoi_F1F2S(float3 p) { return 2.0f*voronoi_F1F2(p) - 1.0f; } -__device float voronoi_CrS(float3 p) { return 2.0f*voronoi_Cr(p) - 1.0f; } +ccl_device float voronoi_F1S(float3 p) { return 2.0f*voronoi_F1(p) - 1.0f; } +ccl_device float voronoi_F2S(float3 p) { return 2.0f*voronoi_F2(p) - 1.0f; } +ccl_device float voronoi_F3S(float3 p) { return 2.0f*voronoi_F3(p) - 1.0f; } +ccl_device float voronoi_F4S(float3 p) { return 2.0f*voronoi_F4(p) - 1.0f; } +ccl_device float voronoi_F1F2S(float3 p) { return 2.0f*voronoi_F1F2(p) - 1.0f; } +ccl_device float voronoi_CrS(float3 p) { return 2.0f*voronoi_Cr(p) - 1.0f; } /* Noise Bases */ -__device float noise_basis(float3 p, NodeNoiseBasis basis) +ccl_device float noise_basis(float3 p, NodeNoiseBasis basis) { /* Only Perlin enabled for now, others break CUDA compile by making kernel * too big, with compile using > 4GB, due to everything being inlined. */ @@ -173,7 +173,7 @@ __device float noise_basis(float3 p, NodeNoiseBasis basis) /* Soft/Hard Noise */ -__device float noise_basis_hard(float3 p, NodeNoiseBasis basis, int hard) +ccl_device float noise_basis_hard(float3 p, NodeNoiseBasis basis, int hard) { float t = noise_basis(p, basis); return (hard)? fabsf(2.0f*t - 1.0f): t; @@ -181,7 +181,7 @@ __device float noise_basis_hard(float3 p, NodeNoiseBasis basis, int hard) /* Turbulence */ -__device_noinline float noise_turbulence(float3 p, NodeNoiseBasis basis, float octaves, int hard) +ccl_device_noinline float noise_turbulence(float3 p, NodeNoiseBasis basis, float octaves, int hard) { float fscale = 1.0f; float amp = 1.0f; diff --git a/intern/cycles/kernel/svm/svm_value.h b/intern/cycles/kernel/svm/svm_value.h index cd5a2e0d871..7beed065288 100644 --- a/intern/cycles/kernel/svm/svm_value.h +++ b/intern/cycles/kernel/svm/svm_value.h @@ -18,12 +18,12 @@ CCL_NAMESPACE_BEGIN /* Value Nodes */ -__device void svm_node_value_f(KernelGlobals *kg, ShaderData *sd, float *stack, uint ivalue, uint out_offset) +ccl_device void svm_node_value_f(KernelGlobals *kg, ShaderData *sd, float *stack, uint ivalue, uint out_offset) { stack_store_float(stack, out_offset, __uint_as_float(ivalue)); } -__device void svm_node_value_v(KernelGlobals *kg, ShaderData *sd, float *stack, uint out_offset, int *offset) +ccl_device void svm_node_value_v(KernelGlobals *kg, ShaderData *sd, float *stack, uint out_offset, int *offset) { /* read extra data */ uint4 node1 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_vector_transform.h b/intern/cycles/kernel/svm/svm_vector_transform.h index 95ef8e3e558..1e3fc2fa03b 100644 --- a/intern/cycles/kernel/svm/svm_vector_transform.h +++ b/intern/cycles/kernel/svm/svm_vector_transform.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Vector Transform */ -__device void svm_node_vector_transform(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) +ccl_device void svm_node_vector_transform(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node) { uint itype, ifrom, ito; uint vector_in, vector_out; diff --git a/intern/cycles/kernel/svm/svm_voronoi.h b/intern/cycles/kernel/svm/svm_voronoi.h index c9ebea2bceb..604fd3404c5 100644 --- a/intern/cycles/kernel/svm/svm_voronoi.h +++ b/intern/cycles/kernel/svm/svm_voronoi.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Voronoi */ -__device_noinline float4 svm_voronoi(NodeVoronoiColoring coloring, float scale, float3 p) +ccl_device_noinline float4 svm_voronoi(NodeVoronoiColoring coloring, float scale, float3 p) { /* compute distance and point coordinate of 4 nearest neighbours */ float4 dpa0 = voronoi_Fn(p*scale, 1.0f, 0, -1); @@ -39,7 +39,7 @@ __device_noinline float4 svm_voronoi(NodeVoronoiColoring coloring, float scale, return make_float4(color.x, color.y, color.z, fac); } -__device void svm_node_tex_voronoi(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_voronoi(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint coloring = node.y; uint scale_offset, co_offset, fac_offset, color_offset; diff --git a/intern/cycles/kernel/svm/svm_wave.h b/intern/cycles/kernel/svm/svm_wave.h index d906266bcf9..3749135f8c7 100644 --- a/intern/cycles/kernel/svm/svm_wave.h +++ b/intern/cycles/kernel/svm/svm_wave.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN /* Wave */ -__device_noinline float svm_wave(NodeWaveType type, float3 p, float scale, float detail, float distortion, float dscale) +ccl_device_noinline float svm_wave(NodeWaveType type, float3 p, float scale, float detail, float distortion, float dscale) { float n; @@ -35,7 +35,7 @@ __device_noinline float svm_wave(NodeWaveType type, float3 p, float scale, float return 0.5f + 0.5f * sinf(n); } -__device void svm_node_tex_wave(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) +ccl_device void svm_node_tex_wave(KernelGlobals *kg, ShaderData *sd, float *stack, uint4 node, int *offset) { uint4 node2 = read_node(kg, offset); diff --git a/intern/cycles/kernel/svm/svm_wavelength.h b/intern/cycles/kernel/svm/svm_wavelength.h index f9dd24dacef..dca4003b89a 100644 --- a/intern/cycles/kernel/svm/svm_wavelength.h +++ b/intern/cycles/kernel/svm/svm_wavelength.h @@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN /* Wavelength to RGB */ -__device void svm_node_wavelength(ShaderData *sd, float *stack, uint wavelength, uint color_out) +ccl_device void svm_node_wavelength(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 diff --git a/intern/cycles/kernel/svm/svm_wireframe.h b/intern/cycles/kernel/svm/svm_wireframe.h index 9ecb81847d7..e560e6303cc 100644 --- a/intern/cycles/kernel/svm/svm_wireframe.h +++ b/intern/cycles/kernel/svm/svm_wireframe.h @@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN /* Wireframe Node */ -__device void svm_node_wireframe(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_size, uint out_fac, uint use_pixel_size) +ccl_device void svm_node_wireframe(KernelGlobals *kg, ShaderData *sd, float *stack, uint in_size, uint out_fac, uint use_pixel_size) { /* Input Data */ float size = stack_load_float(stack, in_size); diff --git a/intern/cycles/util/util_color.h b/intern/cycles/util/util_color.h index 9bdd5b29a8e..8b13a006673 100644 --- a/intern/cycles/util/util_color.h +++ b/intern/cycles/util/util_color.h @@ -22,7 +22,7 @@ CCL_NAMESPACE_BEGIN -__device float color_srgb_to_scene_linear(float c) +ccl_device float color_srgb_to_scene_linear(float c) { if(c < 0.04045f) return (c < 0.0f)? 0.0f: c * (1.0f/12.92f); @@ -30,7 +30,7 @@ __device float color_srgb_to_scene_linear(float c) return powf((c + 0.055f) * (1.0f / 1.055f), 2.4f); } -__device float color_scene_linear_to_srgb(float c) +ccl_device float color_scene_linear_to_srgb(float c) { if(c < 0.0031308f) return (c < 0.0f)? 0.0f: c * 12.92f; @@ -38,7 +38,7 @@ __device float color_scene_linear_to_srgb(float c) return 1.055f * powf(c, 1.0f / 2.4f) - 0.055f; } -__device float3 rgb_to_hsv(float3 rgb) +ccl_device float3 rgb_to_hsv(float3 rgb) { float cmax, cmin, h, s, v, cdelta; float3 c; @@ -77,7 +77,7 @@ __device float3 rgb_to_hsv(float3 rgb) return make_float3(h, s, v); } -__device float3 hsv_to_rgb(float3 hsv) +ccl_device float3 hsv_to_rgb(float3 hsv) { float i, f, p, q, t, h, s, v; float3 rgb; @@ -112,7 +112,7 @@ __device float3 hsv_to_rgb(float3 hsv) return rgb; } -__device float3 xyY_to_xyz(float x, float y, float Y) +ccl_device float3 xyY_to_xyz(float x, float y, float Y) { float X, Z; @@ -125,7 +125,7 @@ __device float3 xyY_to_xyz(float x, float y, float Y) return make_float3(X, Y, Z); } -__device float3 xyz_to_rgb(float x, float y, float z) +ccl_device float3 xyz_to_rgb(float x, float y, float z) { return make_float3(3.240479f * x + -1.537150f * y + -0.498535f * z, -0.969256f * x + 1.875991f * y + 0.041556f * z, @@ -134,7 +134,7 @@ __device float3 xyz_to_rgb(float x, float y, float z) #ifndef __KERNEL_OPENCL__ -__device float3 color_srgb_to_scene_linear(float3 c) +ccl_device float3 color_srgb_to_scene_linear(float3 c) { return make_float3( color_srgb_to_scene_linear(c.x), @@ -142,7 +142,7 @@ __device float3 color_srgb_to_scene_linear(float3 c) color_srgb_to_scene_linear(c.z)); } -__device float3 color_scene_linear_to_srgb(float3 c) +ccl_device float3 color_scene_linear_to_srgb(float3 c) { return make_float3( color_scene_linear_to_srgb(c.x), @@ -152,7 +152,7 @@ __device float3 color_scene_linear_to_srgb(float3 c) #endif -__device float linear_rgb_to_gray(float3 c) +ccl_device float linear_rgb_to_gray(float3 c) { return c.x*0.2126f + c.y*0.7152f + c.z*0.0722f; } diff --git a/intern/cycles/util/util_math.h b/intern/cycles/util/util_math.h index 987ef4f17ac..6db532faf74 100644 --- a/intern/cycles/util/util_math.h +++ b/intern/cycles/util/util_math.h @@ -89,12 +89,12 @@ CCL_NAMESPACE_BEGIN #ifndef __KERNEL_OPENCL__ -__device_inline float fmaxf(float a, float b) +ccl_device_inline float fmaxf(float a, float b) { return (a > b)? a: b; } -__device_inline float fminf(float a, float b) +ccl_device_inline float fminf(float a, float b) { return (a < b)? a: b; } @@ -105,63 +105,63 @@ __device_inline float fminf(float a, float b) #ifndef __KERNEL_GPU__ -__device_inline int max(int a, int b) +ccl_device_inline int max(int a, int b) { return (a > b)? a: b; } -__device_inline int min(int a, int b) +ccl_device_inline int min(int a, int b) { return (a < b)? a: b; } -__device_inline float max(float a, float b) +ccl_device_inline float max(float a, float b) { return (a > b)? a: b; } -__device_inline float min(float a, float b) +ccl_device_inline float min(float a, float b) { return (a < b)? a: b; } -__device_inline double max(double a, double b) +ccl_device_inline double max(double a, double b) { return (a > b)? a: b; } -__device_inline double min(double a, double b) +ccl_device_inline double min(double a, double b) { return (a < b)? a: b; } #endif -__device_inline float min4(float a, float b, float c, float d) +ccl_device_inline float min4(float a, float b, float c, float d) { return min(min(a, b), min(c, d)); } -__device_inline float max4(float a, float b, float c, float d) +ccl_device_inline float max4(float a, float b, float c, float d) { return max(max(a, b), max(c, d)); } #ifndef __KERNEL_OPENCL__ -__device_inline int clamp(int a, int mn, int mx) +ccl_device_inline int clamp(int a, int mn, int mx) { return min(max(a, mn), mx); } -__device_inline float clamp(float a, float mn, float mx) +ccl_device_inline float clamp(float a, float mn, float mx) { return min(max(a, mn), mx); } #endif -__device_inline int float_to_int(float f) +ccl_device_inline int float_to_int(float f) { #if defined(__KERNEL_SSE2__) && !defined(_MSC_VER) return _mm_cvtt_ss2si(_mm_load_ss(&f)); @@ -170,22 +170,22 @@ __device_inline int float_to_int(float f) #endif } -__device_inline int floor_to_int(float f) +ccl_device_inline int floor_to_int(float f) { return float_to_int(floorf(f)); } -__device_inline int ceil_to_int(float f) +ccl_device_inline int ceil_to_int(float f) { return float_to_int(ceilf(f)); } -__device_inline float signf(float f) +ccl_device_inline float signf(float f) { return (f < 0.0f)? -1.0f: 1.0f; } -__device_inline float nonzerof(float f, float eps) +ccl_device_inline float nonzerof(float f, float eps) { if(fabsf(f) < eps) return signf(f)*eps; @@ -193,7 +193,7 @@ __device_inline float nonzerof(float f, float eps) return f; } -__device_inline float smoothstepf(float f) +ccl_device_inline float smoothstepf(float f) { float ff = f*f; return (3.0f*ff - 2.0f*ff*f); @@ -203,7 +203,7 @@ __device_inline float smoothstepf(float f) #ifndef __KERNEL_OPENCL__ -__device_inline bool is_zero(const float2 a) +ccl_device_inline bool is_zero(const float2 a) { return (a.x == 0.0f && a.y == 0.0f); } @@ -212,7 +212,7 @@ __device_inline bool is_zero(const float2 a) #ifndef __KERNEL_OPENCL__ -__device_inline float average(const float2 a) +ccl_device_inline float average(const float2 a) { return (a.x + a.y)*(1.0f/2.0f); } @@ -221,85 +221,85 @@ __device_inline float average(const float2 a) #ifndef __KERNEL_OPENCL__ -__device_inline float2 operator-(const float2 a) +ccl_device_inline float2 operator-(const float2 a) { return make_float2(-a.x, -a.y); } -__device_inline float2 operator*(const float2 a, const float2 b) +ccl_device_inline float2 operator*(const float2 a, const float2 b) { return make_float2(a.x*b.x, a.y*b.y); } -__device_inline float2 operator*(const float2 a, float f) +ccl_device_inline float2 operator*(const float2 a, float f) { return make_float2(a.x*f, a.y*f); } -__device_inline float2 operator*(float f, const float2 a) +ccl_device_inline float2 operator*(float f, const float2 a) { return make_float2(a.x*f, a.y*f); } -__device_inline float2 operator/(float f, const float2 a) +ccl_device_inline float2 operator/(float f, const float2 a) { return make_float2(f/a.x, f/a.y); } -__device_inline float2 operator/(const float2 a, float f) +ccl_device_inline float2 operator/(const float2 a, float f) { float invf = 1.0f/f; return make_float2(a.x*invf, a.y*invf); } -__device_inline float2 operator/(const float2 a, const float2 b) +ccl_device_inline float2 operator/(const float2 a, const float2 b) { return make_float2(a.x/b.x, a.y/b.y); } -__device_inline float2 operator+(const float2 a, const float2 b) +ccl_device_inline float2 operator+(const float2 a, const float2 b) { return make_float2(a.x+b.x, a.y+b.y); } -__device_inline float2 operator-(const float2 a, const float2 b) +ccl_device_inline float2 operator-(const float2 a, const float2 b) { return make_float2(a.x-b.x, a.y-b.y); } -__device_inline float2 operator+=(float2& a, const float2 b) +ccl_device_inline float2 operator+=(float2& a, const float2 b) { return a = a + b; } -__device_inline float2 operator*=(float2& a, const float2 b) +ccl_device_inline float2 operator*=(float2& a, const float2 b) { return a = a * b; } -__device_inline float2 operator*=(float2& a, float f) +ccl_device_inline float2 operator*=(float2& a, float f) { return a = a * f; } -__device_inline float2 operator/=(float2& a, const float2 b) +ccl_device_inline float2 operator/=(float2& a, const float2 b) { return a = a / b; } -__device_inline float2 operator/=(float2& a, float f) +ccl_device_inline float2 operator/=(float2& a, float f) { float invf = 1.0f/f; return a = a * invf; } -__device_inline float dot(const float2 a, const float2 b) +ccl_device_inline float dot(const float2 a, const float2 b) { return a.x*b.x + a.y*b.y; } -__device_inline float cross(const float2 a, const float2 b) +ccl_device_inline float cross(const float2 a, const float2 b) { return (a.x*b.y - a.y*b.x); } @@ -308,58 +308,58 @@ __device_inline float cross(const float2 a, const float2 b) #ifndef __KERNEL_OPENCL__ -__device_inline bool operator==(const int2 a, const int2 b) +ccl_device_inline bool operator==(const int2 a, const int2 b) { return (a.x == b.x && a.y == b.y); } -__device_inline float len(const float2 a) +ccl_device_inline float len(const float2 a) { return sqrtf(dot(a, a)); } -__device_inline float2 normalize(const float2 a) +ccl_device_inline float2 normalize(const float2 a) { return a/len(a); } -__device_inline float2 normalize_len(const float2 a, float *t) +ccl_device_inline float2 normalize_len(const float2 a, float *t) { *t = len(a); return a/(*t); } -__device_inline bool operator==(const float2 a, const float2 b) +ccl_device_inline bool operator==(const float2 a, const float2 b) { return (a.x == b.x && a.y == b.y); } -__device_inline bool operator!=(const float2 a, const float2 b) +ccl_device_inline bool operator!=(const float2 a, const float2 b) { return !(a == b); } -__device_inline float2 min(float2 a, float2 b) +ccl_device_inline float2 min(float2 a, float2 b) { return make_float2(min(a.x, b.x), min(a.y, b.y)); } -__device_inline float2 max(float2 a, float2 b) +ccl_device_inline float2 max(float2 a, float2 b) { return make_float2(max(a.x, b.x), max(a.y, b.y)); } -__device_inline float2 clamp(float2 a, float2 mn, float2 mx) +ccl_device_inline float2 clamp(float2 a, float2 mn, float2 mx) { return min(max(a, mn), mx); } -__device_inline float2 fabs(float2 a) +ccl_device_inline float2 fabs(float2 a) { return make_float2(fabsf(a.x), fabsf(a.y)); } -__device_inline float2 as_float2(const float4 a) +ccl_device_inline float2 as_float2(const float4 a) { return make_float2(a.x, a.y); } @@ -368,7 +368,7 @@ __device_inline float2 as_float2(const float4 a) #ifndef __KERNEL_GPU__ -__device_inline void print_float2(const char *label, const float2& a) +ccl_device_inline void print_float2(const char *label, const float2& a) { printf("%s: %.8f %.8f\n", label, (double)a.x, (double)a.y); } @@ -377,7 +377,7 @@ __device_inline void print_float2(const char *label, const float2& a) #ifndef __KERNEL_OPENCL__ -__device_inline float2 interp(float2 a, float2 b, float t) +ccl_device_inline float2 interp(float2 a, float2 b, float t) { return a + t*(b - a); } @@ -388,84 +388,84 @@ __device_inline float2 interp(float2 a, float2 b, float t) #ifndef __KERNEL_OPENCL__ -__device_inline float3 operator-(const float3 a) +ccl_device_inline float3 operator-(const float3 a) { return make_float3(-a.x, -a.y, -a.z); } -__device_inline float3 operator*(const float3 a, const float3 b) +ccl_device_inline float3 operator*(const float3 a, const float3 b) { return make_float3(a.x*b.x, a.y*b.y, a.z*b.z); } -__device_inline float3 operator*(const float3 a, float f) +ccl_device_inline float3 operator*(const float3 a, float f) { return make_float3(a.x*f, a.y*f, a.z*f); } -__device_inline float3 operator*(float f, const float3 a) +ccl_device_inline float3 operator*(float f, const float3 a) { return make_float3(a.x*f, a.y*f, a.z*f); } -__device_inline float3 operator/(float f, const float3 a) +ccl_device_inline float3 operator/(float f, const float3 a) { return make_float3(f/a.x, f/a.y, f/a.z); } -__device_inline float3 operator/(const float3 a, float f) +ccl_device_inline float3 operator/(const float3 a, float f) { float invf = 1.0f/f; return make_float3(a.x*invf, a.y*invf, a.z*invf); } -__device_inline float3 operator/(const float3 a, const float3 b) +ccl_device_inline float3 operator/(const float3 a, const float3 b) { return make_float3(a.x/b.x, a.y/b.y, a.z/b.z); } -__device_inline float3 operator+(const float3 a, const float3 b) +ccl_device_inline float3 operator+(const float3 a, const float3 b) { return make_float3(a.x+b.x, a.y+b.y, a.z+b.z); } -__device_inline float3 operator-(const float3 a, const float3 b) +ccl_device_inline float3 operator-(const float3 a, const float3 b) { return make_float3(a.x-b.x, a.y-b.y, a.z-b.z); } -__device_inline float3 operator+=(float3& a, const float3 b) +ccl_device_inline float3 operator+=(float3& a, const float3 b) { return a = a + b; } -__device_inline float3 operator*=(float3& a, const float3 b) +ccl_device_inline float3 operator*=(float3& a, const float3 b) { return a = a * b; } -__device_inline float3 operator*=(float3& a, float f) +ccl_device_inline float3 operator*=(float3& a, float f) { return a = a * f; } -__device_inline float3 operator/=(float3& a, const float3 b) +ccl_device_inline float3 operator/=(float3& a, const float3 b) { return a = a / b; } -__device_inline float3 operator/=(float3& a, float f) +ccl_device_inline float3 operator/=(float3& a, float f) { float invf = 1.0f/f; return a = a * invf; } -__device_inline float dot(const float3 a, const float3 b) +ccl_device_inline float dot(const float3 a, const float3 b) { return a.x*b.x + a.y*b.y + a.z*b.z; } -__device_inline float3 cross(const float3 a, const float3 b) +ccl_device_inline float3 cross(const float3 a, const float3 b) { float3 r = make_float3(a.y*b.z - a.z*b.y, a.z*b.x - a.x*b.z, a.x*b.y - a.y*b.x); return r; @@ -473,26 +473,26 @@ __device_inline float3 cross(const float3 a, const float3 b) #endif -__device_inline float len(const float3 a) +ccl_device_inline float len(const float3 a) { return sqrtf(dot(a, a)); } -__device_inline float len_squared(const float3 a) +ccl_device_inline float len_squared(const float3 a) { return dot(a, a); } #ifndef __KERNEL_OPENCL__ -__device_inline float3 normalize(const float3 a) +ccl_device_inline float3 normalize(const float3 a) { return a/len(a); } #endif -__device_inline float3 normalize_len(const float3 a, float *t) +ccl_device_inline float3 normalize_len(const float3 a, float *t) { *t = len(a); return a/(*t); @@ -500,7 +500,7 @@ __device_inline float3 normalize_len(const float3 a, float *t) #ifndef __KERNEL_OPENCL__ -__device_inline bool operator==(const float3 a, const float3 b) +ccl_device_inline bool operator==(const float3 a, const float3 b) { #ifdef __KERNEL_SSE__ return (_mm_movemask_ps(_mm_cmpeq_ps(a.m128, b.m128)) & 7) == 7; @@ -509,12 +509,12 @@ __device_inline bool operator==(const float3 a, const float3 b) #endif } -__device_inline bool operator!=(const float3 a, const float3 b) +ccl_device_inline bool operator!=(const float3 a, const float3 b) { return !(a == b); } -__device_inline float3 min(float3 a, float3 b) +ccl_device_inline float3 min(float3 a, float3 b) { #ifdef __KERNEL_SSE__ return _mm_min_ps(a.m128, b.m128); @@ -523,7 +523,7 @@ __device_inline float3 min(float3 a, float3 b) #endif } -__device_inline float3 max(float3 a, float3 b) +ccl_device_inline float3 max(float3 a, float3 b) { #ifdef __KERNEL_SSE__ return _mm_max_ps(a.m128, b.m128); @@ -532,12 +532,12 @@ __device_inline float3 max(float3 a, float3 b) #endif } -__device_inline float3 clamp(float3 a, float3 mn, float3 mx) +ccl_device_inline float3 clamp(float3 a, float3 mn, float3 mx) { return min(max(a, mn), mx); } -__device_inline float3 fabs(float3 a) +ccl_device_inline float3 fabs(float3 a) { #ifdef __KERNEL_SSE__ __m128 mask = _mm_castsi128_ps(_mm_set1_epi32(0x7fffffff)); @@ -549,29 +549,29 @@ __device_inline float3 fabs(float3 a) #endif -__device_inline float3 float2_to_float3(const float2 a) +ccl_device_inline float3 float2_to_float3(const float2 a) { return make_float3(a.x, a.y, 0.0f); } -__device_inline float3 float4_to_float3(const float4 a) +ccl_device_inline float3 float4_to_float3(const float4 a) { return make_float3(a.x, a.y, a.z); } -__device_inline float4 float3_to_float4(const float3 a) +ccl_device_inline float4 float3_to_float4(const float3 a) { return make_float4(a.x, a.y, a.z, 1.0f); } #ifndef __KERNEL_GPU__ -__device_inline void print_float3(const char *label, const float3& a) +ccl_device_inline void print_float3(const char *label, const float3& a) { printf("%s: %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z); } -__device_inline float3 rcp(const float3& a) +ccl_device_inline float3 rcp(const float3& a) { #ifdef __KERNEL_SSE__ float4 r = _mm_rcp_ps(a.m128); @@ -583,12 +583,12 @@ __device_inline float3 rcp(const float3& a) #endif -__device_inline float3 interp(float3 a, float3 b, float t) +ccl_device_inline float3 interp(float3 a, float3 b, float t) { return a + t*(b - a); } -__device_inline bool is_zero(const float3 a) +ccl_device_inline bool is_zero(const float3 a) { #ifdef __KERNEL_SSE__ return a == make_float3(0.0f); @@ -597,7 +597,7 @@ __device_inline bool is_zero(const float3 a) #endif } -__device_inline float reduce_add(const float3 a) +ccl_device_inline float reduce_add(const float3 a) { #ifdef __KERNEL_SSE__ return (a.x + a.y + a.z); @@ -606,7 +606,7 @@ __device_inline float reduce_add(const float3 a) #endif } -__device_inline float average(const float3 a) +ccl_device_inline float average(const float3 a) { return reduce_add(a)*(1.0f/3.0f); } @@ -639,7 +639,7 @@ template<> __forceinline const float4 shuffle<0, 1, 0, 1>(const float4& b) #ifndef __KERNEL_OPENCL__ -__device_inline float4 operator-(const float4& a) +ccl_device_inline float4 operator-(const float4& a) { #ifdef __KERNEL_SSE__ __m128 mask = _mm_castsi128_ps(_mm_set1_epi32(0x80000000)); @@ -649,7 +649,7 @@ __device_inline float4 operator-(const float4& a) #endif } -__device_inline float4 operator*(const float4& a, const float4& b) +ccl_device_inline float4 operator*(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return _mm_mul_ps(a.m128, b.m128); @@ -658,7 +658,7 @@ __device_inline float4 operator*(const float4& a, const float4& b) #endif } -__device_inline float4 operator*(const float4& a, float f) +ccl_device_inline float4 operator*(const float4& a, float f) { #ifdef __KERNEL_SSE__ return a * make_float4(f); @@ -667,12 +667,12 @@ __device_inline float4 operator*(const float4& a, float f) #endif } -__device_inline float4 operator*(float f, const float4& a) +ccl_device_inline float4 operator*(float f, const float4& a) { return a * f; } -__device_inline float4 rcp(const float4& a) +ccl_device_inline float4 rcp(const float4& a) { #ifdef __KERNEL_SSE__ float4 r = _mm_rcp_ps(a.m128); @@ -682,12 +682,12 @@ __device_inline float4 rcp(const float4& a) #endif } -__device_inline float4 operator/(const float4& a, float f) +ccl_device_inline float4 operator/(const float4& a, float f) { return a * (1.0f/f); } -__device_inline float4 operator/(const float4& a, const float4& b) +ccl_device_inline float4 operator/(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return a * rcp(b); @@ -697,7 +697,7 @@ __device_inline float4 operator/(const float4& a, const float4& b) } -__device_inline float4 operator+(const float4& a, const float4& b) +ccl_device_inline float4 operator+(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return _mm_add_ps(a.m128, b.m128); @@ -706,7 +706,7 @@ __device_inline float4 operator+(const float4& a, const float4& b) #endif } -__device_inline float4 operator-(const float4& a, const float4& b) +ccl_device_inline float4 operator-(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return _mm_sub_ps(a.m128, b.m128); @@ -715,22 +715,22 @@ __device_inline float4 operator-(const float4& a, const float4& b) #endif } -__device_inline float4 operator+=(float4& a, const float4& b) +ccl_device_inline float4 operator+=(float4& a, const float4& b) { return a = a + b; } -__device_inline float4 operator*=(float4& a, const float4& b) +ccl_device_inline float4 operator*=(float4& a, const float4& b) { return a = a * b; } -__device_inline float4 operator/=(float4& a, float f) +ccl_device_inline float4 operator/=(float4& a, float f) { return a = a / f; } -__device_inline int4 operator<(const float4& a, const float4& b) +ccl_device_inline int4 operator<(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return _mm_cvtps_epi32(_mm_cmplt_ps(a.m128, b.m128)); /* todo: avoid cvt */ @@ -739,7 +739,7 @@ __device_inline int4 operator<(const float4& a, const float4& b) #endif } -__device_inline int4 operator>=(float4 a, float4 b) +ccl_device_inline int4 operator>=(float4 a, float4 b) { #ifdef __KERNEL_SSE__ return _mm_cvtps_epi32(_mm_cmpge_ps(a.m128, b.m128)); /* todo: avoid cvt */ @@ -748,7 +748,7 @@ __device_inline int4 operator>=(float4 a, float4 b) #endif } -__device_inline int4 operator<=(const float4& a, const float4& b) +ccl_device_inline int4 operator<=(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return _mm_cvtps_epi32(_mm_cmple_ps(a.m128, b.m128)); /* todo: avoid cvt */ @@ -757,7 +757,7 @@ __device_inline int4 operator<=(const float4& a, const float4& b) #endif } -__device_inline bool operator==(const float4 a, const float4 b) +ccl_device_inline bool operator==(const float4 a, const float4 b) { #ifdef __KERNEL_SSE__ return (_mm_movemask_ps(_mm_cmpeq_ps(a.m128, b.m128)) & 15) == 15; @@ -766,7 +766,7 @@ __device_inline bool operator==(const float4 a, const float4 b) #endif } -__device_inline float4 cross(const float4& a, const float4& b) +ccl_device_inline float4 cross(const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ return (shuffle<1,2,0,0>(a)*shuffle<2,0,1,0>(b)) - (shuffle<2,0,1,0>(a)*shuffle<1,2,0,0>(b)); @@ -775,7 +775,7 @@ __device_inline float4 cross(const float4& a, const float4& b) #endif } -__device_inline bool is_zero(const float4& a) +ccl_device_inline bool is_zero(const float4& a) { #ifdef __KERNEL_SSE__ return a == make_float4(0.0f); @@ -784,7 +784,7 @@ __device_inline bool is_zero(const float4& a) #endif } -__device_inline float reduce_add(const float4& a) +ccl_device_inline float reduce_add(const float4& a) { #ifdef __KERNEL_SSE__ float4 h = shuffle<1,0,3,2>(a) + a; @@ -794,27 +794,27 @@ __device_inline float reduce_add(const float4& a) #endif } -__device_inline float average(const float4& a) +ccl_device_inline float average(const float4& a) { return reduce_add(a) * 0.25f; } -__device_inline float dot(const float4& a, const float4& b) +ccl_device_inline float dot(const float4& a, const float4& b) { return reduce_add(a * b); } -__device_inline float len(const float4 a) +ccl_device_inline float len(const float4 a) { return sqrtf(dot(a, a)); } -__device_inline float4 normalize(const float4 a) +ccl_device_inline float4 normalize(const float4 a) { return a/len(a); } -__device_inline float4 min(float4 a, float4 b) +ccl_device_inline float4 min(float4 a, float4 b) { #ifdef __KERNEL_SSE__ return _mm_min_ps(a.m128, b.m128); @@ -823,7 +823,7 @@ __device_inline float4 min(float4 a, float4 b) #endif } -__device_inline float4 max(float4 a, float4 b) +ccl_device_inline float4 max(float4 a, float4 b) { #ifdef __KERNEL_SSE__ return _mm_max_ps(a.m128, b.m128); @@ -836,7 +836,7 @@ __device_inline float4 max(float4 a, float4 b) #ifndef __KERNEL_GPU__ -__device_inline float4 select(const int4& mask, const float4& a, const float4& b) +ccl_device_inline float4 select(const int4& mask, const float4& a, const float4& b) { #ifdef __KERNEL_SSE__ /* blendv is sse4, and apparently broken on vs2008 */ @@ -846,7 +846,7 @@ __device_inline float4 select(const int4& mask, const float4& a, const float4& b #endif } -__device_inline float4 reduce_min(const float4& a) +ccl_device_inline float4 reduce_min(const float4& a) { #ifdef __KERNEL_SSE__ float4 h = min(shuffle<1,0,3,2>(a), a); @@ -856,7 +856,7 @@ __device_inline float4 reduce_min(const float4& a) #endif } -__device_inline float4 reduce_max(const float4& a) +ccl_device_inline float4 reduce_max(const float4& a) { #ifdef __KERNEL_SSE__ float4 h = max(shuffle<1,0,3,2>(a), a); @@ -867,7 +867,7 @@ __device_inline float4 reduce_max(const float4& a) } #if 0 -__device_inline float4 reduce_add(const float4& a) +ccl_device_inline float4 reduce_add(const float4& a) { #ifdef __KERNEL_SSE__ float4 h = shuffle<1,0,3,2>(a) + a; @@ -878,7 +878,7 @@ __device_inline float4 reduce_add(const float4& a) } #endif -__device_inline void print_float4(const char *label, const float4& a) +ccl_device_inline void print_float4(const char *label, const float4& a) { printf("%s: %.8f %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z, (double)a.w); } @@ -889,7 +889,7 @@ __device_inline void print_float4(const char *label, const float4& a) #ifndef __KERNEL_OPENCL__ -__device_inline int3 min(int3 a, int3 b) +ccl_device_inline int3 min(int3 a, int3 b) { #ifdef __KERNEL_SSE__ return _mm_min_epi32(a.m128, b.m128); @@ -898,7 +898,7 @@ __device_inline int3 min(int3 a, int3 b) #endif } -__device_inline int3 max(int3 a, int3 b) +ccl_device_inline int3 max(int3 a, int3 b) { #ifdef __KERNEL_SSE__ return _mm_max_epi32(a.m128, b.m128); @@ -907,7 +907,7 @@ __device_inline int3 max(int3 a, int3 b) #endif } -__device_inline int3 clamp(const int3& a, int mn, int mx) +ccl_device_inline int3 clamp(const int3& a, int mn, int mx) { #ifdef __KERNEL_SSE__ return min(max(a, make_int3(mn)), make_int3(mx)); @@ -916,7 +916,7 @@ __device_inline int3 clamp(const int3& a, int mn, int mx) #endif } -__device_inline int3 clamp(const int3& a, int3& mn, int mx) +ccl_device_inline int3 clamp(const int3& a, int3& mn, int mx) { #ifdef __KERNEL_SSE__ return min(max(a, mn), make_int3(mx)); @@ -929,7 +929,7 @@ __device_inline int3 clamp(const int3& a, int3& mn, int mx) #ifndef __KERNEL_GPU__ -__device_inline void print_int3(const char *label, const int3& a) +ccl_device_inline void print_int3(const char *label, const int3& a) { printf("%s: %d %d %d\n", label, a.x, a.y, a.z); } @@ -940,7 +940,7 @@ __device_inline void print_int3(const char *label, const int3& a) #ifndef __KERNEL_GPU__ -__device_inline int4 operator+(const int4& a, const int4& b) +ccl_device_inline int4 operator+(const int4& a, const int4& b) { #ifdef __KERNEL_SSE__ return _mm_add_epi32(a.m128, b.m128); @@ -949,12 +949,12 @@ __device_inline int4 operator+(const int4& a, const int4& b) #endif } -__device_inline int4 operator+=(int4& a, const int4& b) +ccl_device_inline int4 operator+=(int4& a, const int4& b) { return a = a + b; } -__device_inline int4 operator>>(const int4& a, int i) +ccl_device_inline int4 operator>>(const int4& a, int i) { #ifdef __KERNEL_SSE__ return _mm_srai_epi32(a.m128, i); @@ -963,7 +963,7 @@ __device_inline int4 operator>>(const int4& a, int i) #endif } -__device_inline int4 min(int4 a, int4 b) +ccl_device_inline int4 min(int4 a, int4 b) { #ifdef __KERNEL_SSE__ return _mm_min_epi32(a.m128, b.m128); @@ -972,7 +972,7 @@ __device_inline int4 min(int4 a, int4 b) #endif } -__device_inline int4 max(int4 a, int4 b) +ccl_device_inline int4 max(int4 a, int4 b) { #ifdef __KERNEL_SSE__ return _mm_max_epi32(a.m128, b.m128); @@ -981,12 +981,12 @@ __device_inline int4 max(int4 a, int4 b) #endif } -__device_inline int4 clamp(const int4& a, const int4& mn, const int4& mx) +ccl_device_inline int4 clamp(const int4& a, const int4& mn, const int4& mx) { return min(max(a, mn), mx); } -__device_inline int4 select(const int4& mask, const int4& a, const int4& b) +ccl_device_inline int4 select(const int4& mask, const int4& a, const int4& b) { #ifdef __KERNEL_SSE__ __m128 m = _mm_cvtepi32_ps(mask); @@ -996,7 +996,7 @@ __device_inline int4 select(const int4& mask, const int4& a, const int4& b) #endif } -__device_inline void print_int4(const char *label, const int4& a) +ccl_device_inline void print_int4(const char *label, const int4& a) { printf("%s: %d %d %d %d\n", label, a.x, a.y, a.z, a.w); } @@ -1007,49 +1007,49 @@ __device_inline void print_int4(const char *label, const int4& a) #ifndef __KERNEL_OPENCL__ -__device_inline int as_int(uint i) +ccl_device_inline int as_int(uint i) { union { uint ui; int i; } u; u.ui = i; return u.i; } -__device_inline uint as_uint(int i) +ccl_device_inline uint as_uint(int i) { union { uint ui; int i; } u; u.i = i; return u.ui; } -__device_inline uint as_uint(float f) +ccl_device_inline uint as_uint(float f) { union { uint i; float f; } u; u.f = f; return u.i; } -__device_inline int __float_as_int(float f) +ccl_device_inline int __float_as_int(float f) { union { int i; float f; } u; u.f = f; return u.i; } -__device_inline float __int_as_float(int i) +ccl_device_inline float __int_as_float(int i) { union { int i; float f; } u; u.i = i; return u.f; } -__device_inline uint __float_as_uint(float f) +ccl_device_inline uint __float_as_uint(float f) { union { uint i; float f; } u; u.f = f; return u.i; } -__device_inline float __uint_as_float(uint i) +ccl_device_inline float __uint_as_float(uint i) { union { uint i; float f; } u; u.i = i; @@ -1065,7 +1065,7 @@ template<class A, class B> A lerp(const A& a, const A& b, const B& t) /* Triangle */ -__device_inline float triangle_area(const float3 v1, const float3 v2, const float3 v3) +ccl_device_inline float triangle_area(const float3 v1, const float3 v2, const float3 v3) { return len(cross(v3 - v2, v1 - v2))*0.5f; } @@ -1074,7 +1074,7 @@ __device_inline float triangle_area(const float3 v1, const float3 v2, const floa /* Orthonormal vectors */ -__device_inline void make_orthonormals(const float3 N, float3 *a, float3 *b) +ccl_device_inline void make_orthonormals(const float3 N, float3 *a, float3 *b) { #if 0 if(fabsf(N.y) >= 0.999f) { @@ -1100,7 +1100,7 @@ __device_inline void make_orthonormals(const float3 N, float3 *a, float3 *b) /* Color division */ -__device_inline float3 safe_divide_color(float3 a, float3 b) +ccl_device_inline float3 safe_divide_color(float3 a, float3 b) { float x, y, z; @@ -1111,7 +1111,7 @@ __device_inline float3 safe_divide_color(float3 a, float3 b) return make_float3(x, y, z); } -__device_inline float3 safe_divide_even_color(float3 a, float3 b) +ccl_device_inline float3 safe_divide_even_color(float3 a, float3 b) { float x, y, z; @@ -1149,7 +1149,7 @@ __device_inline float3 safe_divide_even_color(float3 a, float3 b) /* Rotation of point around axis and angle */ -__device_inline float3 rotate_around_axis(float3 p, float3 axis, float angle) +ccl_device_inline float3 rotate_around_axis(float3 p, float3 axis, float angle) { float costheta = cosf(angle); float sintheta = sinf(angle); @@ -1172,12 +1172,12 @@ __device_inline float3 rotate_around_axis(float3 p, float3 axis, float angle) /* NaN-safe math ops */ -__device_inline float safe_sqrtf(float f) +ccl_device_inline float safe_sqrtf(float f) { return sqrtf(max(f, 0.0f)); } -__device float safe_asinf(float a) +ccl_device float safe_asinf(float a) { if(a <= -1.0f) return -M_PI_2_F; @@ -1187,7 +1187,7 @@ __device float safe_asinf(float a) return asinf(a); } -__device float safe_acosf(float a) +ccl_device float safe_acosf(float a) { if(a <= -1.0f) return M_PI_F; @@ -1197,7 +1197,7 @@ __device float safe_acosf(float a) return acosf(a); } -__device float compatible_powf(float x, float y) +ccl_device float compatible_powf(float x, float y) { /* GPU pow doesn't accept negative x, do manual checks here */ if(x < 0.0f) { @@ -1212,7 +1212,7 @@ __device float compatible_powf(float x, float y) return powf(x, y); } -__device float safe_powf(float a, float b) +ccl_device float safe_powf(float a, float b) { if(b == 0.0f) return 1.0f; @@ -1224,7 +1224,7 @@ __device float safe_powf(float a, float b) return compatible_powf(a, b); } -__device float safe_logf(float a, float b) +ccl_device float safe_logf(float a, float b) { if(a < 0.0f || b < 0.0f) return 0.0f; @@ -1232,19 +1232,19 @@ __device float safe_logf(float a, float b) return logf(a)/logf(b); } -__device float safe_divide(float a, float b) +ccl_device float safe_divide(float a, float b) { return (b != 0.0f)? a/b: 0.0f; } -__device float safe_modulo(float a, float b) +ccl_device float safe_modulo(float a, float b) { return (b != 0.0f)? fmodf(a, b): 0.0f; } /* Ray Intersection */ -__device bool ray_sphere_intersect( +ccl_device bool ray_sphere_intersect( float3 ray_P, float3 ray_D, float ray_t, float3 sphere_P, float sphere_radius, float3 *isect_P, float *isect_t) @@ -1276,7 +1276,7 @@ __device bool ray_sphere_intersect( return false; } -__device bool ray_aligned_disk_intersect( +ccl_device bool ray_aligned_disk_intersect( float3 ray_P, float3 ray_D, float ray_t, float3 disk_P, float disk_radius, float3 *isect_P, float *isect_t) @@ -1305,7 +1305,7 @@ __device bool ray_aligned_disk_intersect( return true; } -__device bool ray_triangle_intersect( +ccl_device bool ray_triangle_intersect( float3 ray_P, float3 ray_D, float ray_t, float3 v0, float3 v1, float3 v2, float3 *isect_P, float *isect_t) @@ -1348,7 +1348,7 @@ __device bool ray_triangle_intersect( return true; } -__device bool ray_quad_intersect( +ccl_device bool ray_quad_intersect( float3 ray_P, float3 ray_D, float ray_t, float3 quad_P, float3 quad_u, float3 quad_v, float3 *isect_P, float *isect_t) diff --git a/intern/cycles/util/util_transform.h b/intern/cycles/util/util_transform.h index 01655bdf9bc..1b6315ba484 100644 --- a/intern/cycles/util/util_transform.h +++ b/intern/cycles/util/util_transform.h @@ -43,7 +43,7 @@ typedef struct Transform { * * For the DecompMotionTransform we drop scale from pre/post. */ -typedef struct __may_alias MotionTransform { +typedef struct ccl_may_alias MotionTransform { Transform pre; Transform mid; Transform post; @@ -57,7 +57,7 @@ typedef struct DecompMotionTransform { /* Functions */ -__device_inline float3 transform_perspective(const Transform *t, const float3 a) +ccl_device_inline float3 transform_perspective(const Transform *t, const float3 a) { float4 b = make_float4(a.x, a.y, a.z, 1.0f); float3 c = make_float3(dot(t->x, b), dot(t->y, b), dot(t->z, b)); @@ -66,7 +66,7 @@ __device_inline float3 transform_perspective(const Transform *t, const float3 a) return (w != 0.0f)? c/w: make_float3(0.0f, 0.0f, 0.0f); } -__device_inline float3 transform_point(const Transform *t, const float3 a) +ccl_device_inline float3 transform_point(const Transform *t, const float3 a) { float3 c = make_float3( a.x*t->x.x + a.y*t->x.y + a.z*t->x.z + t->x.w, @@ -76,7 +76,7 @@ __device_inline float3 transform_point(const Transform *t, const float3 a) return c; } -__device_inline float3 transform_direction(const Transform *t, const float3 a) +ccl_device_inline float3 transform_direction(const Transform *t, const float3 a) { float3 c = make_float3( a.x*t->x.x + a.y*t->x.y + a.z*t->x.z, @@ -86,7 +86,7 @@ __device_inline float3 transform_direction(const Transform *t, const float3 a) return c; } -__device_inline float3 transform_direction_transposed(const Transform *t, const float3 a) +ccl_device_inline float3 transform_direction_transposed(const Transform *t, const float3 a) { float3 x = make_float3(t->x.x, t->y.x, t->z.x); float3 y = make_float3(t->x.y, t->y.y, t->z.y); @@ -95,7 +95,7 @@ __device_inline float3 transform_direction_transposed(const Transform *t, const return make_float3(dot(x, a), dot(y, a), dot(z, a)); } -__device_inline Transform transform_transpose(const Transform a) +ccl_device_inline Transform transform_transpose(const Transform a) { Transform t; @@ -107,7 +107,7 @@ __device_inline Transform transform_transpose(const Transform a) return t; } -__device_inline Transform make_transform(float a, float b, float c, float d, +ccl_device_inline Transform make_transform(float a, float b, float c, float d, float e, float f, float g, float h, float i, float j, float k, float l, float m, float n, float o, float p) @@ -124,7 +124,7 @@ __device_inline Transform make_transform(float a, float b, float c, float d, #ifndef __KERNEL_GPU__ -__device_inline Transform operator*(const Transform a, const Transform b) +ccl_device_inline Transform operator*(const Transform a, const Transform b) { Transform c = transform_transpose(b); Transform t; @@ -137,7 +137,7 @@ __device_inline Transform operator*(const Transform a, const Transform b) return t; } -__device_inline void print_transform(const char *label, const Transform& t) +ccl_device_inline void print_transform(const char *label, const Transform& t) { print_float4(label, t.x); print_float4(label, t.y); @@ -146,7 +146,7 @@ __device_inline void print_transform(const char *label, const Transform& t) printf("\n"); } -__device_inline Transform transform_translate(float3 t) +ccl_device_inline Transform transform_translate(float3 t) { return make_transform( 1, 0, 0, t.x, @@ -155,12 +155,12 @@ __device_inline Transform transform_translate(float3 t) 0, 0, 0, 1); } -__device_inline Transform transform_translate(float x, float y, float z) +ccl_device_inline Transform transform_translate(float x, float y, float z) { return transform_translate(make_float3(x, y, z)); } -__device_inline Transform transform_scale(float3 s) +ccl_device_inline Transform transform_scale(float3 s) { return make_transform( s.x, 0, 0, 0, @@ -169,12 +169,12 @@ __device_inline Transform transform_scale(float3 s) 0, 0, 0, 1); } -__device_inline Transform transform_scale(float x, float y, float z) +ccl_device_inline Transform transform_scale(float x, float y, float z) { return transform_scale(make_float3(x, y, z)); } -__device_inline Transform transform_perspective(float fov, float n, float f) +ccl_device_inline Transform transform_perspective(float fov, float n, float f) { Transform persp = make_transform( 1, 0, 0, 0, @@ -189,7 +189,7 @@ __device_inline Transform transform_perspective(float fov, float n, float f) return scale * persp; } -__device_inline Transform transform_rotate(float angle, float3 axis) +ccl_device_inline Transform transform_rotate(float angle, float3 axis) { float s = sinf(angle); float c = cosf(angle); @@ -216,7 +216,7 @@ __device_inline Transform transform_rotate(float angle, float3 axis) 0.0f, 0.0f, 0.0f, 1.0f); } -__device_inline Transform transform_euler(float3 euler) +ccl_device_inline Transform transform_euler(float3 euler) { return transform_rotate(euler.x, make_float3(1.0f, 0.0f, 0.0f)) * @@ -224,33 +224,33 @@ __device_inline Transform transform_euler(float3 euler) transform_rotate(euler.z, make_float3(0.0f, 0.0f, 1.0f)); } -__device_inline Transform transform_orthographic(float znear, float zfar) +ccl_device_inline Transform transform_orthographic(float znear, float zfar) { return transform_scale(1.0f, 1.0f, 1.0f / (zfar-znear)) * transform_translate(0.0f, 0.0f, -znear); } -__device_inline Transform transform_identity() +ccl_device_inline Transform transform_identity() { return transform_scale(1.0f, 1.0f, 1.0f); } -__device_inline bool operator==(const Transform& A, const Transform& B) +ccl_device_inline bool operator==(const Transform& A, const Transform& B) { return memcmp(&A, &B, sizeof(Transform)) == 0; } -__device_inline bool operator!=(const Transform& A, const Transform& B) +ccl_device_inline bool operator!=(const Transform& A, const Transform& B) { return !(A == B); } -__device_inline float3 transform_get_column(const Transform *t, int column) +ccl_device_inline float3 transform_get_column(const Transform *t, int column) { return make_float3(t->x[column], t->y[column], t->z[column]); } -__device_inline void transform_set_column(Transform *t, int column, float3 value) +ccl_device_inline void transform_set_column(Transform *t, int column, float3 value) { t->x[column] = value.x; t->y[column] = value.y; @@ -259,7 +259,7 @@ __device_inline void transform_set_column(Transform *t, int column, float3 value Transform transform_inverse(const Transform& a); -__device_inline bool transform_uniform_scale(const Transform& tfm, float& scale) +ccl_device_inline bool transform_uniform_scale(const Transform& tfm, float& scale) { /* the epsilon here is quite arbitrary, but this function is only used for * surface area and bump, where we except it to not be so sensitive */ @@ -283,7 +283,7 @@ __device_inline bool transform_uniform_scale(const Transform& tfm, float& scale) return false; } -__device_inline bool transform_negative_scale(const Transform& tfm) +ccl_device_inline bool transform_negative_scale(const Transform& tfm) { float3 c0 = transform_get_column(&tfm, 0); float3 c1 = transform_get_column(&tfm, 1); @@ -292,7 +292,7 @@ __device_inline bool transform_negative_scale(const Transform& tfm) return (dot(cross(c0, c1), c2) < 0.0f); } -__device_inline Transform transform_clear_scale(const Transform& tfm) +ccl_device_inline Transform transform_clear_scale(const Transform& tfm) { Transform ntfm = tfm; @@ -307,7 +307,7 @@ __device_inline Transform transform_clear_scale(const Transform& tfm) /* Motion Transform */ -__device_inline float4 quat_interpolate(float4 q1, float4 q2, float t) +ccl_device_inline float4 quat_interpolate(float4 q1, float4 q2, float t) { /* use simpe nlerp instead of slerp. it's faster and almost the same */ return normalize((1.0f - t)*q1 + t*q2); @@ -333,7 +333,7 @@ __device_inline float4 quat_interpolate(float4 q1, float4 q2, float t) #endif } -__device_inline Transform transform_quick_inverse(Transform M) +ccl_device_inline Transform transform_quick_inverse(Transform M) { /* possible optimization: can we avoid doing this altogether and construct * the inverse matrix directly from negated translation, transposed rotation, @@ -356,7 +356,7 @@ __device_inline Transform transform_quick_inverse(Transform M) return R; } -__device_inline void transform_compose(Transform *tfm, const Transform *decomp) +ccl_device_inline void transform_compose(Transform *tfm, const Transform *decomp) { /* rotation */ float q0, q1, q2, q3, qda, qdb, qdc, qaa, qab, qac, qbb, qbc, qcc; @@ -395,7 +395,7 @@ __device_inline void transform_compose(Transform *tfm, const Transform *decomp) /* Disabled for now, need arc-length parametrization for constant speed motion. * #define CURVED_MOTION_INTERPOLATE */ -__device void transform_motion_interpolate(Transform *tfm, const DecompMotionTransform *motion, float t) +ccl_device void transform_motion_interpolate(Transform *tfm, const DecompMotionTransform *motion, float t) { /* possible optimization: is it worth it adding a check to skip scaling? * it's probably quite uncommon to have scaling objects. or can we skip @@ -447,7 +447,7 @@ __device void transform_motion_interpolate(Transform *tfm, const DecompMotionTra #ifndef __KERNEL_GPU__ -__device_inline bool operator==(const MotionTransform& A, const MotionTransform& B) +ccl_device_inline bool operator==(const MotionTransform& A, const MotionTransform& B) { return (A.pre == B.pre && A.post == B.post); } diff --git a/intern/cycles/util/util_types.h b/intern/cycles/util/util_types.h index a8f514864db..c53d67235f6 100644 --- a/intern/cycles/util/util_types.h +++ b/intern/cycles/util/util_types.h @@ -27,24 +27,22 @@ #ifndef __KERNEL_GPU__ -#define __device static inline -#define __device_noinline static -#define __global -#define __local -#define __shared -#define __constant +#define ccl_device static inline +#define ccl_device_noinline static +#define ccl_global +#define ccl_constant #if defined(_WIN32) && !defined(FREE_WINDOWS) -#define __device_inline static __forceinline -#define __align(...) __declspec(align(__VA_ARGS__)) -#define __may_alias +#define ccl_device_inline static __forceinline +#define ccl_align(...) __declspec(align(__VA_ARGS__)) +#define ccl_may_alias #else -#define __device_inline static inline __attribute__((always_inline)) +#define ccl_device_inline static inline __attribute__((always_inline)) #ifndef FREE_WINDOWS64 #define __forceinline inline __attribute__((always_inline)) #endif -#define __align(...) __attribute__((aligned(__VA_ARGS__))) -#define __may_alias __attribute__((__may_alias__)) +#define ccl_align(...) __attribute__((aligned(__VA_ARGS__))) +#define ccl_may_alias __attribute__((__may_alias__)) #endif #endif @@ -190,7 +188,7 @@ struct int2 { }; #ifdef __KERNEL_SSE__ -struct __align(16) int3 { +struct ccl_align(16) int3 { union { __m128i m128; struct { int x, y, z, w; }; @@ -210,7 +208,7 @@ struct int3 { }; #ifdef __KERNEL_SSE__ -struct __align(16) int4 { +struct ccl_align(16) int4 { union { __m128i m128; struct { int x, y, z, w; }; @@ -258,7 +256,7 @@ struct float2 { }; #ifdef __KERNEL_SSE__ -struct __align(16) float3 { +struct ccl_align(16) float3 { union { __m128 m128; struct { float x, y, z, w; }; @@ -278,7 +276,7 @@ struct float3 { }; #ifdef __KERNEL_SSE__ -struct __align(16) float4 { +struct ccl_align(16) float4 { union { __m128 m128; struct { float x, y, z, w; }; @@ -305,31 +303,31 @@ struct float4 { * * OpenCL does not support C++ class, so we use these instead. */ -__device_inline uchar2 make_uchar2(uchar x, uchar y) +ccl_device_inline uchar2 make_uchar2(uchar x, uchar y) { uchar2 a = {x, y}; return a; } -__device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z) +ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z) { uchar3 a = {x, y, z}; return a; } -__device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w) +ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w) { uchar4 a = {x, y, z, w}; return a; } -__device_inline int2 make_int2(int x, int y) +ccl_device_inline int2 make_int2(int x, int y) { int2 a = {x, y}; return a; } -__device_inline int3 make_int3(int x, int y, int z) +ccl_device_inline int3 make_int3(int x, int y, int z) { #ifdef __KERNEL_SSE__ int3 a; @@ -341,7 +339,7 @@ __device_inline int3 make_int3(int x, int y, int z) return a; } -__device_inline int4 make_int4(int x, int y, int z, int w) +ccl_device_inline int4 make_int4(int x, int y, int z, int w) { #ifdef __KERNEL_SSE__ int4 a; @@ -353,31 +351,31 @@ __device_inline int4 make_int4(int x, int y, int z, int w) return a; } -__device_inline uint2 make_uint2(uint x, uint y) +ccl_device_inline uint2 make_uint2(uint x, uint y) { uint2 a = {x, y}; return a; } -__device_inline uint3 make_uint3(uint x, uint y, uint z) +ccl_device_inline uint3 make_uint3(uint x, uint y, uint z) { uint3 a = {x, y, z}; return a; } -__device_inline uint4 make_uint4(uint x, uint y, uint z, uint w) +ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w) { uint4 a = {x, y, z, w}; return a; } -__device_inline float2 make_float2(float x, float y) +ccl_device_inline float2 make_float2(float x, float y) { float2 a = {x, y}; return a; } -__device_inline float3 make_float3(float x, float y, float z) +ccl_device_inline float3 make_float3(float x, float y, float z) { #ifdef __KERNEL_SSE__ float3 a; @@ -389,7 +387,7 @@ __device_inline float3 make_float3(float x, float y, float z) return a; } -__device_inline float4 make_float4(float x, float y, float z, float w) +ccl_device_inline float4 make_float4(float x, float y, float z, float w) { #ifdef __KERNEL_SSE__ float4 a; @@ -401,12 +399,12 @@ __device_inline float4 make_float4(float x, float y, float z, float w) return a; } -__device_inline int align_up(int offset, int alignment) +ccl_device_inline int align_up(int offset, int alignment) { return (offset + alignment - 1) & ~(alignment - 1); } -__device_inline int3 make_int3(int i) +ccl_device_inline int3 make_int3(int i) { #ifdef __KERNEL_SSE__ int3 a; @@ -418,7 +416,7 @@ __device_inline int3 make_int3(int i) return a; } -__device_inline int4 make_int4(int i) +ccl_device_inline int4 make_int4(int i) { #ifdef __KERNEL_SSE__ int4 a; @@ -430,7 +428,7 @@ __device_inline int4 make_int4(int i) return a; } -__device_inline float3 make_float3(float f) +ccl_device_inline float3 make_float3(float f) { #ifdef __KERNEL_SSE__ float3 a; @@ -442,7 +440,7 @@ __device_inline float3 make_float3(float f) return a; } -__device_inline float4 make_float4(float f) +ccl_device_inline float4 make_float4(float f) { #ifdef __KERNEL_SSE__ float4 a; @@ -454,7 +452,7 @@ __device_inline float4 make_float4(float f) return a; } -__device_inline float4 make_float4(const int4& i) +ccl_device_inline float4 make_float4(const int4& i) { #ifdef __KERNEL_SSE__ float4 a; @@ -466,7 +464,7 @@ __device_inline float4 make_float4(const int4& i) return a; } -__device_inline int4 make_int4(const float3& f) +ccl_device_inline int4 make_int4(const float3& f) { #ifdef __KERNEL_SSE__ int4 a; @@ -489,17 +487,17 @@ __device_inline int4 make_int4(const float3& f) /* faster version for SSSE3 */ typedef __m128i shuffle_swap_t; -__device_inline const shuffle_swap_t shuffle_swap_identity(void) +ccl_device_inline const shuffle_swap_t shuffle_swap_identity(void) { return _mm_set_epi8(15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0); } -__device_inline const shuffle_swap_t shuffle_swap_swap(void) +ccl_device_inline const shuffle_swap_t shuffle_swap_swap(void) { return _mm_set_epi8(7, 6, 5, 4, 3, 2, 1, 0, 15, 14, 13, 12, 11, 10, 9, 8); } -__device_inline const __m128 shuffle_swap(const __m128& a, const shuffle_swap_t& shuf) +ccl_device_inline const __m128 shuffle_swap(const __m128& a, const shuffle_swap_t& shuf) { return _mm_castsi128_ps(_mm_shuffle_epi8(_mm_castps_si128(a), shuf)); } @@ -509,17 +507,17 @@ __device_inline const __m128 shuffle_swap(const __m128& a, const shuffle_swap_t& /* somewhat slower version for SSE2 */ typedef int shuffle_swap_t; -__device_inline const shuffle_swap_t shuffle_swap_identity(void) +ccl_device_inline const shuffle_swap_t shuffle_swap_identity(void) { return 0; } -__device_inline const shuffle_swap_t shuffle_swap_swap(void) +ccl_device_inline const shuffle_swap_t shuffle_swap_swap(void) { return 1; } -__device_inline const __m128 shuffle_swap(const __m128& a, shuffle_swap_t shuf) +ccl_device_inline const __m128 shuffle_swap(const __m128& a, shuffle_swap_t shuf) { /* shuffle value must be a constant, so we need to branch */ if(shuf) @@ -530,12 +528,12 @@ __device_inline const __m128 shuffle_swap(const __m128& a, shuffle_swap_t shuf) #endif -template<size_t i0, size_t i1, size_t i2, size_t i3> __device_inline const __m128 shuffle(const __m128& a, const __m128& b) +template<size_t i0, size_t i1, size_t i2, size_t i3> ccl_device_inline const __m128 shuffle(const __m128& a, const __m128& b) { return _mm_shuffle_ps(a, b, _MM_SHUFFLE(i3, i2, i1, i0)); } -template<size_t i0, size_t i1, size_t i2, size_t i3> __device_inline const __m128 shuffle(const __m128& b) +template<size_t i0, size_t i1, size_t i2, size_t i3> ccl_device_inline const __m128 shuffle(const __m128& b) { return _mm_castsi128_ps(_mm_shuffle_epi32(_mm_castps_si128(b), _MM_SHUFFLE(i3, i2, i1, i0))); } @@ -554,7 +552,7 @@ struct half4 { half x, y, z, w; }; #ifdef __KERNEL_CUDA__ -__device_inline void float4_store_half(half *h, const float4 *f, float scale) +ccl_device_inline void float4_store_half(half *h, const float4 *f, float scale) { h[0] = __float2half_rn(f->x * scale); h[1] = __float2half_rn(f->y * scale); @@ -564,7 +562,7 @@ __device_inline void float4_store_half(half *h, const float4 *f, float scale) #else -__device_inline void float4_store_half(half *h, const float4 *f, float scale) +ccl_device_inline void float4_store_half(half *h, const float4 *f, float scale) { #ifndef __KERNEL_SSE2__ for(int i = 0; i < 4; i++) { |