[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [39256] branches/cycles/intern/cycles/ kernel: Cycles: amd opencl compatibility fixes.

[Brecht Van Lommel]

Revision: 39256
          http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=39256
Author:   blendix
Date:     2011-08-10 14:26:51 +0000 (Wed, 10 Aug 2011)
Log Message:
-----------
Cycles: amd opencl compatibility fixes.

Modified Paths:
--------------
    branches/cycles/intern/cycles/kernel/kernel.cl
    branches/cycles/intern/cycles/kernel/kernel_shader.h
    branches/cycles/intern/cycles/kernel/svm/bsdf_ward.h
    branches/cycles/intern/cycles/kernel/svm/svm_math.h
    branches/cycles/intern/cycles/kernel/svm/svm_texture.h

Modified: branches/cycles/intern/cycles/kernel/kernel.cl
===================================================================
--- branches/cycles/intern/cycles/kernel/kernel.cl	2011-08-10 13:59:03 UTC (rev 39255)
+++ branches/cycles/intern/cycles/kernel/kernel.cl	2011-08-10 14:26:51 UTC (rev 39256)
@@ -51,7 +51,6 @@
 
 	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);

Modified: branches/cycles/intern/cycles/kernel/kernel_shader.h
===================================================================
--- branches/cycles/intern/cycles/kernel/kernel_shader.h	2011-08-10 13:59:03 UTC (rev 39255)
+++ branches/cycles/intern/cycles/kernel/kernel_shader.h	2011-08-10 14:26:51 UTC (rev 39256)
@@ -140,14 +140,18 @@
 #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) {

Modified: branches/cycles/intern/cycles/kernel/svm/bsdf_ward.h
===================================================================
--- branches/cycles/intern/cycles/kernel/svm/bsdf_ward.h	2011-08-10 13:59:03 UTC (rev 39255)
+++ branches/cycles/intern/cycles/kernel/svm/bsdf_ward.h	2011-08-10 14:26:51 UTC (rev 39256)
@@ -166,9 +166,7 @@
 		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) {

Modified: branches/cycles/intern/cycles/kernel/svm/svm_math.h
===================================================================
--- branches/cycles/intern/cycles/kernel/svm/svm_math.h	2011-08-10 13:59:03 UTC (rev 39255)
+++ branches/cycles/intern/cycles/kernel/svm/svm_math.h	2011-08-10 14:26:51 UTC (rev 39256)
@@ -21,7 +21,7 @@
 __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 @@
 	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);

Modified: branches/cycles/intern/cycles/kernel/svm/svm_texture.h
===================================================================
--- branches/cycles/intern/cycles/kernel/svm/svm_texture.h	2011-08-10 13:59:03 UTC (rev 39255)
+++ branches/cycles/intern/cycles/kernel/svm/svm_texture.h	2011-08-10 14:26:51 UTC (rev 39256)
@@ -194,15 +194,15 @@
     	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)));