Cr8-xyz/blender
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Cycles: amd opencl compatibility fixes.

Browse Source
This commit is contained in:
Brecht Van Lommel
2011-08-10 14:26:51 +00:00
parent abc601d10e
commit b98ccf6998
5 changed files with 16 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
intern/cycles/kernel/kernel.cl
View File

@@ -51,7 +51,6 @@ __kernel void kernel_ocl_path_trace(
int x = sx + get_global_id(0); int x = sx + get_global_id(0);
int y = sy + get_global_id(1); int y = sy + get_global_id(1);
int w = kernel_data.cam.width;
if(x < sx + sw && y < sy + sh) if(x < sx + sw && y < sy + sh)
kernel_path_trace(kg, buffer, rng_state, pass, x, y); kernel_path_trace(kg, buffer, rng_state, pass, x, y);

4
intern/cycles/kernel/kernel_shader.h
View File

@@ -140,14 +140,18 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd,
#endif #endif
/* detect instancing, for non-instanced the object index is -object-1 */ /* detect instancing, for non-instanced the object index is -object-1 */
#ifdef __INSTANCING__
bool instanced = false; bool instanced = false;
if(sd->prim != ~0) { if(sd->prim != ~0) {
if(sd->object >= 0) if(sd->object >= 0)
instanced = true; instanced = true;
else else
#endif
sd->object = -sd->object-1; sd->object = -sd->object-1;
#ifdef __INSTANCING__
} }
#endif
/* smooth normal */ /* smooth normal */
if(sd->shader < 0) { if(sd->shader < 0) {

4
intern/cycles/kernel/svm/bsdf_ward.h
View File

@@ -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; h = h.x * X + h.y * Y + h.z * m_N;
// generate the final sample // generate the final sample
float oh = dot(h, sd->I); float oh = dot(h, sd->I);
omega_in->x = 2 * oh * h.x - sd->I.x; *omega_in = 2.0f * oh * h - sd->I;
omega_in->y = 2 * oh * h.y - sd->I.y;
omega_in->z = 2 * oh * h.z - sd->I.z;
if(dot(sd->Ng, *omega_in) > 0) { if(dot(sd->Ng, *omega_in) > 0) {
float cosNI = dot(m_N, *omega_in); float cosNI = dot(m_N, *omega_in);
if(cosNI > 0) { if(cosNI > 0) {

11
intern/cycles/kernel/svm/svm_math.h
View File

@@ -21,7 +21,7 @@ CCL_NAMESPACE_BEGIN
__device float safe_asinf(float a) __device float safe_asinf(float a)
{ {
if(a <= -1.0f) if(a <= -1.0f)
return -M_PI_2; return -M_PI_2_F;
else if(a >= 1.0f) else if(a >= 1.0f)
return M_PI_2_F; return M_PI_2_F;
@@ -114,15 +114,20 @@ __device float svm_math(NodeMath type, float Fac1, float Fac2)
return Fac; 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) __device void svm_vector_math(float *Fac, float3 *Vector, NodeVectorMath type, float3 Vector1, float3 Vector2)
{ {
if(type == NODE_VECTOR_MATH_ADD) { if(type == NODE_VECTOR_MATH_ADD) {
*Vector = Vector1 + Vector2; *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) { else if(type == NODE_VECTOR_MATH_SUBTRACT) {
*Vector = Vector1 - Vector2; *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) { else if(type == NODE_VECTOR_MATH_AVERAGE) {
*Fac = len(Vector1 + Vector2); *Fac = len(Vector1 + Vector2);

6
intern/cycles/kernel/svm/svm_texture.h
View File

@@ -194,15 +194,15 @@ __device float noise_wave(NodeWaveType wave, float a)
return 0.5f + 0.5f*sin(a); return 0.5f + 0.5f*sin(a);
} }
else if(wave == NODE_WAVE_SAW) { else if(wave == NODE_WAVE_SAW) {
float b = 2*M_PI; float b = 2.0f*M_PI_F;
int n = (int)(a / b); int n = (int)(a / b);
a -= n*b; a -= n*b;
if(a < 0) a += b; if(a < 0.0f) a += b;
return a / b; return a / b;
} }
else if(wave == NODE_WAVE_TRI) { else if(wave == NODE_WAVE_TRI) {
float b = 2*M_PI; float b = 2.0f*M_PI_F;
float rmax = 1.0f; float rmax = 1.0f;
return rmax - 2.0f*fabsf(floorf((a*(1.0f/b))+0.5f) - (a*(1.0f/b))); return rmax - 2.0f*fabsf(floorf((a*(1.0f/b))+0.5f) - (a*(1.0f/b)));