/****************************************************************************\
*      --- Practical Course: GPU Programming in Computer Vision ---
*
* time:    winter term 2012/13 / March 11-18, 2013
*
* project: diffusion
* file:    diffusion.cu
*
* 
\******* PLEASE ENTER YOUR CORRECT STUDENT LOGIN, NAME AND ID BELOW *********/
const char* studentLogin = "p116";
const char* studentName  = "Arash Bakhtiari";
const int   studentID    = 03625141;
/****************************************************************************\
*
* In this file the following methods have to be edited or completed:
*
* diffuse_linear_isotrop_shared(const float  *d_input, ... )
* diffuse_linear_isotrop_shared(const float3 *d_input, ... )
* diffuse_nonlinear_isotrop_shared(const float  *d_input, ... )
* diffuse_nonlinear_isotrop_shared(const float3 *d_input, ... )
* compute_tv_diffusivity_shared
* compute_tv_diffusivity_joined_shared
* compute_tv_diffusivity_separate_shared
* jacobi_shared(float  *d_output, ... )
* jacobi_shared(float3 *d_output, ... )
* sor_shared(float  *d_output, ... )
* sor_shared(float3 *d_output, ... )
*
\****************************************************************************/


#define DIFF_BW 16
#define DIFF_BH 16

#define TV_EPSILON 0.1f


#include "diffusion.cuh"



const char* getStudentLogin() { return studentLogin; };
const char* getStudentName()  { return studentName; };
int         getStudentID()    { return studentID; };
bool checkStudentData() { return strcmp(studentLogin, "p010") != 0 && strcmp(studentName, "John Doe") != 0 && studentID != 1234567; };
bool checkStudentNameAndID() { return strcmp(studentName, "John Doe") != 0 && studentID != 1234567; };


//----------------------------------------------------------------------------
// Linear Diffusion
//----------------------------------------------------------------------------


// mode 0 gray: linear diffusion
__global__ void diffuse_linear_isotrop_shared(
  const float *d_input,
  float *d_output,
  float timeStep, 
  int nx, int ny,
  size_t pitch)
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const int idx = y*pitch + x;
  //d_output[idx] = 0;
  __shared__ float u[DIFF_BW+2][DIFF_BH+2];

  // load data into shared memory
  if (x < nx && y < ny) {

    u[tx][ty] = d_input[idx];

    if (x == 0)  u[0][ty] = u[tx][ty];
    else if (threadIdx.x == 0) u[0][ty] = d_input[idx-1];
    if (x == nx-1) u[tx+1][ty] = u[tx][ty];
    else if (threadIdx.x == blockDim.x-1) u[tx+1][ty] = d_input[idx+1];

    if (y == 0)  u[tx][0] = u[tx][ty];
    else if (threadIdx.y == 0) u[tx][0] = d_input[idx-pitch];
    if (y == ny-1) u[tx][ty+1] = u[tx][ty];
    else if (threadIdx.y == blockDim.y-1) u[tx][ty+1] = d_input[idx+pitch];
  }

  __syncthreads();

  // ### implement me ###
  if (x < nx && y < ny) {
	d_output[idx] = u[tx][ty] + timeStep * ( u[tx + 1][ty]
			+ u[tx - 1][ty] + u[tx][ty + 1]
			+ u[tx][ty - 1] - 4 * u[tx][ty]);
  }
}



// mode 0 interleaved: linear diffusion
__global__ void diffuse_linear_isotrop_shared
(
 const float3 *d_input,
 float3 *d_output,
 float timeStep,
 int nx, int ny,
 size_t pitchBytes
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const char* imgP = (char*)d_input + y*pitchBytes + x*sizeof(float3);

  __shared__ float3 u[DIFF_BW+2][DIFF_BH+2];
  float3 imgValue;

  // load data into shared memory
  if (x < nx && y < ny) {

    imgValue = *( (float3*)imgP );
    u[tx][ty] = imgValue;

    if (x == 0)  u[0][ty] = imgValue;
    else if (threadIdx.x == 0) u[0][ty] = *( ((float3*)imgP)-1 );
    if (x == nx-1) u[tx+1][ty] = imgValue;
    else if (threadIdx.x == blockDim.x-1) u[tx+1][ty] = *( ((float3*)imgP)+1 );

    if (y == 0)  u[tx][0] = imgValue;
    else if (threadIdx.y == 0) u[tx][0] = *( (float3*)(imgP-pitchBytes) );
    if (y == ny-1) u[tx][ty+1] = imgValue;
    else if (threadIdx.y == blockDim.y-1) u[tx][ty+1] = *( (float3*)(imgP+pitchBytes) );
  }

  __syncthreads();

  float3 tmpValue;
  tmpValue.x = u[tx][ty].x + timeStep * (u[tx + 1][ty].x + u[tx - 1][ty].x
			+ u[tx][ty + 1].x + u[tx][ty - 1].x - 4 * u[tx][ty].x);
  
  tmpValue.y = u[tx][ty].y + timeStep * (u[tx + 1][ty].y + u[tx - 1][ty].y
			+ u[tx][ty + 1].y + u[tx][ty - 1].y - 4 * u[tx][ty].y);

  tmpValue.z = u[tx][ty].z + timeStep * (u[tx + 1][ty].z + u[tx - 1][ty].z
			+ u[tx][ty + 1].z + u[tx][ty - 1].z - 4 * u[tx][ty].z);
  if (x < nx && y < ny) 
  *((float3*)(((char*)d_output) + y*pitchBytes) + x) = tmpValue;
  
}




//----------------------------------------------------------------------------
// Non-linear Diffusion - explicit scheme
//----------------------------------------------------------------------------



// mode 1 gray: nonlinear diffusion
__global__ void diffuse_nonlinear_isotrop_shared
(
 const float *d_input,
 const float *d_diffusivity,
 float *d_output,
 float timeStep,
 int   nx,
 int   ny,
 size_t   pitch
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const int idx = y*pitch + x;

  __shared__ float u[DIFF_BW+2][DIFF_BH+2];
  __shared__ float g[DIFF_BW+2][DIFF_BH+2];


  // load data into shared memory
  if (x < nx && y < ny) {
    u[tx][ty] = d_input[idx];
    g[tx][ty] = d_diffusivity[idx];

    if (x == 0) {
      u[0][ty] = u[tx][ty];
      g[0][ty] = g[tx][ty];
    }
    else if (threadIdx.x == 0) {
      u[0][ty] = d_input[idx-1];
      g[0][ty] = d_diffusivity[idx-1];
    }
      
    if (x == nx-1) {
      u[tx+1][ty] = u[tx][ty];
      g[tx+1][ty] = g[tx][ty];
    }
    else if (threadIdx.x == blockDim.x-1) {
      u[tx+1][ty] = d_input[idx+1];
      g[tx+1][ty] = d_diffusivity[idx+1];
    }


    if (y == 0) {
      u[tx][0] = u[tx][ty];
      g[tx][0] = g[tx][ty];
    }
    else if (threadIdx.y == 0) {
      u[tx][0] = d_input[idx-pitch];
      g[tx][0] = d_diffusivity[idx-pitch];
    }
      
    if (y == ny-1) {
      u[tx][ty+1] = u[tx][ty];
      g[tx][ty+1] = g[tx][ty];
    } 
    else if (threadIdx.y == blockDim.y-1) {
      u[tx][ty+1] = d_input[idx+pitch];
      g[tx][ty+1] = d_diffusivity[idx+pitch];
    }
  }

  __syncthreads();

  float phiR = 0.5 * (g[tx+1][ty] + g[tx][ty]);
  float phiL = 0.5 * (g[tx-1][ty] + g[tx][ty]);
  float phiU = 0.5 * (g[tx][ty+1] + g[tx][ty]);
  float phiD = 0.5 * (g[tx][ty-1] + g[tx][ty]);
  
  // ### implement me ###
	if (x < nx && y < ny) {
		d_output[idx] = u[tx][ty] + timeStep * (
						u[tx + 1][ty]*phiR +
						u[tx - 1][ty]*phiL + 
						u[tx][ty + 1]*phiU + 
						u[tx][ty - 1]*phiD - 
						u[tx][ty]*(phiR+phiL+phiU+phiD) );
	}

}



// mode 1 interleaved: nonlinear diffusion
__global__ void diffuse_nonlinear_isotrop_shared
(
 const float3 *d_input,
 const float3 *d_diffusivity,
 float3 *d_output,
 float timeStep,
 int   nx,
 int   ny,
 size_t   pitchBytes
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const char* imgP = (char*)d_input + y*pitchBytes + x*sizeof(float3);
  const char* diffP = (char*)d_diffusivity + y*pitchBytes + x*sizeof(float3);

  __shared__ float3 u[DIFF_BW+2][DIFF_BH+2];
  __shared__ float3 g[DIFF_BW+2][DIFF_BH+2];


  // load data into shared memory
  if (x < nx && y < ny) {
    u[tx][ty] = *( (float3*)imgP );
    g[tx][ty] = *( (float3*)diffP );

    if (x == 0) {
      u[0][ty] = u[tx][ty];
      g[0][ty] = g[tx][ty];
    }
    else if (threadIdx.x == 0) {
      u[0][ty] = *( ((float3*)imgP)-1 );
      g[0][ty] = *( ((float3*)diffP)-1 );
    }
    if (x == nx-1) {
      u[tx+1][ty] = u[tx][ty];
      g[tx+1][ty] = g[tx][ty];
    } 
    else if (threadIdx.x == blockDim.x-1) {
      u[tx+1][ty] = *( ((float3*)imgP)+1 );
      g[tx+1][ty] = *( ((float3*)diffP)+1 );
    }

    if (y == 0) {
      u[tx][0] = u[tx][ty];
      g[tx][0] = g[tx][ty];
    } 
    else if (threadIdx.y == 0) {
      u[tx][0] = *( (float3*)(imgP-pitchBytes) );
      g[tx][0] = *( (float3*)(diffP-pitchBytes) );
    }
    if (y == ny-1) {
      u[tx][ty+1] = u[tx][ty];
      g[tx][ty+1] = g[tx][ty];
    }
    else if (threadIdx.y == blockDim.y-1) {
      u[tx][ty+1] = *( (float3*)(imgP+pitchBytes) );
      g[tx][ty+1] = *( (float3*)(diffP+pitchBytes) );
    }
  }

  __syncthreads();
  float3 phiR,phiL,phiU, phiD;
  
    phiR.x = 0.5 * (g[tx+1][ty].x+ g[tx][ty].x);
	phiL.x = 0.5 * (g[tx-1][ty].x+ g[tx][ty].x);
	phiU.x = 0.5 * (g[tx][ty+1].x+ g[tx][ty].x);
	phiD.x = 0.5 * (g[tx][ty-1].x+ g[tx][ty].x);
	
    phiR.y= 0.5 * (g[tx+1][ty].y+ g[tx][ty].y);
	phiL.y= 0.5 * (g[tx-1][ty].y+ g[tx][ty].y);
	phiU.y= 0.5 * (g[tx][ty+1].y+ g[tx][ty].y);
	phiD.y= 0.5 * (g[tx][ty-1].y+ g[tx][ty].y);
	
    phiR.z= 0.5 * (g[tx+1][ty].z+ g[tx][ty].z);
	phiL.z= 0.5 * (g[tx-1][ty].z+ g[tx][ty].z);
	phiU.z= 0.5 * (g[tx][ty+1].z+ g[tx][ty].z);
	phiD.z= 0.5 * (g[tx][ty-1].z+ g[tx][ty].z);
	
	// ### implement me ###
	float3 res;
	if (x < nx && y < ny) {
		res.x = u[tx][ty].x + timeStep * (
						u[tx + 1][ty].x*phiR.x +
						u[tx - 1][ty].x*phiL.x + 
						u[tx][ty + 1].x*phiU.x + 
						u[tx][ty - 1].x*phiD.x - 
						u[tx][ty].x*(phiR.x+phiL.x+phiU.x+phiD.x) );
		
		res.y = u[tx][ty].y + timeStep * (
						u[tx + 1][ty].y*phiR.y +
						u[tx - 1][ty].y*phiL.y + 
						u[tx][ty + 1].y*phiU.y + 
						u[tx][ty - 1].y*phiD.y - 
						u[tx][ty].y*(phiR.y+phiL.y+phiU.y+phiD.y) );
		
		res.z = u[tx][ty].z + timeStep * (
						u[tx + 1][ty].z*phiR.z +
						u[tx - 1][ty].z*phiL.z + 
						u[tx][ty + 1].z*phiU.z + 
						u[tx][ty - 1].z*phiD.z - 
						u[tx][ty].z*(phiR.z+phiL.z+phiU.z+phiD.z) );
		
		 *((float3*)(((char*)d_output) + y*pitchBytes) + x) = res;
	}
	
	
}


