[Bf-blender-cvs] [11f3c94] experimental-build: Cycles: Use array data storage for Fermi cards as well.

Thomas Dinges noreply at git.blender.org
Mon May 16 14:58:03 CEST 2016


Commit: 11f3c9470d3b11180aa7ce14ac89eb0829353469
Author: Thomas Dinges
Date:   Mon May 16 14:57:09 2016 +0200
Branches: experimental-build
https://developer.blender.org/rB11f3c9470d3b11180aa7ce14ac89eb0829353469

Cycles: Use array data storage for Fermi cards as well.

This is for testing purposes, to see if we can use array storage on all devices.

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

M	intern/cycles/device/device_cuda.cpp
M	intern/cycles/kernel/kernel_compat_cuda.h
M	intern/cycles/kernel/kernel_globals.h

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

diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 69015e7..2ce2c2e 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -86,9 +86,7 @@ public:
 	CUmodule cuModule;
 	map<device_ptr, bool> tex_interp_map;
 	int cuDevId;
-	int cuDevArchitecture;
 	bool first_error;
-	bool use_texture_storage;
 
 	struct PixelMem {
 		GLuint cuPBO;
@@ -176,7 +174,6 @@ public:
 	{
 		first_error = true;
 		background = background_;
-		use_texture_storage = true;
 
 		cuDevId = info.num;
 		cuDevice = 0;
@@ -207,15 +204,6 @@ public:
 		if(cuda_error_(result, "cuCtxCreate"))
 			return;
 
-		int major, minor;
-		cuDeviceComputeCapability(&major, &minor, cuDevId);
-		cuDevArchitecture = major*100 + minor*10;
-
-		/* In order to use full 6GB of memory on Titan cards, use arrays instead
-		 * of textures. On earlier cards this seems slower, but on Titan it is
-		 * actually slightly faster in tests. */
-		use_texture_storage = (cuDevArchitecture < 300);
-
 		cuda_pop_context();
 	}
 
@@ -495,7 +483,7 @@ public:
 		CUarray_format_enum format;
 		size_t dsize = datatype_size(mem.data_type);
 		size_t size = mem.memory_size();
-		bool use_texture = (interpolation != INTERPOLATION_NONE) || use_texture_storage;
+		bool use_texture = (interpolation != INTERPOLATION_NONE);
 
 		if(use_texture) {
 
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index d10d325..167c994 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -67,18 +67,7 @@ typedef texture<uchar4, 2, cudaReadModeNormalizedFloat> texture_image_uchar4;
 
 /* Macros to handle different memory storage on different devices */
 
-/* In order to use full 6GB of memory on Titan cards, use arrays instead
- * of textures. On earlier cards this seems slower, but on Titan it is
- * actually slightly faster in tests. */
-#if __CUDA_ARCH__ < 300
-#  define __KERNEL_CUDA_TEX_STORAGE__
-#endif
-
-#ifdef __KERNEL_CUDA_TEX_STORAGE__
-#  define kernel_tex_fetch(t, index) tex1Dfetch(t, index)
-#else
-#  define kernel_tex_fetch(t, index) t[(index)]
-#endif
+#define kernel_tex_fetch(t, index) t[(index)]
 #define kernel_tex_image_interp(t, x, y) tex2D(t, x, y)
 #define kernel_tex_image_interp_3d(t, x, y, z) tex3D(t, x, y, z)
 
diff --git a/intern/cycles/kernel/kernel_globals.h b/intern/cycles/kernel/kernel_globals.h
index c44ea1b..cb8ce4c 100644
--- a/intern/cycles/kernel/kernel_globals.h
+++ b/intern/cycles/kernel/kernel_globals.h
@@ -65,12 +65,9 @@ typedef struct KernelGlobals {
 __constant__ KernelData __data;
 typedef struct KernelGlobals {} KernelGlobals;
 
-#  ifdef __KERNEL_CUDA_TEX_STORAGE__
-#    define KERNEL_TEX(type, ttype, name) ttype name;
-#  else
-#    define KERNEL_TEX(type, ttype, name) const __constant__ __device__ type *name;
-#  endif
+#  define KERNEL_TEX(type, ttype, name) const __constant__ __device__ type *name;
 #  define KERNEL_IMAGE_TEX(type, ttype, name) ttype name;
+
 #  include "kernel_textures.h"
 
 #endif  /* __KERNEL_CUDA__ */




More information about the Bf-blender-cvs mailing list