some examples tuning

examples_ptx/aobench/Makefile_gpu

@@ -4,7 +4,7 @@ CU_SRC=ao.cu
 CXX_SRC=ao.cpp  ao_serial.cpp
 PTXCC_REGMAX=64
 
-# LLVM_GPU=1
+LLVM_GPU=1
 NVVM_GPU=1
 
 include ../common_gpu.mk

examples_ptx/deferred/kernels.cu

@@ -133,7 +133,7 @@ struct Uniform
     data[chunkIdx] = shdata[programIndex];
   }
 };
-#elif 0
+#elif 1
 template<typename T, int N>
 struct Uniform
 {

examples_ptx/mandelbrot_tasks/mandelbrot_tasks.cu

@@ -35,10 +35,10 @@
 
 __device__
 static inline int
-mandel(float c_re, float c_im, int count) {
+mandel(float c_re, float c_im, float count) {
     float z_re = c_re, z_im = c_im;
-    int i;
-    for (i = 0; i < count; ++i) {
+    float i = 0.0f;
+    for (; i < count; ++i) {
         if (z_re * z_re + z_im * z_im > 4.0f)
             break;
 

examples_ptx/mandelbrot_tasks/mandelbrot_tasks.ispc

@@ -32,10 +32,10 @@
 */
 
 static inline int
-mandel(float c_re, float c_im, int count) {
+mandel(float c_re, float c_im, float count) {
   float z_re = c_re, z_im = c_im;
-  int i;
-  for (i = 0; i < count; ++i) {
+  float i = 0.0f;
+  for (;i < count; ++i) {
     if (z_re * z_re + z_im * z_im > 4.0f)
       break;
 

examples_ptx/options/Makefile_gpu

@@ -2,7 +2,7 @@ PROG=options
 ISPC_SRC=options.ispc
 CU_SRC=options.cu
 CXX_SRC=options.cpp  options_serial.cpp
-PTXCC_REGMAX=32
+PTXCC_REGMAX=128
 
 
 LLVM_GPU=1

examples_ptx/options/options.cu

@@ -327,6 +327,8 @@ __host__ void
 binomial_put_ispc_tasks( float Sa[],  float Xa[],  float Ta[],
                           float ra[],  float va[], 
                           float result[],  int count) {
+
+  cudaDeviceSetCacheConfig (cudaFuncCachePreferL1);
   binomial_put_ispc_tasks___export<<<1,32>>>(Sa,Xa,Ta,ra,va,result,count);
   cudaDeviceSynchronize();
 }