// diffusivity computation for modes 1-3 gray
__global__ void compute_tv_diffusivity_shared
(
 const float *d_input,
 float *d_output,
 int   nx,
 int   ny,
 size_t   pitch
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const int idx = y*pitch + x;

  __shared__ float u[DIFF_BW+2][DIFF_BH+2];

  // load data into shared memory
  if (x < nx && y < ny) {

    u[tx][ty] = d_input[idx];

    if (x == 0)  u[0][ty] = u[tx][ty];
    else if (threadIdx.x == 0) u[0][ty] = d_input[idx-1];      
    if (x == nx-1) u[tx+1][ty] = u[tx][ty];
    else if (threadIdx.x == blockDim.x-1) u[tx+1][ty] = d_input[idx+1];

    if (y == 0)  u[tx][0] = u[tx][ty];
    else if (threadIdx.y == 0) u[tx][0] = d_input[idx-pitch];
    if (y == ny-1) u[tx][ty+1] = u[tx][ty];
    else if (threadIdx.y == blockDim.y-1) u[tx][ty+1] = d_input[idx+pitch];
  }

  __syncthreads();

 
  // make use of the constant TV_EPSILON
  float tempDerX;
  float tempDerY;
  float tmpGrad;
  if (x < nx && y < ny) {
	 tempDerX = 0.5f*(u[threadIdx.x + 2][threadIdx.y+1]-u[threadIdx.x][threadIdx.y+1]);
	 tempDerY = 0.5f*(u[threadIdx.x+1][threadIdx.y + 2] - u[threadIdx.x+1][threadIdx.y]);
	 tmpGrad = sqrt( tempDerX*tempDerX + tempDerY*tempDerY );
	 d_output[idx] = 1.0 / sqrt(tmpGrad*tmpGrad + TV_EPSILON);
  }

}


/*! Computes a joined diffusivity for an RGB Image:
 *  (g_R,g_G,g_B)(R,G,B) := 
 *  (g((R+G+B)/3),g((R+G+B)/3),g((R+G+B)/3))
 * */
__global__ void compute_tv_diffusivity_joined_shared
(
 const float3 *d_input,
 float3 *d_output,
 int   nx,
 int   ny,
 size_t   pitchBytes
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const char* imgP = (char*)d_input + y*pitchBytes + x*sizeof(float3);

  __shared__ float3 u[DIFF_BW+2][DIFF_BH+2];

  // load data into shared memory
  if (x < nx && y < ny) {

    u[tx][ty] = *( (float3*)imgP );

    if (x == 0)  u[0][ty] = u[tx][ty];
    else if (threadIdx.x == 0) u[0][ty] = *( ((float3*)imgP)-1 );
    if (x == nx-1) u[tx+1][ty] = u[tx][ty];
    else if (threadIdx.x == blockDim.x-1) u[tx+1][ty] = *( ((float3*)imgP)+1 );

    if (y == 0)  u[tx][0] = u[tx][ty];
    else if (threadIdx.y == 0) u[tx][0] = *( (float3*)(imgP-pitchBytes) );
    if (y == ny-1) u[tx][ty+1] = u[tx][ty];
    else if (threadIdx.y == blockDim.y-1) u[tx][ty+1] = *( (float3*)(imgP+pitchBytes) );
  }

  __syncthreads();
  
  
  // make use of the constant TV_EPSILON
	float3 tmpGrad;
	float3 xValue;
	float3 yValue;
	if (x < nx && y < ny) {
		float avgr = ( u[threadIdx.x + 2][threadIdx.y+1].x + u[threadIdx.x + 2][threadIdx.y+1].y + u[threadIdx.x + 2][threadIdx.y+1].z ) / 3.0;
		float avgu = ( u[threadIdx.x+1][threadIdx.y + 2].x + u[threadIdx.x+1][threadIdx.y + 2].y + u[threadIdx.x+1][threadIdx.y + 2].z ) / 3.0;
		float avgl = ( u[threadIdx.x][threadIdx.y+1].x + u[threadIdx.x][threadIdx.y+1].y + u[threadIdx.x][threadIdx.y+1].z) / 3.0;
		float avgd = ( u[threadIdx.x+1][threadIdx.y].x + u[threadIdx.x+1][threadIdx.y].y + u[threadIdx.x+1][threadIdx.y].z) / 3.0;
		// x derivatives 
		xValue.x = 0.5f * (avgr	- avgl);
		xValue.y = xValue.x;
		xValue.z = xValue.x;
		
		// y derivatives 
		yValue.x = 0.5f * (avgu - avgd);
		yValue.y = yValue.x;
		yValue.z = yValue.x;
		
		tmpGrad.x = 1.0 / sqrt(xValue.x*xValue.x + yValue.x*yValue.x + TV_EPSILON);
		tmpGrad.y = 1.0 / sqrt(xValue.x*xValue.x + yValue.x*yValue.x + TV_EPSILON);
		tmpGrad.z = 1.0 / sqrt(xValue.x*xValue.x + yValue.x*yValue.x + TV_EPSILON);

		*((float3*)(((char*)d_output) + y*pitchBytes) + x) = tmpGrad;
	
	}
}


/*! Computes a separate diffusivity for an RGB Image:
 *  (g_R,g_G,g_B)(R,G,B) := 
 *  (g(R),g(G),g(B))
 * */
__global__ void compute_tv_diffusivity_separate_shared
(
 const float3 *d_input,
 float3 *d_output,
 int   nx,
 int   ny,
 size_t   pitchBytes
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;
  const char* imgP = (char*)d_input + y*pitchBytes + x*sizeof(float3);

  __shared__ float3 u[DIFF_BW+2][DIFF_BH+2];

  // load data into shared memory
  if (x < nx && y < ny) {

    u[tx][ty] = *( (float3*)imgP );

    if (x == 0)  u[threadIdx.x][ty] = u[tx][ty];
    else if (threadIdx.x == 0) u[0][ty] = *( ((float3*)imgP)-1 );
    if (x == nx-1) u[tx+1][ty] = u[tx][ty];
    else if (threadIdx.x == blockDim.x-1) u[tx+1][ty] = *( ((float3*)imgP)+1 );

    if (y == 0)  u[tx][0] = u[tx][ty];
    else if (threadIdx.y == 0) u[tx][0] = *( (float3*)(imgP-pitchBytes) );
    if (y == ny-1) u[tx][ty+1] = u[tx][ty];
    else if (threadIdx.y == blockDim.y-1) u[tx][ty+1] = *( (float3*)(imgP+pitchBytes) );
  }

  __syncthreads();

  
  // make use of the constant TV_EPSILON
	float3 tmpGrad;
	float3 xValue;
	float3 yValue;
	if (x < nx && y < ny) {
		// x derivatives 
		xValue.x = 0.5f * (u[threadIdx.x + 2][threadIdx.y+1].x
				- u[threadIdx.x][threadIdx.y+1].x);
		xValue.y = 0.5f * (u[threadIdx.x + 2][threadIdx.y+1].y
				- u[threadIdx.x][threadIdx.y+1].y);
		xValue.z = 0.5f * (u[threadIdx.x + 2][threadIdx.y+1].z
				- u[threadIdx.x][threadIdx.y+1].z);
		// y derivatives 
		yValue.x = 0.5f * (u[threadIdx.x+1][threadIdx.y + 2].x
				- u[threadIdx.x+1][threadIdx.y].x);
		yValue.y = 0.5f * (u[threadIdx.x+1][threadIdx.y + 2].y
				- u[threadIdx.x+1][threadIdx.y].y);
		yValue.z = 0.5f * (u[threadIdx.x+1][threadIdx.y + 2].z
				- u[threadIdx.x+1][threadIdx.y].z);

		
		tmpGrad.x = 1.0 / sqrt(xValue.x*xValue.x + yValue.x*yValue.x + TV_EPSILON);
		tmpGrad.y = 1.0 / sqrt(xValue.x*xValue.x + yValue.x*yValue.x + TV_EPSILON);
		tmpGrad.z = 1.0 / sqrt(xValue.x*xValue.x + yValue.x*yValue.x + TV_EPSILON);

		*((float3*)(((char*)d_output) + y*pitchBytes) + x) = tmpGrad;	
	}

}




//----------------------------------------------------------------------------
// Non-linear Diffusion - Jacobi scheme
//----------------------------------------------------------------------------



