diff options
author | Brecht Van Lommel <brechtvanlommel@pandora.be> | 2011-08-10 18:26:51 +0400 |
---|---|---|
committer | Brecht Van Lommel <brechtvanlommel@pandora.be> | 2011-08-10 18:26:51 +0400 |
commit | b98ccf699878d70b7cbfdd7863440c34326a5569 (patch) | |
tree | 7dde3d3deddedb5d54f2a45ffbb517f2999a83dc /intern/cycles/kernel | |
parent | abc601d10ee51aa8de83004094e2e6c73e96e2eb (diff) |
Cycles: amd opencl compatibility fixes.
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/kernel.cl | 1 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_shader.h | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/svm/bsdf_ward.h | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/svm/svm_math.h | 11 | ||||
-rw-r--r-- | intern/cycles/kernel/svm/svm_texture.h | 6 |
5 files changed, 16 insertions, 10 deletions
diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index e1a9b3a0696..007a5f6de6a 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -51,7 +51,6 @@ __kernel void kernel_ocl_path_trace( int x = sx + get_global_id(0); int y = sy + get_global_id(1); - int w = kernel_data.cam.width; if(x < sx + sw && y < sy + sh) kernel_path_trace(kg, buffer, rng_state, pass, x, y); diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index f1abbda7ae5..c4fc65596b2 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -140,14 +140,18 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, #endif /* detect instancing, for non-instanced the object index is -object-1 */ +#ifdef __INSTANCING__ bool instanced = false; if(sd->prim != ~0) { if(sd->object >= 0) instanced = true; else +#endif sd->object = -sd->object-1; +#ifdef __INSTANCING__ } +#endif /* smooth normal */ if(sd->shader < 0) { diff --git a/intern/cycles/kernel/svm/bsdf_ward.h b/intern/cycles/kernel/svm/bsdf_ward.h index 9f857b32468..c54418afa77 100644 --- a/intern/cycles/kernel/svm/bsdf_ward.h +++ b/intern/cycles/kernel/svm/bsdf_ward.h @@ -166,9 +166,7 @@ __device int bsdf_ward_sample(const ShaderData *sd, float randu, float randv, fl h = h.x * X + h.y * Y + h.z * m_N; // generate the final sample float oh = dot(h, sd->I); - omega_in->x = 2 * oh * h.x - sd->I.x; - omega_in->y = 2 * oh * h.y - sd->I.y; - omega_in->z = 2 * oh * h.z - sd->I.z; + *omega_in = 2.0f * oh * h - sd->I; if(dot(sd->Ng, *omega_in) > 0) { float cosNI = dot(m_N, *omega_in); if(cosNI > 0) { diff --git a/intern/cycles/kernel/svm/svm_math.h b/intern/cycles/kernel/svm/svm_math.h index 401bd22b45a..bc2f774097e 100644 --- a/intern/cycles/kernel/svm/svm_math.h +++ b/intern/cycles/kernel/svm/svm_math.h @@ -21,7 +21,7 @@ CCL_NAMESPACE_BEGIN __device float safe_asinf(float a) { if(a <= -1.0f) - return -M_PI_2; + return -M_PI_2_F; else if(a >= 1.0f) return M_PI_2_F; @@ -114,15 +114,20 @@ __device float svm_math(NodeMath type, float Fac1, float Fac2) return Fac; } +__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) { if(type == NODE_VECTOR_MATH_ADD) { *Vector = Vector1 + Vector2; - *Fac = (fabsf(Vector->x) + fabsf(Vector->y) + fabsf(Vector->z))/3.0f; + *Fac = average_fac(*Vector); } else if(type == NODE_VECTOR_MATH_SUBTRACT) { *Vector = Vector1 - Vector2; - *Fac = (fabsf(Vector->x) + fabsf(Vector->y) + fabsf(Vector->z))/3.0f; + *Fac = average_fac(*Vector); } else if(type == NODE_VECTOR_MATH_AVERAGE) { *Fac = len(Vector1 + Vector2); diff --git a/intern/cycles/kernel/svm/svm_texture.h b/intern/cycles/kernel/svm/svm_texture.h index d4765cca384..0bf0cf8c2cd 100644 --- a/intern/cycles/kernel/svm/svm_texture.h +++ b/intern/cycles/kernel/svm/svm_texture.h @@ -194,15 +194,15 @@ __device float noise_wave(NodeWaveType wave, float a) return 0.5f + 0.5f*sin(a); } else if(wave == NODE_WAVE_SAW) { - float b = 2*M_PI; + float b = 2.0f*M_PI_F; int n = (int)(a / b); a -= n*b; - if(a < 0) a += b; + if(a < 0.0f) a += b; return a / b; } else if(wave == NODE_WAVE_TRI) { - float b = 2*M_PI; + float b = 2.0f*M_PI_F; float rmax = 1.0f; return rmax - 2.0f*fabsf(floorf((a*(1.0f/b))+0.5f) - (a*(1.0f/b))); |