From a401eb5a3b15dbb864e7b47067ef9c233cba339d Mon Sep 17 00:00:00 2001
From: Evghenii <egaburov@dds.nl>
Date: Wed, 8 Jan 2014 15:38:02 +0100
Subject: [PATCH] added DeviceCacheSettings

---
 examples_ptx/deferred/kernels.cu |  4 ++--
 examples_ptx/ptxcc/PTXParser.h   | 10 ++++++++++
 2 files changed, 12 insertions(+), 2 deletions(-)

diff --git a/examples_ptx/deferred/kernels.cu b/examples_ptx/deferred/kernels.cu
index f6bb02f2..3b353acf 100644
--- a/examples_ptx/deferred/kernels.cu
+++ b/examples_ptx/deferred/kernels.cu
@@ -133,7 +133,7 @@ struct Uniform
     data[chunkIdx] = shdata[programIndex];
   }
 };
-#elif 1
+#elif 0
 template<typename T, int N>
 struct Uniform
 {
@@ -174,7 +174,7 @@ __shared__ int shdata_full[4*MAX_LIGHTS];
 template<typename T, int N>
 struct Uniform
 {
-  volatile T *shdata;
+  /* volatile */ T *shdata;
 
   __device__ Uniform()
   {
diff --git a/examples_ptx/ptxcc/PTXParser.h b/examples_ptx/ptxcc/PTXParser.h
index f0b6b055..5ae0d903 100644
--- a/examples_ptx/ptxcc/PTXParser.h
+++ b/examples_ptx/ptxcc/PTXParser.h
@@ -141,6 +141,16 @@ namespace parser
         s << printArgumentList();
         s << "\n )\n";
         s << "{\n   ";
+//        s << " cudaFuncSetCacheConfig (" << calleeName << ", ";
+        s << " cudaDeviceSetCacheConfig (";
+#if 1
+        s << " cudaFuncCachePreferEqual ";
+#elif 1
+        s << " cudaFuncCachePreferL1 ";
+#else
+        s << " cudaFuncCachePreferShared ";
+#endif
+        s << ");\n";
         s << calleeName;
         s << "<<<1,32>>>(\n";
         s << printArgumentList(false);