// mode 2 gray: Jacobi solver
__global__ void jacobi_shared
(
 float *d_output,
 const float *d_input,
 const float *d_original,
 const float *d_diffusivity,
 float weight,
 int   nx,
 int   ny,
 size_t   pitch
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int idx = y*pitch + x;

  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;

  __shared__ float u[DIFF_BW+2][DIFF_BH+2];
  __shared__ float g[DIFF_BW+2][DIFF_BH+2];


  // load data into shared memory
  if (x < nx && y < ny) {
    u[tx][ty] = d_input[idx];
    g[tx][ty] = d_diffusivity[idx];

    if (x == 0)  {
      u[0][ty] = u[tx][ty];
      g[0][ty] = g[tx][ty];
    }
    else if (threadIdx.x == 0) {
      u[0][ty] = d_input[idx-1];
      g[0][ty] = d_diffusivity[idx-1];
    }
    if (x == nx-1) {
      u[tx+1][ty] = u[tx][ty];
      g[tx+1][ty] = g[tx][ty];
    }
    else if (threadIdx.x == blockDim.x-1) {
      u[tx+1][ty] = d_input[idx+1];
      g[tx+1][ty] = d_diffusivity[idx+1];
    }

    if (y == 0) {
      u[tx][0] = u[tx][ty];
      g[tx][0] = g[tx][ty];
    }
    else if (threadIdx.y == 0) {
      u[tx][0] = d_input[idx-pitch];
      g[tx][0] = d_diffusivity[idx-pitch];
    }
    if (y == ny-1) {
      u[tx][ty+1] = u[tx][ty];
      g[tx][ty+1] = g[tx][ty];
    }
    else if (threadIdx.y == blockDim.y-1) {
      u[tx][ty+1] = d_input[idx+pitch];
      g[tx][ty+1] = d_diffusivity[idx+pitch];
    }
  }

  __syncthreads();

  
  // ### implement me ###
  float phiR = 0.5 * (g[tx+1][ty] + g[tx][ty]);
  float phiL = 0.5 * (g[tx-1][ty] + g[tx][ty]);
  float phiU = 0.5 * (g[tx][ty+1] + g[tx][ty]);
  float phiD = 0.5 * (g[tx][ty-1] + g[tx][ty]);
  
  if (x == 0)
	  phiL = 0;
  else if (x == nx - 1)
	  phiR = 0;
  
  if (y == 0)
	  phiD = 0;
   else if (y == ny - 1)
	   phiU = 0;
  
  // ### implement me ###
	if (x < nx && y < ny) {
		d_output[idx] = (1 / ( 1 + weight * (phiR+phiL+phiU+phiD ) ) )* 
				( d_original[idx] + weight*( u[tx + 1][ty]*phiR + u[tx - 1][ty]*phiL + u[tx][ty + 1]*phiU + u[tx][ty - 1]*phiD ) );
	}
				
  
}



// mode 2 interleaved: Jacobi solver
__global__ void jacobi_shared
(
 float3 *d_output,
 const float3 *d_input,
 const float3 *d_original,
 const float3 *d_diffusivity,
 float weight,
 int   nx,
 int   ny,
 size_t   pitchBytes
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const char* imgP = (char*)d_input + y*pitchBytes + x*sizeof(float3);
  const char* diffP = (char*)d_diffusivity + y*pitchBytes + x*sizeof(float3);

  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;

  __shared__ float3 u[DIFF_BW+2][DIFF_BH+2];
  __shared__ float3 g[DIFF_BW+2][DIFF_BH+2];


  // load data into shared memory
  if (x < nx && y < ny) {
    u[tx][ty] = *( (float3*)imgP );
    g[tx][ty] = *( (float3*)diffP );

    if (x == 0)  {
      u[0][ty] = u[tx][ty];
      g[0][ty] = g[tx][ty];
    }
    else if (threadIdx.x == 0) {
      u[0][ty] = *( ((float3*)imgP)-1 );
      g[0][ty] = *( ((float3*)diffP)-1 );
    }
    if (x == nx-1) {
      u[tx+1][ty] = u[tx][ty];
      g[tx+1][ty] = g[tx][ty];
    }
    else if (threadIdx.x == blockDim.x-1) {
      u[tx+1][ty] = *( ((float3*)imgP)+1 );
      g[tx+1][ty] = *( ((float3*)diffP)+1 );
    }

    if (y == 0) {
      u[tx][0] = u[tx][ty];
      g[tx][0] = g[tx][ty];
    }
    else if (threadIdx.y == 0) {
      u[tx][0] = *( (float3*)(imgP-pitchBytes) );
      g[tx][0] = *( (float3*)(diffP-pitchBytes) );
    }
    if (y == ny-1) {
      u[tx][ty+1] = u[tx][ty];
      g[tx][ty+1] = g[tx][ty];
    }
    else if (threadIdx.y == blockDim.y-1) {
      u[tx][ty+1] = *( (float3*)(imgP+pitchBytes) );
      g[tx][ty+1] = *( (float3*)(diffP+pitchBytes) );
    }
  }

  __syncthreads();

  
  	float3 phiR,phiL,phiU, phiD;
  
    phiR.x = 0.5 * (g[tx+1][ty].x+ g[tx][ty].x);
	phiL.x = 0.5 * (g[tx-1][ty].x+ g[tx][ty].x);
	phiU.x = 0.5 * (g[tx][ty+1].x+ g[tx][ty].x);
	phiD.x = 0.5 * (g[tx][ty-1].x+ g[tx][ty].x);
	
    phiR.y= 0.5 * (g[tx+1][ty].y+ g[tx][ty].y);
	phiL.y= 0.5 * (g[tx-1][ty].y+ g[tx][ty].y);
	phiU.y= 0.5 * (g[tx][ty+1].y+ g[tx][ty].y);
	phiD.y= 0.5 * (g[tx][ty-1].y+ g[tx][ty].y);
	
    phiR.z= 0.5 * (g[tx+1][ty].z+ g[tx][ty].z);
	phiL.z= 0.5 * (g[tx-1][ty].z+ g[tx][ty].z);
	phiU.z= 0.5 * (g[tx][ty+1].z+ g[tx][ty].z);
	phiD.z= 0.5 * (g[tx][ty-1].z+ g[tx][ty].z);
	
	
	if (x == 0) {
		phiL.x = 0;
		phiL.y = 0;
		phiL.z = 0;
	}
	else if (x == nx - 1) {
		phiR.x = 0;
		phiR.y = 0;
		phiR.z = 0;
	}
	
	if (y == 0) {
		phiD.x = 0;
		phiD.y = 0;
		phiD.z = 0;
	}
	else if (y == ny - 1) {
		phiU.x = 0;
		phiU.y = 0;
		phiU.z = 0;
		
	}
	// ### implement me ###
	float3 res;
	const char* original = (char*)d_original + y*pitchBytes + x*sizeof(float3);
	const float3 origVal = *((float3*)original);
	
	if (x < nx && y < ny) {
		res.x = (1 / ( 1 + weight * (phiR.x+phiL.x+phiU.x+phiD.x ) ) )* 
					( origVal.x + weight*( u[tx + 1][ty].x*phiR.x + u[tx - 1][ty].x*phiL.x + u[tx][ty + 1].x*phiU.x + u[tx][ty - 1].x*phiD.x ) );
		res.y = (1 / ( 1 + weight * (phiR.y+phiL.y+phiU.y+phiD.y ) ) )* 
						( origVal.y + weight*( u[tx + 1][ty].y*phiR.y + u[tx - 1][ty].y*phiL.y + u[tx][ty + 1].y*phiU.y + u[tx][ty - 1].y*phiD.y ) );
		res.z = (1 / ( 1 + weight * (phiR.z+phiL.z+phiU.z+phiD.z ) ) )* 
						( origVal.z + weight*( u[tx + 1][ty].z*phiR.z + u[tx - 1][ty].z*phiL.z + u[tx][ty + 1].z*phiU.z + u[tx][ty - 1].z*phiD.z ) );

		
		 *((float3*)(((char*)d_output) + y*pitchBytes) + x) = res;
	}
	


}



//----------------------------------------------------------------------------
// Non-linear Diffusion - Successive Over-Relaxation (SOR)
//----------------------------------------------------------------------------



// mode 3 gray: SOR solver
__global__ void sor_shared
(
 float *d_output,
 const float *d_input,
 const float *d_original,
 const float *d_diffusivity,
 float weight,
 float overrelaxation,
 int   nx,
 int   ny,
 size_t   pitch,
 int   red
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const int idx = y*pitch + x;
  
  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;

  __shared__ float u[DIFF_BW+2][DIFF_BH+2];
  __shared__ float g[DIFF_BW+2][DIFF_BH+2];


  // load data into shared memory
  if (x < nx && y < ny) {
    u[tx][ty] = d_input[idx];
    g[tx][ty] = d_diffusivity[idx];

    if (x == 0)  {
      u[0][ty] = u[tx][ty];
      g[0][ty] = g[tx][ty];
    }
    else if (threadIdx.x == 0) {
      u[0][ty] = d_input[idx-1];
      g[0][ty] = d_diffusivity[idx-1];
    }
    if (x == nx-1) {
      u[tx+1][ty] = u[tx][ty];
      g[tx+1][ty] = g[tx][ty];
    }
    else if (threadIdx.x == blockDim.x-1) {
      u[tx+1][ty] = d_input[idx+1];
      g[tx+1][ty] = d_diffusivity[idx+1];
    }

    if (y == 0) {
      u[tx][0] = u[tx][ty];
      g[tx][0] = g[tx][ty];
    }
    else if (threadIdx.y == 0) {
      u[tx][0] = d_input[idx-pitch];
      g[tx][0] = d_diffusivity[idx-pitch];
    }
    if (y == ny-1) {
      u[tx][ty+1] = u[tx][ty];
      g[tx][ty+1] = g[tx][ty];
    }
    else if (threadIdx.y == blockDim.y-1) {
      u[tx][ty+1] = d_input[idx+pitch];
      g[tx][ty+1] = d_diffusivity[idx+pitch];
    }
  }

  __syncthreads();


  float phiR = 0.5 * (g[tx+1][ty] + g[tx][ty]);
  float phiL = 0.5 * (g[tx-1][ty] + g[tx][ty]);
  float phiU = 0.5 * (g[tx][ty+1] + g[tx][ty]);
  float phiD = 0.5 * (g[tx][ty-1] + g[tx][ty]);
  
  if (x == 0)
		phiL = 0;
	else if (x == nx - 1)
		phiR = 0;

	if (y == 0)
		phiD = 0;
	else if (y == ny - 1)
		phiU = 0;
   
	if ((x < nx && y < ny && red && (x+y)%2==0) || (x < nx && y < ny && !red && (x+y)%2!=0)) {
		d_output[idx] = overrelaxation*(1 / ( 1 + weight * (phiR+phiL+phiU+phiD ) ) )* 
				( d_original[idx] + weight*( u[tx + 1][ty]*phiR + u[tx - 1][ty]*phiL + u[tx][ty + 1]*phiU + u[tx][ty - 1]*phiD ) )
				+ (1 - overrelaxation)*u[tx][ty];
	}
		
}



