[Bf-blender-cvs] [36b80b1] master: Cycles: Fix/workaround for missing AA on Pascal (GTX10x0) cards

[Sergey Sharybin]

Commit: 36b80b185f2c01e72ba251b852b517b02c806c73
Author: Sergey Sharybin
Date:   Tue Aug 23 16:32:09 2016 +0200
Branches: master
https://developer.blender.org/rB36b80b185f2c01e72ba251b852b517b02c806c73

Cycles: Fix/workaround for missing AA on Pascal (GTX10x0) cards

===================================================================

M	intern/cycles/kernel/kernel_random.h

===================================================================

diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h
index 731dc04..b534d99 100644
--- a/intern/cycles/kernel/kernel_random.h
+++ b/intern/cycles/kernel/kernel_random.h
@@ -132,7 +132,13 @@ ccl_device_inline float path_rng_1D(KernelGlobals *kg, ccl_addr_space RNG *rng,
 #endif
 }
 
-ccl_device_inline void path_rng_2D(KernelGlobals *kg, ccl_addr_space RNG *rng, int sample, int num_samples, int dimension, float *fx, float *fy)
+/* Temporary workaround for Pascal cards, otherwise AA does not work properly. */
+#if defined(__KERNEL_GPU__) && __CUDA_ARCH__ >= 600
+__device__ __forceinline__
+#else
+ccl_device_inline
+#endif
+void path_rng_2D(KernelGlobals *kg, ccl_addr_space RNG *rng, int sample, int num_samples, int dimension, float *fx, float *fy)
 {
 #ifdef __CMJ__
 	if(kernel_data.integrator.sampling_pattern == SAMPLING_PATTERN_CMJ) {