 #ifndef M_PI
 #define M_PI (3.1415926535897932384626)
 #endif
+
+  
+#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
 
 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_));

  
#define GET(x) (x)

#define diffusion_backward()                                            \
  do {                                                                  \
    sb[ps] = s2;                                                        \