// mode 3 interleaved: SOR solver
__global__ void sor_shared
(
 float3 *d_output,
 const float3 *d_input,
 const float3 *d_original,
 const float3 *d_diffusivity,
 float weight,
 float overrelaxation,
 int   nx,
 int   ny,
 size_t   pitchBytes,
 int   red
 )
{
  const int x = blockIdx.x * blockDim.x + threadIdx.x;
  const int y = blockIdx.y * blockDim.y + threadIdx.y;
  const char* imgP = (char*)d_input + y*pitchBytes + x*sizeof(float3);
  const char* diffP = (char*)d_diffusivity + y*pitchBytes + x*sizeof(float3);

  const int tx = threadIdx.x+1;
  const int ty = threadIdx.y+1;

  __shared__ float3 u[DIFF_BW+2][DIFF_BH+2];
  __shared__ float3 g[DIFF_BW+2][DIFF_BH+2];



  // load data into shared memory
  if (x < nx && y < ny) {
    u[tx][ty] = *( (float3*)imgP );
    g[tx][ty] = *( (float3*)diffP );

    if (x == 0)  {
      u[0][ty] = u[tx][ty];
      g[0][ty] = g[tx][ty];
    }
    else if (threadIdx.x == 0) {
      u[0][ty] = *( ((float3*)imgP)-1 );
      g[0][ty] = *( ((float3*)diffP)-1 );
    }
    if (x == nx-1) {
      u[tx+1][ty] = u[tx][ty];
      g[tx+1][ty] = g[tx][ty];
    }
    else if (threadIdx.x == blockDim.x-1) {
      u[tx+1][ty] = *( ((float3*)imgP)+1 );
      g[tx+1][ty] = *( ((float3*)diffP)+1 );
    }

    if (y == 0) {
      u[tx][0] = u[tx][ty];
      g[tx][0] = g[tx][ty];
    }
    else if (threadIdx.y == 0) {
      u[tx][0] = *( (float3*)(imgP-pitchBytes) );
      g[tx][0] = *( (float3*)(diffP-pitchBytes) );
    }
    if (y == ny-1) {
      u[tx][ty+1] = u[tx][ty];
      g[tx][ty+1] = g[tx][ty];
    }
    else if (threadIdx.y == blockDim.y-1) {
      u[tx][ty+1] = *( (float3*)(imgP+pitchBytes) );
      g[tx][ty+1] = *( (float3*)(diffP+pitchBytes) );
    }
  }

  __syncthreads();

  
	float3 phiR,phiL,phiU, phiD;

    phiR.x = 0.5 * (g[tx+1][ty].x+ g[tx][ty].x);
	phiL.x = 0.5 * (g[tx-1][ty].x+ g[tx][ty].x);
	phiU.x = 0.5 * (g[tx][ty+1].x+ g[tx][ty].x);
	phiD.x = 0.5 * (g[tx][ty-1].x+ g[tx][ty].x);
	
    phiR.y= 0.5 * (g[tx+1][ty].y+ g[tx][ty].y);
	phiL.y= 0.5 * (g[tx-1][ty].y+ g[tx][ty].y);
	phiU.y= 0.5 * (g[tx][ty+1].y+ g[tx][ty].y);
	phiD.y= 0.5 * (g[tx][ty-1].y+ g[tx][ty].y);
	
    phiR.z= 0.5 * (g[tx+1][ty].z+ g[tx][ty].z);
	phiL.z= 0.5 * (g[tx-1][ty].z+ g[tx][ty].z);
	phiU.z= 0.5 * (g[tx][ty+1].z+ g[tx][ty].z);
	phiD.z= 0.5 * (g[tx][ty-1].z+ g[tx][ty].z);
	
	if (x == 0) {
		phiL.x = 0;
		phiL.y = 0;
		phiL.z = 0;
	}
	else if (x == nx - 1) {
		phiR.x = 0;
		phiR.y = 0;
		phiR.z = 0;
	}
	
	if (y == 0) {
		phiD.x = 0;
		phiD.y = 0;
		phiD.z = 0;
	}
	else if (y == ny - 1) {
		phiU.x = 0;
		phiU.y = 0;
		phiU.z = 0;
		
	}
	// ### implement me ###
	float3 res;
	const char* original = (char*)d_original + y*pitchBytes + x*sizeof(float3);
	const float3 origVal = *((float3*)original);
			
	if ((x < nx && y < ny && red && (x+y)%2==0) || (x < nx && y < ny && !red && (x+y)%2!=0)) {
		res.x = overrelaxation*(1 / ( 1 + weight * (phiR.x+phiL.x+phiU.x+phiD.x ) ) )* 
					( origVal.x + weight*( u[tx + 1][ty].x*phiR.x + u[tx - 1][ty].x*phiL.x + u[tx][ty + 1].x*phiU.x + u[tx][ty - 1].x*phiD.x ) ) + (1 - overrelaxation)*u[tx][ty].x;
		res.y = (1 / ( 1 + weight * (phiR.y+phiL.y+phiU.y+phiD.y ) ) )* 
						( origVal.y + weight*( u[tx + 1][ty].y*phiR.y + u[tx - 1][ty].y*phiL.y + u[tx][ty + 1].y*phiU.y + u[tx][ty - 1].y*phiD.y ) ) + (1 - overrelaxation)*u[tx][ty].y;
		res.z = (1 / ( 1 + weight * (phiR.z+phiL.z+phiU.z+phiD.z ) ) )* 
						( origVal.z + weight*( u[tx + 1][ty].z*phiR.z + u[tx - 1][ty].z*phiL.z + u[tx][ty + 1].z*phiU.z + u[tx][ty - 1].z*phiD.z ) ) + (1 - overrelaxation)*u[tx][ty].z;

		
		 *((float3*)(((char*)d_output) + y*pitchBytes) + x) = res;
	}
	

}




//----------------------------------------------------------------------------
// Host function
//----------------------------------------------------------------------------



