[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