Code cleanup: remove some more unused code after recent CUDA changes.

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;

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__

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__ */

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
-
 

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);

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

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;