   
 #define GET(x) (x)
 
-#define bdimx (BLOCK_X)
-#define bdimy (BLOCK_Y)
-#define SHIFT3(x, y, z) x = y; y = z
-#define SHIFT4(x, y, z, k) x = y; y = z; z = k
-
-
 #define diffusion_backward()                                            \
   do {                                                                  \
     sb[ps] = s2;                                                        \

#ifndef M_PI
#define M_PI (3.1415926535897932384626)
#endif

using std::vector;
using std::string;

  CUDA_SAFE_CALL(cudaMalloc((void**)&f1_d, s));
  CUDA_SAFE_CALL(cudaMalloc((void**)&f2_d, s));
  //CUDA_SAFE_CALL(cudaFuncSetCacheConfig(diffusion_kernel_shared,
                                        //cudaFuncCachePreferL1));
  CUDA_SAFE_CALL(cudaFuncSetCacheConfig(diffusion_kernel_shared6,
                                        cudaFuncCachePreferShared));
  CUDA_SAFE_CALL(cudaEventCreate(&ev1_));