void gpu_diffusion
(
 const float *input,
 float *output,
 int nx, int ny, int nc, 
 float timeStep,
 int iterations,
 float weight,
 int lagged_iterations,
 float overrelaxation,
 int mode,
 bool jointDiffusivity
 )
{
  int i,j;
  size_t pitchF1, pitchBytesF1, pitchBytesF3;
  float *d_input = 0;
  float *d_output = 0;
  float *d_diffusivity = 0;
  float *d_original = 0;
  float *temp = 0;

  dim3 dimGrid((int)ceil((float)nx/DIFF_BW), (int)ceil((float)ny/DIFF_BH));
  dim3 dimBlock(DIFF_BW,DIFF_BH);

  // Allocation of GPU Memory
  if (nc == 1) {

    cutilSafeCall( cudaMallocPitch( (void**)&(d_input), &pitchBytesF1, nx*sizeof(float), ny ) );
    cutilSafeCall( cudaMallocPitch( (void**)&(d_output), &pitchBytesF1, nx*sizeof(float), ny ) );
    if (mode) cutilSafeCall( cudaMallocPitch( (void**)&(d_diffusivity), &pitchBytesF1, nx*sizeof(float), ny ) );
    if (mode >= 2) cutilSafeCall( cudaMallocPitch( (void**)&(d_original), &pitchBytesF1, nx*sizeof(float), ny ) );

    cutilSafeCall( cudaMemcpy2D(d_input, pitchBytesF1, input, nx*sizeof(float), nx*sizeof(float), ny, cudaMemcpyHostToDevice) );
    if (mode >= 2) cutilSafeCall( cudaMemcpy2D(d_original, pitchBytesF1, d_input, pitchBytesF1, nx*sizeof(float), ny, cudaMemcpyDeviceToDevice) );

    pitchF1 = pitchBytesF1/sizeof(float);

  } else if (nc == 3) {

    cutilSafeCall( cudaMallocPitch( (void**)&(d_input), &pitchBytesF3, nx*sizeof(float3), ny ) );
    cutilSafeCall( cudaMallocPitch( (void**)&(d_output), &pitchBytesF3, nx*sizeof(float3), ny ) );
    if (mode) cutilSafeCall( cudaMallocPitch( (void**)&(d_diffusivity), &pitchBytesF3, nx*sizeof(float3), ny ) );
    if (mode >= 2) cutilSafeCall( cudaMallocPitch( (void**)&(d_original), &pitchBytesF3, nx*sizeof(float3), ny ) );

    cutilSafeCall( cudaMemcpy2D(d_input, pitchBytesF3, input, nx*sizeof(float3), nx*sizeof(float3), ny, cudaMemcpyHostToDevice) );
    if (mode >= 2) cutilSafeCall( cudaMemcpy2D(d_original, pitchBytesF3, d_input, pitchBytesF3, nx*sizeof(float3), ny, cudaMemcpyDeviceToDevice) );

  }


  // Execution of the Diffusion Kernel

  if (mode == 0) {   // linear isotropic diffision
    if (nc == 1) {
      for (i=0;i<iterations;i++) {
        diffuse_linear_isotrop_shared<<<dimGrid,dimBlock>>>(d_input, d_output, timeStep, nx, ny, pitchF1);

        cutilSafeCall( cudaThreadSynchronize() );

        temp = d_input;
        d_input = d_output;
        d_output = temp;
      }
    }
    else if (nc == 3) {
      for (i=0;i<iterations;i++) {
        diffuse_linear_isotrop_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_output,timeStep,nx,ny,pitchBytesF3);

        cutilSafeCall( cudaThreadSynchronize() );

        temp = d_input;
        d_input = d_output;
        d_output = temp;
      }
    }
  }
  else if (mode == 1) {  // nonlinear isotropic diffusion
    if (nc == 1) {

      for (i=0;i<iterations;i++) {
        compute_tv_diffusivity_shared<<<dimGrid,dimBlock>>>(d_input,d_diffusivity,nx,ny,pitchF1);

        cutilSafeCall( cudaThreadSynchronize() );

        diffuse_nonlinear_isotrop_shared<<<dimGrid,dimBlock>>>(d_input,d_diffusivity,d_output,timeStep,nx,ny,pitchF1);

        cutilSafeCall( cudaThreadSynchronize() );

        temp = d_input;
        d_input = d_output;
        d_output = temp;
      }
    }
    else if (nc == 3) {
      for (i=0;i<iterations;i++) {
        if (jointDiffusivity)
          compute_tv_diffusivity_joined_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_diffusivity,nx,ny,pitchBytesF3);
        else
          compute_tv_diffusivity_separate_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_diffusivity,nx,ny,pitchBytesF3);


        cutilSafeCall( cudaThreadSynchronize() );

        diffuse_nonlinear_isotrop_shared<<<dimGrid,dimBlock>>>
          ((float3*)d_input,(float3*)d_diffusivity,(float3*)d_output,timeStep,nx,ny,pitchBytesF3);

        cutilSafeCall( cudaThreadSynchronize() );

        temp = d_input;
        d_input = d_output;
        d_output = temp;
      }
    }
  }
  else if (mode == 2) {    // Jacobi-method
    if (nc == 1) {
      for (i=0;i<iterations;i++) {
        compute_tv_diffusivity_shared<<<dimGrid,dimBlock>>>(d_input,d_diffusivity,nx,ny,pitchF1);

        cutilSafeCall( cudaThreadSynchronize() );

        for (j=0;j<lagged_iterations;j++) {
          jacobi_shared<<<dimGrid,dimBlock>>> (d_output,d_input,d_original,
            d_diffusivity,weight,nx,ny,pitchF1);

          cutilSafeCall( cudaThreadSynchronize() );

          temp = d_input;
          d_input = d_output;
          d_output = temp;
        }
      }
    }
    else if (nc == 3) {
      for (i=0;i<iterations;i++) {
        if (jointDiffusivity)
          compute_tv_diffusivity_joined_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_diffusivity,nx,ny,pitchBytesF3);
        else
          compute_tv_diffusivity_separate_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_diffusivity,nx,ny,pitchBytesF3);

        cutilSafeCall( cudaThreadSynchronize() );

        for (j=0;j<lagged_iterations;j++) {
          jacobi_shared<<<dimGrid,dimBlock>>>
            ((float3*)d_output,(float3*)d_input,
            (float3*)d_original,(float3*)d_diffusivity,
            weight,nx,ny,pitchBytesF3);

          cutilSafeCall( cudaThreadSynchronize() );

          temp = d_input;
          d_input = d_output;
          d_output = temp;
        }
      }
    }    
  }
  else if (mode == 3) {    // Successive Over Relaxation (Gauss-Seidel with extrapolation)
    if (nc == 1) {
      for (i=0;i<iterations;i++) {
        compute_tv_diffusivity_shared<<<dimGrid,dimBlock>>>(d_input,d_diffusivity,nx,ny,pitchF1);

        cutilSafeCall( cudaThreadSynchronize() );

        for(j=0;j<lagged_iterations;j++) {					
          sor_shared<<<dimGrid,dimBlock>>>(d_input,d_input,d_original,
            d_diffusivity,weight,overrelaxation,nx,ny,pitchF1, 0);

          cutilSafeCall( cudaThreadSynchronize() );

          sor_shared<<<dimGrid,dimBlock>>>(d_input,d_input,d_original,
            d_diffusivity,weight,overrelaxation,nx,ny,pitchF1, 1);

          cutilSafeCall( cudaThreadSynchronize() );
        }
      }
    }
    if (nc == 3) {
      for (i=0;i<iterations;i++) {
        if (jointDiffusivity)
          compute_tv_diffusivity_joined_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_diffusivity,nx,ny,pitchBytesF3);
        else
          compute_tv_diffusivity_separate_shared<<<dimGrid,dimBlock>>>((float3*)d_input,(float3*)d_diffusivity,nx,ny,pitchBytesF3);

        cutilSafeCall( cudaThreadSynchronize() );

        for (j=0;j<lagged_iterations;j++) {
          sor_shared<<<dimGrid,dimBlock>>>
            ((float3*)d_input,(float3*)d_input,
            (float3*)d_original,(float3*)d_diffusivity,
            weight,overrelaxation,nx,ny,pitchBytesF3, 0);

          cutilSafeCall( cudaThreadSynchronize() );

          sor_shared<<<dimGrid,dimBlock>>>
            ((float3*)d_input,(float3*)d_input,
            (float3*)d_original,(float3*)d_diffusivity,
            weight,overrelaxation,nx,ny,pitchBytesF3, 1);

          cutilSafeCall( cudaThreadSynchronize() );
        }
      }
    }
  }


  if (nc == 1) {
    if (mode == 3) cutilSafeCall( cudaMemcpy2D(output, nx*sizeof(float), d_input, pitchBytesF1, nx*sizeof(float), ny, cudaMemcpyDeviceToHost) );
    else cutilSafeCall( cudaMemcpy2D(output, nx*sizeof(float), d_output, pitchBytesF1, nx*sizeof(float), ny, cudaMemcpyDeviceToHost) );
  } else if (nc == 3) {
    if (mode == 3) cutilSafeCall( cudaMemcpy2D(output, nx*sizeof(float3), d_input, pitchBytesF3, nx*sizeof(float3), ny, cudaMemcpyDeviceToHost) );
    else cutilSafeCall( cudaMemcpy2D(output, nx*sizeof(float3), d_output, pitchBytesF3, nx*sizeof(float3), ny, cudaMemcpyDeviceToHost) );
  }


  // clean up
  if (d_original) cutilSafeCall( cudaFree(d_original) );
  if (d_diffusivity) cutilSafeCall( cudaFree(d_diffusivity) );
  if (d_output) cutilSafeCall( cudaFree(d_output) );
  if (d_input)  cutilSafeCall( cudaFree(d_input) );
}
	.file	"tmpxft_00255177_00000000-6_diffusion.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2036:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2036:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._Z14__cudaSafeCall9cudaErrorPKci.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%s(%i) : cutilSafeCall() Runtime API error : %s.\n"
	.section	.text._Z14__cudaSafeCall9cudaErrorPKci,"axG",@progbits,_Z14__cudaSafeCall9cudaErrorPKci,comdat
	.weak	_Z14__cudaSafeCall9cudaErrorPKci
	.type	_Z14__cudaSafeCall9cudaErrorPKci, @function
