 #define T1 16
 #define T2 16
 #define Tcomm 16
-__global__ void sd_t_s1_7_kernel(int h1d,int h2d,int h3d,int p4d,int p6d,int p4ld_t2,int h1ld_t2,int h3ld_v2,int h2ld_v2,int p6ld_v2,int h3ld_t3,int h2ld_t3,int h1ld_t3,int p6ld_t3,int p4ld_t3,double *t3d, double *t2_d, double *v2_d,int p4, int total_x) {
-  int h1,h2,h3,p6;
+__global__ void sd_t_s1_7_kernel(size_t h1d,size_t h2d,size_t h3d,size_t p4d,size_t p6d,size_t p4ld_t2,size_t h1ld_t2,size_t h3ld_v2,size_t h2ld_v2,size_t p6ld_v2,size_t h3ld_t3,size_t h2ld_t3,size_t h1ld_t3,size_t p6ld_t3,size_t p4ld_t3,double *t3d, double *t2_d, double *v2_d,size_t p4, size_t total_x) {
+  size_t h1,h2,h3,p6;
   __shared__ double t2_shm[T1*4*Tcomm];
   
-  for(int i=threadIdx.x;i<h1d*p4d;i+=blockDim.x)
+  for(size_t i=threadIdx.x;i<h1d*p4d;i+=blockDim.x)
   if(i<h1d*p4d)
   t2_shm[i] = t2_d[i];
-  int rest_x=blockIdx.x;
-  int thread_x = T2*T1 * rest_x + threadIdx.x;
+  size_t rest_x=blockIdx.x;
+  size_t thread_x = T2*T1 * rest_x + threadIdx.x;
   rest_x = thread_x;
     __syncthreads();
 /* the following computation may need to happen inside the loop */
-  for(int i=0;i<total_x;i+=gridDim.x*blockDim.x)
+  for(size_t i=0;i<total_x;i+=gridDim.x*blockDim.x)
   {
     rest_x += i;
     h3=rest_x%h3d;

  p6ld_t3 = h1d * h2d * h3d;
  p4ld_t3 = p6d * h1d * h2d * h3d;
  p5ld_t3 = p4d * p6d * h1d * h2d * h3d;
  int total_x = h3d*h2d*p6d*p5d;
  dim3 dimBlock(T2*T1);dim3 dimGrid(DIV_UB(total_x,T2*T1), 1);
  for(i=0;i<nstreams;++i){
    sd_t_s1_5_kernel<<<dimGrid,dimBlock,0,streams[i]>>>(h1d,h2d,h3d,p4d,p5d,p6d,p4ld_t2,h1ld_t2,h3ld_v2,h2ld_v2,p6ld_v2,p5ld_v2,h3ld_t3,h2ld_t3,h1ld_t3,p6ld_t3,p5ld_t3,p4ld_t3,t3_s_d,t2_d,v2_d,i,total_x);

#define T1 16
#define T2 16
#define Tcomm 16
__global__ void sd_t_d2_9_kernel(size_t h1d,size_t h2d,size_t h3d,size_t p4d,size_t p5d,size_t p6d,size_t p7d,size_t p7ld_t2,size_t p4ld_t2,size_t h1ld_t2,size_t h2ld_t2,size_t p7ld_v2,size_t h3ld_v2,size_t p6ld_v2,size_t p5ld_v2,size_t h2ld_t3,size_t h3ld_t3,size_t h1ld_t3,size_t p4ld_t3,size_t p6ld_t3,size_t p5ld_t3,double *t3d, double *t2_d, double *v2_d,size_t unused_idx, size_t total_x, size_t total_y) {
  size_t h1_0,h1_1,h1_2,h1_3,h2_0,h2_1,h2_2,h2_3,h3_0,h3_1,h3_2,h3_3,p4_0,p4_1,p4_2,p4_3,p5_0,p5_1,p5_2,p5_3,p6_0,p6_1,p6_2,p6_3,p7;
  double a1,b1;
  double a2,b2;
  double a3,b3;
  double a4,b4;
  size_t in1_idxl,in2_idxl,p7l,p7T;
  __shared__ double t2_shm[4*T1][Tcomm];
  __shared__ double v2_shm[Tcomm][4*T2];
  size_t rest_x=blockIdx.x;
  size_t rest_y=blockIdx.y;
  size_t thread_x = T2*4 * rest_x + threadIdx.x;
  size_t thread_y = T1*4 * rest_y + threadIdx.y;
  in1_idxl=threadIdx.y;
  in2_idxl=threadIdx.x ;
  double tlocal1=0;
