From e1ef902058149b6feee96d87e58b26582c522b2d Mon Sep 17 00:00:00 2001
From: Thomas Dinges <blender@dingto.org>
Date: Sat, 17 Feb 2018 22:19:54 +0100
Subject: [PATCH] Cycles: Remove fermi related defines from the code.

Did not touch Texture related defines, that comes next.
---
 intern/cycles/kernel/geom/geom_curve_intersect.h  | 10 ++--------
 intern/cycles/kernel/kernel_compat_cuda.h         |  5 +----
 intern/cycles/kernel/kernels/cuda/kernel_config.h | 14 +-------------
 intern/cycles/util/util_math_intersect.h          |  7 +------
 4 files changed, 5 insertions(+), 31 deletions(-)

diff --git a/intern/cycles/kernel/geom/geom_curve_intersect.h b/intern/cycles/kernel/geom/geom_curve_intersect.h
index e9a149ea1ab..7f24aea5d28 100644
--- a/intern/cycles/kernel/geom/geom_curve_intersect.h
+++ b/intern/cycles/kernel/geom/geom_curve_intersect.h
@@ -18,12 +18,6 @@ CCL_NAMESPACE_BEGIN
 
 #ifdef __HAIR__
 
-#if defined(__KERNEL_CUDA__) && (__CUDA_ARCH__ < 300)
-#  define ccl_device_curveintersect ccl_device
-#else
-#  define ccl_device_curveintersect ccl_device_forceinline
-#endif
-
 #ifdef __KERNEL_SSE2__
 ccl_device_inline ssef transform_point_T3(const ssef t[3], const ssef &a)
 {
@@ -32,7 +26,7 @@ ccl_device_inline ssef transform_point_T3(const ssef t[3], const ssef &a)
 #endif
 
 /* On CPU pass P and dir by reference to aligned vector. */
-ccl_device_curveintersect bool cardinal_curve_intersect(
+ccl_device_forceinline bool cardinal_curve_intersect(
         KernelGlobals *kg,
         Intersection *isect,
         const float3 ccl_ref P,
@@ -505,7 +499,7 @@ ccl_device_curveintersect bool cardinal_curve_intersect(
 	return hit;
 }
 
-ccl_device_curveintersect bool curve_intersect(KernelGlobals *kg,
+ccl_device_forceinline bool curve_intersect(KernelGlobals *kg,
                                                Intersection *isect,
                                                float3 P,
                                                float3 direction,
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index 900f7fe6a2c..1daa7f0db16 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -50,10 +50,7 @@ __device__ half __float2half(const float f)
 /* Qualifier wrappers for different names on different devices */
 
 #define ccl_device  __device__ __inline__
-#if __CUDA_ARCH__ < 300
-#  define ccl_device_inline  __device__ __inline__
-#  define ccl_device_forceinline  __device__ __forceinline__
-#elif __CUDA_ARCH__ < 500
+#if __CUDA_ARCH__ < 500
 #  define ccl_device_inline  __device__ __forceinline__
 #  define ccl_device_forceinline  __device__ __forceinline__
 #else
diff --git a/intern/cycles/kernel/kernels/cuda/kernel_config.h b/intern/cycles/kernel/kernels/cuda/kernel_config.h
index 94f59ff38d9..f3d0d721c5c 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel_config.h
+++ b/intern/cycles/kernel/kernels/cuda/kernel_config.h
@@ -16,20 +16,8 @@
 
 /* device data taken from CUDA occupancy calculator */
 
-/* 2.0 and 2.1 */
-#if __CUDA_ARCH__ == 200 || __CUDA_ARCH__ == 210
-#  define CUDA_MULTIPRESSOR_MAX_REGISTERS 32768
-#  define CUDA_MULTIPROCESSOR_MAX_BLOCKS 8
-#  define CUDA_BLOCK_MAX_THREADS 1024
-#  define CUDA_THREAD_MAX_REGISTERS 63
-
-/* tunable parameters */
-#  define CUDA_THREADS_BLOCK_WIDTH 16
-#  define CUDA_KERNEL_MAX_REGISTERS 32
-#  define CUDA_KERNEL_BRANCHED_MAX_REGISTERS 40
-
 /* 3.0 and 3.5 */
-#elif __CUDA_ARCH__ == 300 || __CUDA_ARCH__ == 350
+#if __CUDA_ARCH__ == 300 || __CUDA_ARCH__ == 350
 #  define CUDA_MULTIPRESSOR_MAX_REGISTERS 65536
 #  define CUDA_MULTIPROCESSOR_MAX_BLOCKS 16
 #  define CUDA_BLOCK_MAX_THREADS 1024
diff --git a/intern/cycles/util/util_math_intersect.h b/intern/cycles/util/util_math_intersect.h
index 498c21b9706..61ddcc38f50 100644
--- a/intern/cycles/util/util_math_intersect.h
+++ b/intern/cycles/util/util_math_intersect.h
@@ -79,12 +79,7 @@ ccl_device bool ray_aligned_disk_intersect(
 	return true;
 }
 
-#if defined(__KERNEL_CUDA__) && __CUDA_ARCH__ < 300
-ccl_device_inline
-#else
-ccl_device_forceinline
-#endif
-bool ray_triangle_intersect(
+ccl_device_forceinline bool ray_triangle_intersect(
         float3 ray_P, float3 ray_dir, float ray_t,
 #if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
         const ssef *ssef_verts,