From 1dcd7db73d13443c59dd824abd9cacbf6bc88997 Mon Sep 17 00:00:00 2001
From: Brecht Van Lommel <brechtvanlommel@gmail.com>
Date: Sun, 18 Feb 2018 00:51:46 +0100
Subject: [PATCH] Code cleanup: remove some more unused code after recent CUDA
 changes.

---
 intern/cycles/device/device_cuda.cpp          |  2 --
 .../cycles/kernel/geom/geom_curve_intersect.h | 22 +++++++++----------
 intern/cycles/kernel/kernel_globals.h         |  2 --
 intern/cycles/kernel/kernel_textures.h        |  6 -----
 intern/cycles/kernel/kernels/cpu/kernel.cpp   |  1 -
 intern/cycles/kernel/svm/svm.h                |  3 +--
 intern/cycles/render/image.cpp                |  2 +-
 7 files changed, 13 insertions(+), 25 deletions(-)

diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 42e78e50540..d28080c667a 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -1069,8 +1069,6 @@ public:
 		}
 
 		/* Image Texture Storage */
-		CUtexref texref = NULL;
-
 		CUarray_format_enum format;
 		switch(mem.data_type) {
 			case TYPE_UCHAR: format = CU_AD_FORMAT_UNSIGNED_INT8; break;
diff --git a/intern/cycles/kernel/geom/geom_curve_intersect.h b/intern/cycles/kernel/geom/geom_curve_intersect.h
index 7f24aea5d28..faf3e3cdf2b 100644
--- a/intern/cycles/kernel/geom/geom_curve_intersect.h
+++ b/intern/cycles/kernel/geom/geom_curve_intersect.h
@@ -500,17 +500,17 @@ ccl_device_forceinline bool cardinal_curve_intersect(
 }
 
 ccl_device_forceinline bool curve_intersect(KernelGlobals *kg,
-                                               Intersection *isect,
-                                               float3 P,
-                                               float3 direction,
-                                               uint visibility,
-                                               int object,
-                                               int curveAddr,
-                                               float time,
-                                               int type,
-                                               uint *lcg_state,
-                                               float difl,
-                                               float extmax)
+                                            Intersection *isect,
+                                            float3 P,
+                                            float3 direction,
+                                            uint visibility,
+                                            int object,
+                                            int curveAddr,
+                                            float time,
+                                            int type,
+                                            uint *lcg_state,
+                                            float difl,
+                                            float extmax)
 {
 	/* define few macros to minimize code duplication for SSE */
 #ifndef __KERNEL_SSE2__
diff --git a/intern/cycles/kernel/kernel_globals.h b/intern/cycles/kernel/kernel_globals.h
index 97d4726407b..74cfacb5bc1 100644
--- a/intern/cycles/kernel/kernel_globals.h
+++ b/intern/cycles/kernel/kernel_globals.h
@@ -47,7 +47,6 @@ struct VolumeStep;
 
 typedef struct KernelGlobals {
 #  define KERNEL_TEX(type, name) texture<type> name;
-#  define KERNEL_IMAGE_TEX(type, ttype, name)
 #  include "kernel/kernel_textures.h"
 
 	KernelData __data;
@@ -93,7 +92,6 @@ typedef struct KernelGlobals {
 } KernelGlobals;
 
 #  define KERNEL_TEX(type, name) const __constant__ __device__ type *name;
-#  define KERNEL_IMAGE_TEX(type, ttype, name) ttype name;
 #  include "kernel/kernel_textures.h"
 
 #endif  /* __KERNEL_CUDA__ */
diff --git a/intern/cycles/kernel/kernel_textures.h b/intern/cycles/kernel/kernel_textures.h
index 7735a7bb355..74b659557e5 100644
--- a/intern/cycles/kernel/kernel_textures.h
+++ b/intern/cycles/kernel/kernel_textures.h
@@ -18,10 +18,6 @@
 #  define KERNEL_TEX(type, name)
 #endif
 
-#ifndef KERNEL_IMAGE_TEX
-#  define KERNEL_IMAGE_TEX(type, ttype, name)
-#endif
-
 /* bvh */
 KERNEL_TEX(float4, __bvh_nodes)
 KERNEL_TEX(float4, __bvh_leaf_nodes)
@@ -82,6 +78,4 @@ KERNEL_TEX(uint, __sobol_directions)
 KERNEL_TEX(TextureInfo, __texture_info)
 
 #undef KERNEL_TEX
-#undef KERNEL_IMAGE_TEX
-
 
diff --git a/intern/cycles/kernel/kernels/cpu/kernel.cpp b/intern/cycles/kernel/kernels/cpu/kernel.cpp
index aa67262f36b..de487f6123f 100644
--- a/intern/cycles/kernel/kernels/cpu/kernel.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel.cpp
@@ -85,7 +85,6 @@ void kernel_tex_copy(KernelGlobals *kg,
 		kg->tname.data = (type*)mem; \
 		kg->tname.width = size; \
 	}
-#define KERNEL_IMAGE_TEX(type, tname)
 #include "kernel/kernel_textures.h"
 	else {
 		assert(0);
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index a8f99d23b7d..fae9f783483 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -30,8 +30,7 @@
  * in local memory on the GPU, as it would take too many register and indexes in
  * ways not known at compile time. This seems the only solution even though it
  * may be slow, with two positive factors. If the same shader is being executed,
- * memory access will be coalesced, and on fermi cards, memory will actually be
- * cached.
+ * memory access will be coalesced and cached.
  *
  * The result of shader execution will be a single closure. This means the
  * closure type, associated label, data and weight. Sampling from multiple
diff --git a/intern/cycles/render/image.cpp b/intern/cycles/render/image.cpp
index 29f37afd676..dbe15a67b9e 100644
--- a/intern/cycles/render/image.cpp
+++ b/intern/cycles/render/image.cpp
@@ -254,7 +254,7 @@ int ImageManager::add_image(const string& filename,
 	/* Check whether it's a float texture. */
 	is_float = (type == IMAGE_DATA_TYPE_FLOAT || type == IMAGE_DATA_TYPE_FLOAT4);
 
-	/* No half textures on OpenCL, use available slots */
+	/* No half textures on OpenCL, use full float instead. */
 	if(!has_half_images) {
 		if(type == IMAGE_DATA_TYPE_HALF4) {
 			type = IMAGE_DATA_TYPE_FLOAT4;