_Z14__cudaSafeCall9cudaErrorPKci:
.LFB2027:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L2
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	movl	%edx, 12(%rsp)
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movq	%rbx, %rcx
	movl	12(%rsp), %r8d
	movq	%rax, %r9
	leaq	.LC0(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L2:
	.cfi_def_cfa_offset 8
	.cfi_restore 3
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z14__cudaSafeCall9cudaErrorPKci, .-_Z14__cudaSafeCall9cudaErrorPKci
	.text
	.globl	_Z15getStudentLoginv
	.type	_Z15getStudentLoginv, @function
_Z15getStudentLoginv:
.LFB2028:
	.cfi_startproc
	endbr64
	movq	studentLogin(%rip), %rax
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z15getStudentLoginv, .-_Z15getStudentLoginv
	.globl	_Z14getStudentNamev
	.type	_Z14getStudentNamev, @function
_Z14getStudentNamev:
.LFB2029:
	.cfi_startproc
	endbr64
	movq	studentName(%rip), %rax
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z14getStudentNamev, .-_Z14getStudentNamev
	.globl	_Z12getStudentIDv
	.type	_Z12getStudentIDv, @function
_Z12getStudentIDv:
.LFB2030:
	.cfi_startproc
	endbr64
	movl	$993889, %eax
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z12getStudentIDv, .-_Z12getStudentIDv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"p010"
.LC2:
	.string	"John Doe"
	.text
	.globl	_Z16checkStudentDatav
	.type	_Z16checkStudentDatav, @function
_Z16checkStudentDatav:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%rcx
	.cfi_def_cfa_offset 16
	movq	studentLogin(%rip), %rdi
	leaq	.LC1(%rip), %rsi
	call	strcmp@PLT
	movl	%eax, %edx
	xorl	%eax, %eax
	testl	%edx, %edx
	je	.L11
	movq	studentName(%rip), %rdi
	leaq	.LC2(%rip), %rsi
	call	strcmp@PLT
	testl	%eax, %eax
	setne	%al
.L11:
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z16checkStudentDatav, .-_Z16checkStudentDatav
	.globl	_Z21checkStudentNameAndIDv
	.type	_Z21checkStudentNameAndIDv, @function
_Z21checkStudentNameAndIDv:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	studentName(%rip), %rdi
	leaq	.LC2(%rip), %rsi
	call	strcmp@PLT
	popq	%rdx
	.cfi_def_cfa_offset 8
	testl	%eax, %eax
	setne	%al
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z21checkStudentNameAndIDv, .-_Z21checkStudentNameAndIDv
	.globl	_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim
	.type	_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim, @function
_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim:
.LFB2058:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movl	%edx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movss	%xmm0, 28(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z29diffuse_linear_isotrop_sharedPKfPffiim(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L19:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2058:
	.size	_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim, .-_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim
	.globl	_Z29diffuse_linear_isotrop_sharedPKfPffiim
	.type	_Z29diffuse_linear_isotrop_sharedPKfPffiim, @function
_Z29diffuse_linear_isotrop_sharedPKfPffiim:
.LFB2059:
	.cfi_startproc
	endbr64
	jmp	_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim
	.cfi_endproc
.LFE2059:
	.size	_Z29diffuse_linear_isotrop_sharedPKfPffiim, .-_Z29diffuse_linear_isotrop_sharedPKfPffiim
	.globl	_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim
	.type	_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim, @function
_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim:
.LFB2060:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movl	%edx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movss	%xmm0, 28(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L24:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2060:
	.size	_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim, .-_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim
	.globl	_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim
	.type	_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim, @function
_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim:
.LFB2061:
	.cfi_startproc
	endbr64
	jmp	_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim
	.cfi_endproc
.LFE2061:
	.size	_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim, .-_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim
	.globl	_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim
	.type	_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim, @function
_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim:
.LFB2062:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movq	%r9, (%rsp)
	movss	%xmm0, 20(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	movq	%rsp, %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L29:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2062:
	.size	_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim, .-_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim
	.globl	_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim
	.type	_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim, @function
_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim:
.LFB2063:
	.cfi_startproc
	endbr64
	jmp	_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim
	.cfi_endproc
.LFE2063:
	.size	_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim, .-_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim
	.globl	_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim
	.type	_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim, @function
_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim:
.LFB2064:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movq	%r9, (%rsp)
	movss	%xmm0, 20(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	movq	%rsp, %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L34:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2064:
	.size	_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim, .-_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim
	.globl	_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim
	.type	_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim, @function
_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim:
.LFB2065:
	.cfi_startproc
	endbr64
	jmp	_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim
	.cfi_endproc
.LFE2065:
	.size	_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim, .-_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim
	.globl	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim
	.type	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim, @function
_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim:
.LFB2066:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z29compute_tv_diffusivity_sharedPKfPfiim(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L39:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L41
	call	__stack_chk_fail@PLT
.L41:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2066:
	.size	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim, .-_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim
	.globl	_Z29compute_tv_diffusivity_sharedPKfPfiim
	.type	_Z29compute_tv_diffusivity_sharedPKfPfiim, @function
_Z29compute_tv_diffusivity_sharedPKfPfiim:
.LFB2067:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim
	.cfi_endproc
.LFE2067:
	.size	_Z29compute_tv_diffusivity_sharedPKfPfiim, .-_Z29compute_tv_diffusivity_sharedPKfPfiim
	.globl	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim
	.type	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim, @function
_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim:
.LFB2068:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L44
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L44:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L46
	call	__stack_chk_fail@PLT
.L46:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2068:
	.size	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim, .-_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim
	.globl	_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim
	.type	_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim, @function
_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim:
.LFB2069:
	.cfi_startproc
	endbr64
	jmp	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim
	.cfi_endproc
.LFE2069:
	.size	_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim, .-_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim
	.globl	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim
	.type	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim, @function
_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim:
.LFB2070:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L49
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L49:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L51
	call	__stack_chk_fail@PLT
.L51:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2070:
	.size	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim, .-_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim
	.globl	_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim
	.type	_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim, @function
_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim:
.LFB2071:
	.cfi_startproc
	endbr64
	jmp	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim
	.cfi_endproc
.LFE2071:
	.size	_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim, .-_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim
	.globl	_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim
	.type	_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim, @function
_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	208(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	pushq	64(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z13jacobi_sharedPfPKfS1_S1_fiim(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 224
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L54:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
	call	__stack_chk_fail@PLT
.L56:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim, .-_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim
	.globl	_Z13jacobi_sharedPfPKfS1_S1_fiim
	.type	_Z13jacobi_sharedPfPKfS1_S1_fiim, @function
_Z13jacobi_sharedPfPKfS1_S1_fiim:
.LFB2073:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim
	.cfi_endproc
.LFE2073:
	.size	_Z13jacobi_sharedPfPKfS1_S1_fiim, .-_Z13jacobi_sharedPfPKfS1_S1_fiim
	.globl	_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim
	.type	_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim, @function
_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim:
.LFB2074:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	208(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L59
	pushq	64(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z13jacobi_sharedP6float3PKS_S2_S2_fiim(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 224
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L59:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L61
	call	__stack_chk_fail@PLT
.L61:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2074:
	.size	_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim, .-_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim
	.globl	_Z13jacobi_sharedP6float3PKS_S2_S2_fiim
	.type	_Z13jacobi_sharedP6float3PKS_S2_S2_fiim, @function
_Z13jacobi_sharedP6float3PKS_S2_S2_fiim:
.LFB2075:
	.cfi_startproc
	endbr64
	jmp	_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim
	.cfi_endproc
.LFE2075:
	.size	_Z13jacobi_sharedP6float3PKS_S2_S2_fiim, .-_Z13jacobi_sharedP6float3PKS_S2_S2_fiim
	.globl	_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi
	.type	_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi, @function
_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi:
.LFB2076:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	movss	%xmm0, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 168(%rsp)
	movq	%rsp, %rax
	movq	%rax, 176(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L64
	pushq	64(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z10sor_sharedPfPKfS1_S1_ffiimi(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 240
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L64:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L66
	call	__stack_chk_fail@PLT
.L66:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2076:
	.size	_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi, .-_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi
	.globl	_Z10sor_sharedPfPKfS1_S1_ffiimi
	.type	_Z10sor_sharedPfPKfS1_S1_ffiimi, @function
_Z10sor_sharedPfPKfS1_S1_ffiimi:
.LFB2077:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi
	.cfi_endproc
.LFE2077:
	.size	_Z10sor_sharedPfPKfS1_S1_ffiimi, .-_Z10sor_sharedPfPKfS1_S1_ffiimi
	.globl	_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi
	.type	_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi, @function
_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi:
.LFB2078:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	movss	%xmm0, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 168(%rsp)
	movq	%rsp, %rax
	movq	%rax, 176(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L69
	pushq	64(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z10sor_sharedP6float3PKS_S2_S2_ffiimi(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 240
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L69:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L71
	call	__stack_chk_fail@PLT
.L71:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2078:
	.size	_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi, .-_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi
	.globl	_Z10sor_sharedP6float3PKS_S2_S2_ffiimi
	.type	_Z10sor_sharedP6float3PKS_S2_S2_ffiimi, @function
_Z10sor_sharedP6float3PKS_S2_S2_ffiimi:
.LFB2079:
	.cfi_startproc
	endbr64
	jmp	_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi
	.cfi_endproc
.LFE2079:
	.size	_Z10sor_sharedP6float3PKS_S2_S2_ffiimi, .-_Z10sor_sharedP6float3PKS_S2_S2_ffiimi
	.section	.rodata.str1.1
.LC4:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/arashb/gpulab/master/Diffusion/src/diffusion.cu"
	.text
	.globl	_Z13gpu_diffusionPKfPfiiifififib
	.type	_Z13gpu_diffusionPKfPfiiifififib, @function
_Z13gpu_diffusionPKfPfiiifififib:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%r9d, %r14d
	xorl	%r9d, %r9d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%r8d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	subq	$152, %rsp
	.cfi_def_cfa_offset 208
	movss	%xmm0, 36(%rsp)
	movb	224(%rsp), %al
	cvtsi2ssl	%ecx, %xmm0
	mulss	.LC3(%rip), %xmm0
	movq	%rdi, 56(%rsp)
	movl	216(%rsp), %ebp
	movq	%rsi, 24(%rsp)
	movl	%edx, 12(%rsp)
	movb	%al, 55(%rsp)
	movss	%xmm1, 32(%rsp)
	movss	%xmm2, 48(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	movq	%r9, 80(%rsp)
	movq	%r9, 88(%rsp)
	movq	%r9, 96(%rsp)
	movq	%r9, 104(%rsp)
	call	ceilf@PLT
	movd	%xmm0, %r15d
	cvtsi2ssl	12(%rsp), %xmm0
	mulss	.LC3(%rip), %xmm0
	call	ceilf@PLT
	movd	%r15d, %xmm3
	movl	$1, 120(%rsp)
	movl	$1, 132(%rsp)
	cvttss2sil	%xmm0, %eax
	movl	%eax, 112(%rsp)
	cvttss2sil	%xmm3, %eax
	movl	%eax, 116(%rsp)
	movabsq	$68719476752, %rax
	movq	%rax, 124(%rsp)
	cmpl	$1, %r12d
	jne	.L75
	movslq	12(%rsp), %rax
	leaq	64(%rsp), %r15
	movslq	%ebx, %rcx
	leaq	80(%rsp), %rdi
	movq	%r15, %rsi
	leaq	.LC4(%rip), %r13
	salq	$2, %rax
	movq	%rax, %rdx
	movq	%rax, 16(%rsp)
	call	cudaMallocPitch@PLT
	movl	$1049, %edx
	movq	%r13, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	16(%rsp), %rdx
	movq	%r15, %rsi
	movslq	%ebx, %rcx
	leaq	88(%rsp), %rdi
	call	cudaMallocPitch@PLT
	movl	$1050, %edx
	movq	%r13, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	testl	%ebp, %ebp
	je	.L77
	movq	16(%rsp), %rdx
	leaq	96(%rsp), %rdi
	movq	%r15, %rsi
	movslq	%ebx, %rcx
	call	cudaMallocPitch@PLT
	movl	$1051, %edx
	movq	%r13, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	cmpl	$1, %ebp
	jle	.L77
	movq	16(%rsp), %rdx
	leaq	104(%rsp), %rdi
	movq	%r15, %rsi
	movslq	%ebx, %rcx
	call	cudaMallocPitch@PLT
	movl	$1052, %edx
	movq	%r13, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L77:
	pushq	%rcx
	.cfi_def_cfa_offset 216
	movslq	%ebx, %r9
	leaq	.LC4(%rip), %r13
	pushq	$1
	.cfi_def_cfa_offset 224
	movq	32(%rsp), %r8
	movq	72(%rsp), %rdx
	movq	96(%rsp), %rdi
	movq	80(%rsp), %rsi
	movq	%r8, %rcx
	call	cudaMemcpy2D@PLT
	popq	%rsi
	.cfi_def_cfa_offset 216
	movl	$1054, %edx
	popq	%r8
	.cfi_def_cfa_offset 208
	movl	%eax, %edi
	movq	%r13, %rsi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	cmpl	$1, %ebp
	jle	.L79
	movq	64(%rsp), %rsi
	pushq	%rax
	.cfi_def_cfa_offset 216
	movslq	%ebx, %r9
	pushq	$3
	.cfi_def_cfa_offset 224
	movq	120(%rsp), %rdi
	movq	32(%rsp), %r8
	movq	96(%rsp), %rdx
	movq	%rsi, %rcx
	call	cudaMemcpy2D@PLT
	movq	%r13, %rsi
	movl	%eax, %edi
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
	movl	$1055, %edx
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L79:
	movq	64(%rsp), %r13
	shrq	$2, %r13
	jmp	.L80
.L75:
	cmpl	$3, %r12d
	jne	.L81
	movslq	12(%rsp), %rax
	movslq	%ebx, %rcx
	leaq	80(%rsp), %rdi
	leaq	.LC4(%rip), %r15
	imulq	$12, %rax, %rax
	movq	%rax, 16(%rsp)
	leaq	72(%rsp), %rax
	movq	16(%rsp), %rdx
	movq	%rax, %rsi
	movq	%rax, 40(%rsp)
	call	cudaMallocPitch@PLT
	movl	$1061, %edx
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	16(%rsp), %rdx
	movq	40(%rsp), %rsi
	movslq	%ebx, %rcx
	leaq	88(%rsp), %rdi
	call	cudaMallocPitch@PLT
	movl	$1062, %edx
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	testl	%ebp, %ebp
	je	.L83
	movq	16(%rsp), %rdx
	movq	40(%rsp), %rsi
	leaq	96(%rsp), %rdi
	movslq	%ebx, %rcx
	call	cudaMallocPitch@PLT
	movl	$1063, %edx
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	cmpl	$1, %ebp
	jle	.L83
	movq	16(%rsp), %rdx
	movq	40(%rsp), %rsi
	leaq	104(%rsp), %rdi
	movslq	%ebx, %rcx
	call	cudaMallocPitch@PLT
	movl	$1064, %edx
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L83:
	pushq	%r11
	.cfi_def_cfa_offset 216
	movslq	%ebx, %r9
	pushq	$1
	.cfi_def_cfa_offset 224
	movq	32(%rsp), %r8
	movq	72(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	96(%rsp), %rdi
	movq	%r8, %rcx
	call	cudaMemcpy2D@PLT
	popq	%r15
	.cfi_def_cfa_offset 216
	leaq	.LC4(%rip), %r15
	movl	$1066, %edx
	movl	%eax, %edi
	movq	%r15, %rsi
	popq	%rax
	.cfi_def_cfa_offset 208
	call	_Z14__cudaSafeCall9cudaErrorPKci
	cmpl	$1, %ebp
	jle	.L80
	movq	72(%rsp), %rsi
	pushq	%r8
	.cfi_def_cfa_offset 216
	movslq	%ebx, %r9
	pushq	$3
	.cfi_def_cfa_offset 224
	movq	96(%rsp), %rdx
	movq	120(%rsp), %rdi
	movq	32(%rsp), %r8
	movq	%rsi, %rcx
	call	cudaMemcpy2D@PLT
	popq	%r9
	.cfi_def_cfa_offset 216
	movl	$1067, %edx
	popq	%r10
	.cfi_def_cfa_offset 208
	movl	%eax, %edi
	movq	%r15, %rsi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	jmp	.L95
.L80:
	testl	%ebp, %ebp
	jne	.L86
	decl	%r12d
	leaq	.LC4(%rip), %r12
	je	.L87
	xorl	%r13d, %r13d
	jmp	.L88
.L90:
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L89
	movl	12(%rsp), %edx
	movss	36(%rsp), %xmm0
	movq	%r13, %r8
	movl	%ebx, %ecx
	movq	88(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z56__device_stub__Z29diffuse_linear_isotrop_sharedPKfPffiimPKfPffiim
.L89:
	call	cudaThreadSynchronize@PLT
	movl	$1079, %edx
	movq	%r12, %rsi
	incl	%ebp
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	80(%rsp), %rax
	movq	88(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	%rdx, 80(%rsp)
.L87:
	cmpl	%r14d, %ebp
	jl	.L90
	jmp	.L150
.L93:
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L92
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movss	36(%rsp), %xmm0
	movq	88(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z63__device_stub__Z29diffuse_linear_isotrop_sharedPK6float3PS_fiimPK6float3PS_fiim
.L92:
	call	cudaThreadSynchronize@PLT
	movl	$1090, %edx
	movq	%r12, %rsi
	incl	%r13d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	80(%rsp), %rax
	movq	88(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	%rdx, 80(%rsp)
.L88:
	cmpl	%r14d, %r13d
	jl	.L93
	jmp	.L153
.L86:
	cmpl	$1, %ebp
	jne	.L95
	xorl	%r15d, %r15d
	decl	%r12d
	leaq	.LC4(%rip), %r12
	je	.L96
	xorl	%r13d, %r13d
	jmp	.L97
.L100:
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L98
	movl	12(%rsp), %edx
	movq	96(%rsp), %rsi
	movq	%r13, %r8
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim
.L98:
	call	cudaThreadSynchronize@PLT
	movl	$1104, %edx
	movq	%r12, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	124(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L99
	movl	12(%rsp), %ecx
	movss	36(%rsp), %xmm0
	movq	%r13, %r9
	movl	%ebx, %r8d
	movq	88(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z62__device_stub__Z32diffuse_nonlinear_isotrop_sharedPKfS0_PffiimPKfS0_Pffiim
.L99:
	call	cudaThreadSynchronize@PLT
	movl	$1108, %edx
	movq	%r12, %rsi
	incl	%r15d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	80(%rsp), %rax
	movq	88(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	%rdx, 80(%rsp)
.L96:
	cmpl	%r14d, %r15d
	jl	.L100
	jmp	.L150
.L106:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	cmpb	$0, 55(%rsp)
	movq	124(%rsp), %rdx
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	je	.L101
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L103
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim
.L103:
	call	cudaThreadSynchronize@PLT
	movl	$1123, %edx
	movq	%r12, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	124(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L105
	movq	72(%rsp), %r9
	movl	12(%rsp), %ecx
	movl	%ebx, %r8d
	movss	36(%rsp), %xmm0
	movq	88(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z69__device_stub__Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiimPK6float3S1_PS_fiim
.L105:
	call	cudaThreadSynchronize@PLT
	movl	$1128, %edx
	movq	%r12, %rsi
	incl	%r13d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	80(%rsp), %rax
	movq	88(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	%rdx, 80(%rsp)
.L97:
	cmpl	%r14d, %r13d
	jl	.L106
	jmp	.L153
.L101:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L103
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim
	jmp	.L103
.L81:
	testl	%ebp, %ebp
	je	.L144
	cmpl	$1, %ebp
	je	.L144
.L95:
	cmpl	$2, %ebp
	jne	.L107
	cmpl	$1, %r12d
	jne	.L187
	xorl	%ebp, %ebp
	leaq	.LC4(%rip), %r15
.L108:
	cmpl	%r14d, %ebp
	jge	.L150
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L110
	movl	12(%rsp), %edx
	movq	96(%rsp), %rsi
	movq	%r13, %r8
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim
.L110:
	call	cudaThreadSynchronize@PLT
	movl	$1141, %edx
	movq	%r15, %rsi
	xorl	%r12d, %r12d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L111:
	cmpl	208(%rsp), %r12d
	jge	.L195
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L112
	pushq	%rcx
	.cfi_def_cfa_offset 216
	movl	%ebx, %r9d
	pushq	%r13
	.cfi_def_cfa_offset 224
	movl	28(%rsp), %r8d
	movss	48(%rsp), %xmm0
	movq	112(%rsp), %rcx
	movq	120(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	104(%rsp), %rdi
	call	_Z46__device_stub__Z13jacobi_sharedPfPKfS1_S1_fiimPfPKfS1_S1_fiim
	popq	%rsi
	.cfi_def_cfa_offset 216
	popq	%rdi
	.cfi_def_cfa_offset 208
.L112:
	call	cudaThreadSynchronize@PLT
	movl	$1147, %edx
	movq	%r15, %rsi
	incl	%r12d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	80(%rsp), %rax
	movq	88(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	%rdx, 80(%rsp)
	jmp	.L111
.L195:
	incl	%ebp
	jmp	.L108
.L187:
	cmpl	$3, %r12d
	jne	.L144
	xorl	%r12d, %r12d
	leaq	.LC4(%rip), %r15
.L115:
	cmpl	%r14d, %r12d
	jge	.L153
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	cmpb	$0, 55(%rsp)
	movq	124(%rsp), %rdx
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	je	.L117
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L119
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim
	jmp	.L119
.L117:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L119
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim
.L119:
	call	cudaThreadSynchronize@PLT
	movl	$1162, %edx
	movq	%r15, %rsi
	xorl	%r13d, %r13d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L121:
	cmpl	208(%rsp), %r13d
	jge	.L196
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L122
	pushq	%r11
	.cfi_def_cfa_offset 216
	movl	%ebx, %r9d
	pushq	80(%rsp)
	.cfi_def_cfa_offset 224
	movl	28(%rsp), %r8d
	movss	48(%rsp), %xmm0
	movq	112(%rsp), %rcx
	movq	120(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	104(%rsp), %rdi
	call	_Z53__device_stub__Z13jacobi_sharedP6float3PKS_S2_S2_fiimP6float3PKS_S2_S2_fiim
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L122:
	call	cudaThreadSynchronize@PLT
	movl	$1170, %edx
	movq	%r15, %rsi
	incl	%r13d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	80(%rsp), %rax
	movq	88(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	%rdx, 80(%rsp)
	jmp	.L121
.L196:
	incl	%r12d
	jmp	.L115
.L107:
	cmpl	$3, %ebp
	jne	.L125
	cmpl	$1, %r12d
	je	.L158
	xorl	%r13d, %r13d
	leaq	.LC4(%rip), %r15
	cmpl	$3, %r12d
	je	.L127
	jmp	.L144
.L158:
	xorl	%ebp, %ebp
	leaq	.LC4(%rip), %r12
.L126:
	cmpl	%r14d, %ebp
	jge	.L197
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L128
	movl	12(%rsp), %edx
	movq	96(%rsp), %rsi
	movq	%r13, %r8
	movl	%ebx, %ecx
	movq	80(%rsp), %rdi
	call	_Z55__device_stub__Z29compute_tv_diffusivity_sharedPKfPfiimPKfPfiim
.L128:
	call	cudaThreadSynchronize@PLT
	movl	$1184, %edx
	movq	%r12, %rsi
	xorl	%r15d, %r15d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L129:
	cmpl	208(%rsp), %r15d
	jge	.L198
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L130
	movq	80(%rsp), %rdi
	pushq	$0
	.cfi_def_cfa_offset 216
	movl	%ebx, %r9d
	pushq	%r13
	.cfi_def_cfa_offset 224
	movl	28(%rsp), %r8d
	movss	64(%rsp), %xmm1
	movss	48(%rsp), %xmm0
	movq	%rdi, %rsi
	movq	112(%rsp), %rcx
	movq	120(%rsp), %rdx
	call	_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi
	popq	%r9
	.cfi_def_cfa_offset 216
	popq	%r10
	.cfi_def_cfa_offset 208
.L130:
	call	cudaThreadSynchronize@PLT
	movl	$1190, %edx
	movq	%r12, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	124(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L131
	movq	80(%rsp), %rdi
	pushq	$1
	.cfi_def_cfa_offset 216
	movl	%ebx, %r9d
	pushq	%r13
	.cfi_def_cfa_offset 224
	movl	28(%rsp), %r8d
	movss	64(%rsp), %xmm1
	movss	48(%rsp), %xmm0
	movq	%rdi, %rsi
	movq	112(%rsp), %rcx
	movq	120(%rsp), %rdx
	call	_Z45__device_stub__Z10sor_sharedPfPKfS1_S1_ffiimiPfPKfS1_S1_ffiimi
	popq	%rdi
	.cfi_def_cfa_offset 216
	popq	%r8
	.cfi_def_cfa_offset 208
.L131:
	call	cudaThreadSynchronize@PLT
	movl	$1195, %edx
	movq	%r12, %rsi
	incl	%r15d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	jmp	.L129
.L198:
	incl	%ebp
	jmp	.L126
.L143:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	cmpb	$0, 55(%rsp)
	movq	124(%rsp), %rdx
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	je	.L135
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L137
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z69__device_stub__Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iimPK6float3PS_iim
.L137:
	call	cudaThreadSynchronize@PLT
	movl	$1206, %edx
	movq	%r15, %rsi
	xorl	%r12d, %r12d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L139:
	cmpl	208(%rsp), %r12d
	jl	.L142
	incl	%r13d
.L127:
	cmpl	%r14d, %r13d
	jl	.L143
	jmp	.L153
.L135:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L137
	movq	72(%rsp), %r8
	movl	12(%rsp), %edx
	movl	%ebx, %ecx
	movq	96(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z71__device_stub__Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iimPK6float3PS_iim
	jmp	.L137
.L142:
	movq	124(%rsp), %rdx
	movq	112(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L140
	movq	80(%rsp), %rdi
	pushq	$0
	.cfi_def_cfa_offset 216
	movl	%ebx, %r9d
	pushq	80(%rsp)
	.cfi_def_cfa_offset 224
	movl	28(%rsp), %r8d
	movss	64(%rsp), %xmm1
	movq	%rdi, %rsi
	movss	48(%rsp), %xmm0
	movq	112(%rsp), %rcx
	movq	120(%rsp), %rdx
	call	_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi
	popq	%rcx
	.cfi_def_cfa_offset 216
	popq	%rsi
	.cfi_def_cfa_offset 208
.L140:
	call	cudaThreadSynchronize@PLT
	movl	$1214, %edx
	movq	%r15, %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	movq	124(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L141
	movq	80(%rsp), %rdi
	pushq	$1
	.cfi_def_cfa_offset 216
	movl	%ebx, %r9d
	pushq	80(%rsp)
	.cfi_def_cfa_offset 224
	movl	28(%rsp), %r8d
	movss	64(%rsp), %xmm1
	movq	%rdi, %rsi
	movss	48(%rsp), %xmm0
	movq	112(%rsp), %rcx
	movq	120(%rsp), %rdx
	call	_Z52__device_stub__Z10sor_sharedP6float3PKS_S2_S2_ffiimiP6float3PKS_S2_S2_ffiimi
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L141:
	call	cudaThreadSynchronize@PLT
	movl	$1221, %edx
	movq	%r15, %rsi
	incl	%r12d
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	jmp	.L139
.L125:
	cmpl	$1, %r12d
	jne	.L199
.L150:
	movslq	12(%rsp), %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 216
	movslq	%ebx, %r9
	pushq	$2
	.cfi_def_cfa_offset 224
	movq	104(%rsp), %rdx
	movq	40(%rsp), %rdi
	movq	80(%rsp), %rcx
	salq	$2, %rsi
	movq	%rsi, %r8
	call	cudaMemcpy2D@PLT
	popq	%r12
	.cfi_def_cfa_offset 216
	movl	$1230, %edx
	popq	%r13
	.cfi_def_cfa_offset 208
	movl	%eax, %edi
	jmp	.L194
.L199:
	cmpl	$3, %r12d
	jne	.L144
.L153:
	movslq	12(%rsp), %rsi
	movq	72(%rsp), %rcx
	movslq	%ebx, %r9
	imulq	$12, %rsi, %rsi
	cmpl	$3, %ebp
	leaq	.LC4(%rip), %rbp
	jne	.L145
	pushq	%r10
	.cfi_def_cfa_offset 216
	movq	%rsi, %r8
	pushq	$2
	.cfi_def_cfa_offset 224
	movq	96(%rsp), %rdx
	movq	40(%rsp), %rdi
	call	cudaMemcpy2D@PLT
	popq	%r11
	.cfi_def_cfa_offset 216
	movl	$1232, %edx
	popq	%rbx
	.cfi_def_cfa_offset 208
	movl	%eax, %edi
	jmp	.L193
.L145:
	pushq	%rdi
	.cfi_def_cfa_offset 216
	movq	%rsi, %r8
	pushq	$2
	.cfi_def_cfa_offset 224
	movq	104(%rsp), %rdx
	movq	40(%rsp), %rdi
	call	cudaMemcpy2D@PLT
	popq	%r8
	.cfi_def_cfa_offset 216
	movl	$1233, %edx
	popq	%r9
	.cfi_def_cfa_offset 208
	movl	%eax, %edi
.L193:
	movq	%rbp, %rsi
.L192:
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L144:
	movq	104(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L146
	call	cudaFree@PLT
	movl	$1238, %edx
	leaq	.LC4(%rip), %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L146:
	movq	96(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L147
	call	cudaFree@PLT
	movl	$1239, %edx
	leaq	.LC4(%rip), %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L147:
	movq	88(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L148
	call	cudaFree@PLT
	movl	$1240, %edx
	leaq	.LC4(%rip), %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
.L148:
	movq	80(%rsp), %rdi
	testq	%rdi, %rdi
	je	.L74
	call	cudaFree@PLT
	movl	$1241, %edx
	leaq	.LC4(%rip), %rsi
	movl	%eax, %edi
	call	_Z14__cudaSafeCall9cudaErrorPKci
	jmp	.L74
.L197:
	movslq	12(%rsp), %rsi
	pushq	%rax
	.cfi_def_cfa_offset 216
	movslq	%ebx, %r9
	pushq	$2
	.cfi_def_cfa_offset 224
	movq	40(%rsp), %rdi
	movq	80(%rsp), %rcx
	movq	96(%rsp), %rdx
	salq	$2, %rsi
	movq	%rsi, %r8
	call	cudaMemcpy2D@PLT
	popq	%rdx
	.cfi_def_cfa_offset 216
	movl	$1229, %edx
	popq	%rcx
	.cfi_def_cfa_offset 208
	movl	%eax, %edi
.L194:
	leaq	.LC4(%rip), %rsi
	jmp	.L192
.L74:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L154
	call	__stack_chk_fail@PLT
.L154:
	addq	$152, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z13gpu_diffusionPKfPfiiifififib, .-_Z13gpu_diffusionPKfPfiiifififib
	.section	.rodata.str1.1
.LC5:
	.string	"_Z10sor_sharedP6float3PKS_S2_S2_ffiimi"
.LC6:
	.string	"_Z10sor_sharedPfPKfS1_S1_ffiimi"
.LC7:
	.string	"_Z13jacobi_sharedP6float3PKS_S2_S2_fiim"
.LC8:
	.string	"_Z13jacobi_sharedPfPKfS1_S1_fiim"
.LC9:
	.string	"_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim"
.LC10:
	.string	"_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim"
.LC11:
	.string	"_Z29compute_tv_diffusivity_sharedPKfPfiim"
.LC12:
	.string	"_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim"
.LC13:
	.string	"_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim"
.LC14:
	.string	"_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim"
.LC15:
	.string	"_Z29diffuse_linear_isotrop_sharedPKfPffiim"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2081:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10sor_sharedP6float3PKS_S2_S2_ffiimi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z10sor_sharedPfPKfS1_S1_ffiimi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z13jacobi_sharedP6float3PKS_S2_S2_fiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC8(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z13jacobi_sharedPfPKfS1_S1_fiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC9(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z38compute_tv_diffusivity_separate_sharedPK6float3PS_iim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC10(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z36compute_tv_diffusivity_joined_sharedPK6float3PS_iim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z29compute_tv_diffusivity_sharedPKfPfiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC12(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z32diffuse_nonlinear_isotrop_sharedPK6float3S1_PS_fiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	leaq	.LC13(%rip), %rdx
	pushq	$0
	.cfi_def_cfa_offset 24
	movq	%rdx, %rcx
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	_Z32diffuse_nonlinear_isotrop_sharedPKfS0_Pffiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC14(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z29diffuse_linear_isotrop_sharedPK6float3PS_fiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC15(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z29diffuse_linear_isotrop_sharedPKfPffiim(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2081:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	studentName
	.section	.rodata.str1.1
.LC16:
	.string	"Arash Bakhtiari"
	.section	.data.rel.local,"aw"
	.align 8
	.type	studentName, @object
	.size	studentName, 8
studentName:
	.quad	.LC16
	.globl	studentLogin
	.section	.rodata.str1.1
.LC17:
	.string	"p116"
	.section	.data.rel.local
	.align 8
	.type	studentLogin, @object
	.size	studentLogin, 8
studentLogin:
	.quad	.LC17
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC3:
	.long	1031798784
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

﻿/****************************************************************************\
*      --- Practical Course: GPU Programming in Computer Vision ---
*
* time:    winter term 2012/13 / March 11-18, 2013
*
* project: diffusion
* file:    diffusion.cuh
*
*
*             !THIS FILE IS SUPPOSED TO REMAIN UNCHANGED!
\****************************************************************************/


#ifndef __DIFFUSION_CUH
#define __DIFFUSION_CUH

#include <stdio.h>
#include <hip/hip_runtime.h>
#ifndef cutilSafeCall
#define cutilSafeCall(err)  __cudaSafeCall(err,__FILE__,__LINE__)
inline void __cudaSafeCall(hipError_t err,
                           const char *file, const int line){
  if(hipSuccess != err) {
    fprintf(stderr,"%s(%i) : cutilSafeCall() Runtime API error : %s.\n",
           file, line, hipGetErrorString(err) );
    exit(-1);
  }
}
#endif



const char* getStudentLogin();
const char* getStudentName();
int         getStudentID();
bool        checkStudentData();
bool        checkStudentNameAndID();


__global__ void diffuse_linear_isotrop_shared(const float *d_input,float *d_output,float timeStep, int nx,int ny, size_t pitch);

__global__ void diffuse_linear_isotrop_shared(const float3 *d_input,float3 *d_output,float timeStep,int nx, int ny, size_t pitchBytes);
 
__global__ void diffuse_nonlinear_isotrop_shared(const float *d_input,const float *d_diffusivity,float *d_output,float timeStep,int nx,int ny, size_t pitch);
 
__global__ void diffuse_nonlinear_isotrop_shared(const float3 *d_input,const float3 *d_diffusivity,float3 *d_output,float timeStep, int nx, int ny, size_t pitchBytes);

__global__ void compute_tv_diffusivity_shared(const float *d_input,float *d_output,int nx,int ny,size_t pitch);

__global__ void compute_tv_diffusivity_joined_shared(const float3 *d_input,float3 *d_output,int nx,int ny,size_t pitchBytes);

__global__ void compute_tv_diffusivity_separate_shared(const float3 *d_input,float3 *d_output,int nx,int ny,size_t pitchBytes);
 
__global__ void jacobi_shared(float *d_output,const float *d_input,const float *d_original,const float *d_diffusivity,float weight,int nx,int ny,size_t pitch);
 
__global__ void jacobi_shared(float3 *d_output,const float3 *d_input,const float3 *d_original,const float3 *d_diffusivity,float weight,int nx,int ny,size_t pitchBytes);
 
__global__ void sor_shared(float *d_output,const float *d_input,const float *d_original,const float *d_diffusivity,float weight,float overrelaxation,int nx,int ny,size_t pitch,int red);
 
__global__ void sor_shared(float3 *d_output,const float3 *d_input,const float3 *d_original,const float3 *d_diffusivity,float weight,float overrelaxation,int nx,int ny,size_t pitchBytes,int red);
 
void gpu_diffusion(const float *input,float *output,int nx, int ny, int nc,float timeStep,int iterations,float weight,int lagged_iterations,float overrelaxation,int mode,bool jointDiffusivity=false);


#endif // #ifndef __DIFFUSION_CUH

	.text
	.file	"diffusion.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
