/*
   Program name: MatrixTranspose.cu
    Author name: Dr. Nileshchandra Pikle
           Email: nilesh.pikle@gmail.com
  Contact Number: 7276834418  
         Webpage: https://piklenileshchandra.wixsite.com/personal
    
    Purpose: To perform Matrix Transpose using CUDA
             
    Description: Matrix transpose program is considered to demonstrate 2D thread block and 2D grid.

    Three versions of matrix transpose are created here
    1. Matrix transpose using single thread (sequential)
    2. Matrix transpose using N threads (1D thread block) where N is #rows in matrix
    3. Matrix transpose using N x N threads (2D thread block) where N is #rows in matrix
    
   *In first version a single thread on GPU responsible to perform entire transpose operation hence it 
    is a sequetial operation. This version performs worst than sequential CPU as GPU core is lightweight
   *In second version each thread is responsible to take a row of matrix and store into new matrix in 
    column. Consider this as coarse grained. 
   *In third version each thread is responsible to take a single element of matrix and store into new 
    matrix at transposed position. Consider this as coarse grained. 

    Note: 3rd version is NOT optimized. However it outperforms first and second version. Shared memory 
          level optimizations can be performed to accelerate execution even further. To see how to 
          optimize second version refer link below
          Link: https://devblogs.nvidia.com/efficient-matrix-transpose-cuda-cc/


*/

#include<stdio.h>
#include<stdio.h>
#include<math.h>
#include<time.h>

#define CHECK(call)                                                     \
{									\
  const cudaError_t error = call;					\
  if(error != cudaSuccess)						\
  {									\
    printf("Error %s %d", __FILE__, __LINE__);				\
    printf("\n Code %d Reason %s \n",error, cudaGetErrorString(error));	\
    exit(1);								\
  }									\
}	
 __global__ void transposeKernel_single_thread(int *d_input, int *d_output, int MAT_SIZE)
 {
   int i,j;
   for(i = 0; i < MAT_SIZE; i++)
   {
     for(j = 0; j < MAT_SIZE; j++)
     {
       d_output[j * MAT_SIZE + i] = d_input[i * MAT_SIZE + j];
     }
   } 
 }

 __global__ void transposeKernel_thread_per_row(int *d_input, int *d_output, int MAT_SIZE)
 {
    int j;
    int gid = threadIdx.x + blockIdx.x *blockDim.x;
    if(gid < MAT_SIZE)
    {
      for(j = 0; j < MAT_SIZE; j++)
      {
        d_output[gid * MAT_SIZE + j] = d_input[j * MAT_SIZE + gid];
      }
    }
 }

  __global__ void transposeKernel_thread_per_element(int *d_input, int *d_output, int MAT_SIZE)
  {
    int row,col;
    col = threadIdx.x + blockIdx.x * blockDim.x;
    row = threadIdx.y + blockIdx.y * blockDim.y;

    if( (row < MAT_SIZE) && (col < MAT_SIZE) )
    { 
     d_output[row * MAT_SIZE + col] = d_input[col * MAT_SIZE + row];
    }
  }

  void initMatrix(int *matrix, int MAT_SIZE, int flag)
  {
    int i,j;
    if(flag == 0)
    {
      for(i = 0; i < MAT_SIZE; i++)
      {
        for(j = 0; j < MAT_SIZE; j++)
        {
          matrix[i * MAT_SIZE + j] = 0;
        }
      }
     }else{
            for(i = 0; i < MAT_SIZE; i++)
            {
             for(j = 0; j < MAT_SIZE; j++)
             {
               matrix[i * MAT_SIZE + j] = i;
               //printf(" %d ",matrix[i * MAT_SIZE + j]);
             }
            //printf("\n");
            }
           }
    }

    int checkResult(int *hg_output, int *h_output, int MAT_SIZE)
    {
       int i,j, flag = 1;
       for(i = 0; i < MAT_SIZE; i++)
       {
         for(j = 0; j < MAT_SIZE; j++)
         {
           if(hg_output[i*MAT_SIZE+j] != h_output[i*MAT_SIZE+j])
           {
             flag = 0;
           }
         }
       }
       return flag;
    }
    
    __global__ void init_kernel(int *d_output, int MAT_SIZE)
    {
       int i,j;
       for(i = 0; i < MAT_SIZE; i++)
       {
         for(j = 0; j < MAT_SIZE; j++)
         {
            d_output[i * MAT_SIZE + j] = 0;
         }
       }
    }

  int main()
  {
    int i,j; 
    int MAT_SIZE = 2048;   // Matrix size MAT_SIZE x MAT_SIZE
    int *h_input  = (int *)malloc(MAT_SIZE*MAT_SIZE*sizeof(int));
    int *h_output = (int *)malloc(MAT_SIZE*MAT_SIZE*sizeof(int));
    int *hg_output = (int *)malloc(MAT_SIZE*MAT_SIZE*sizeof(int));
    
    if(!h_input)
    {
     printf("\n Error: Allocating memory to h_input!!!");
    }
    if(!h_output)
    {
     printf("\n Error: Allocating memory to h_output!!!");
    }
    if(!hg_output)
    {
     printf("\n Error: Allocating memory to h_output!!!");
    }
    /*
      Data initialization for input matrix h_input and h_output matrix
       
    */

     initMatrix(h_output,MAT_SIZE,0);  // matrix initialization function 0 indicates initialize with 0
     initMatrix(hg_output,MAT_SIZE,0); // matrix initialization function 0 indicates initialize with 0
     initMatrix(h_input,MAT_SIZE,1);  // matrix initialization function 0 indicates initialize with row index (refer function)

    /*
      Matrix transpose code: h_output stores the transposed output of matrix of h_input       
    */
    clock_t t;     // clock function to calculate execution time of sequential program
    t = clock();   // record start time

    for(i = 0; i < MAT_SIZE; i++)
    {
      for(j = 0; j < MAT_SIZE; j++)
      {
       h_output[i * MAT_SIZE + j] = h_input[j * MAT_SIZE + i];
       //printf(" %d ", h_output[i * MAT_SIZE + j]);
      }
      //printf("\n");
    }
   t = clock() - t;
   double CPUtime_taken = ((double)t)/CLOCKS_PER_SEC;
   printf("\n     Time required for sequential execution is %.3f ms\n", CPUtime_taken*1000);

    /*
       Parallel CUDA programs start here
    */

    /* Allocate memory to d_input and d_output array on device */
    int *d_input, *d_output;
    CHECK(cudaMalloc((void **)&d_input, MAT_SIZE*MAT_SIZE*sizeof(int)));
    CHECK(cudaMalloc((void **)&d_output, MAT_SIZE*MAT_SIZE*sizeof(int)));
    
    /* Transfer data from host memory to device memory h_input to d_input and h_output to d_output array on device*/ 

    CHECK(cudaMemcpy(d_input, h_input, MAT_SIZE*MAT_SIZE*sizeof(int), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_output, h_output, MAT_SIZE*MAT_SIZE*sizeof(int), cudaMemcpyHostToDevice));

/**********************************************************************************************/
    /* 1. Transpose kerene launch using single thread
          A thread block of single thread is configured to perform matrix transpose
          This is exactly like sequential version of matrix transpose as only one 
          thread is responsible to perform entire matrix transpose operation          
    */
     
  /* Declaring Time variables to measure GPU time*/
  float GPUelapsed = 0.0;  // To store final kernel execution time 
  cudaEvent_t start, stop; // Variables to record start and stop of kernel
  cudaEventCreate(&start); // Event create start
  cudaEventCreate(&stop);  // Event create stop
  cudaEventRecord(start, 0); // Record time at start variables

/**/transposeKernel_single_thread<<<1,1>>>(d_input,d_output,MAT_SIZE); // krenel for matrix transpose

  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&GPUelapsed, start, stop);
  cudaEventDestroy(start); // Event destroy start
  cudaEventDestroy(stop);  // Event destroy stop
  printf("\n 1. Parallel runtime Single thread %0.3f ms",GPUelapsed);
  printf("\n    Parallel Speedup Single thread = (CPU time)/(GPU time) = %0.3f\n", (CPUtime_taken*1000)/GPUelapsed);
    
  CHECK(cudaMemcpy(hg_output, d_output, MAT_SIZE*MAT_SIZE*sizeof(int), cudaMemcpyDeviceToHost)); // copy results back on host

    int check = checkResult(hg_output, h_output, MAT_SIZE); // This function checks two matrices are same or not
    if(check == 0)printf("\n Something went wrong => Results are not matching for 'transposeKernel_single_thread' !!!");
    initMatrix(hg_output,MAT_SIZE,0); // Initialize hg_output matrix to 0
    init_kernel<<<1,1>>>(d_output, MAT_SIZE); // Initialize d_output matrix to 0
/**********************************************************************************************/

    /* 2. Transpose kerene launch using thread per row
          In this configuration, # threads launched are equalt to # rows of matrix i.e MAT_SIZE
          Where each thread is responsible to perform transpose operation.
          Each thread reads the matrix entries row major order from d_input matrix and stores in 
          d_output matrix in column major order.    

    # threads launched = # rows in matrix
    numT = # threads per thread block
    numB = # thread blocks in the grid
  */

   int numT, numB;
   numT = 128;
   numB = ceil(MAT_SIZE/(float)numT);
  /* Declaring Time variables to measure GPU time*/
  float GPUelapsed2 = 0.0;  // To store final kernel execution time 
  cudaEvent_t start2, stop2; // Variables to record start and stop of kernel
  cudaEventCreate(&start2); // Event create start
  cudaEventCreate(&stop2);  // Event create stop
  cudaEventRecord(start2, 0); // Record time at start variables

/**/ transposeKernel_thread_per_row<<<numB,numT>>>(d_input,d_output,MAT_SIZE); // krenel for matrix transpose

  cudaEventRecord(stop2, 0);
  cudaEventSynchronize(stop2);
  cudaEventElapsedTime(&GPUelapsed2, start2, stop2);
  cudaEventDestroy(start2); // Event destroy start
  cudaEventDestroy(stop2);  // Event destroy stop
  printf("\n 2. Parallel runtime thread per row is %0.3f ms",GPUelapsed2);
  printf("\n    Parallel Speedup thread per row = (CPU time)/(GPU time) = %0.3f\n", (CPUtime_taken*1000)/GPUelapsed2);
    
  CHECK(cudaMemcpy(hg_output, d_output, MAT_SIZE*MAT_SIZE*sizeof(int), cudaMemcpyDeviceToHost)); // copy results back on host

    check = checkResult(hg_output, h_output, MAT_SIZE); // This function checks two matrices are same or not
    if(check == 0)printf("\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_row'!!!");
    initMatrix(hg_output,MAT_SIZE,0); // Initialize hg_output matrix to 0
    init_kernel<<<1,1>>>(d_output, MAT_SIZE); // Initialize d_output matrix to 0


/**********************************************************************************************/
    /* 3. Transpose kerene launch using thread per element of the matrix
          In this configuration, # threads launched are equalt to # of elements in the matrix i.e MAT_SIZE x MAT_SIZE
          Where each thread is responsible to perform transpose operation on a single element.
          Each thread reads a single corresponding matrix element from d_input matrix and stores in 
          d_output matrix in transposed indices.    

    # threads launched = # elements in matrix
    numT = # threads per thread block in 2D
    numB = # thread blocks in the grid in 2D
  */

   dim3 num2T(8, 8,1);
   dim3 num2B(ceil(MAT_SIZE/(float)num2T.x), ceil(MAT_SIZE/(float)num2T.y ),1 );
   //printf("\n num2B.x = %d num2B.y = %d ", num2B.x,num2B.x );
  /* Declaring Time variables to measure GPU time */
  float GPUelapsed3 = 0.0;  // To store final kernel execution time 
  cudaEvent_t start3, stop3; // Variables to record start and stop of kernel
  cudaEventCreate(&start3); // Event create start
  cudaEventCreate(&stop3);  // Event create stop
  cudaEventRecord(start3, 0); // Record time at start variables

/**/ transposeKernel_thread_per_element<<<num2B,num2T>>>(d_input,d_output,MAT_SIZE); // krenel for matrix transpose

  cudaEventRecord(stop3, 0);
  cudaEventSynchronize(stop3);
  cudaEventElapsedTime(&GPUelapsed3, start3, stop3);
  cudaEventDestroy(start3); // Event destroy start
  cudaEventDestroy(stop3);  // Event destroy stop
  printf("\n 3. Parallel runtime thread per element is %0.3f ms",GPUelapsed3);
  printf("\n    Parallel Speedup thread per element = (CPU time)/(GPU time) = %0.3f\n", (CPUtime_taken*1000)/GPUelapsed3);
    
  CHECK(cudaMemcpy(hg_output, d_output, MAT_SIZE*MAT_SIZE*sizeof(int), cudaMemcpyDeviceToHost)); // copy results back on host

    check = checkResult(hg_output, h_output, MAT_SIZE); // This function checks two matrices are same or not
    if(check == 0)
   {
    printf("\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_element'!!!");
   }

   free(h_input);
   free(h_output);
   free(hg_output);
   cudaFree(d_input);
   cudaFree(d_output);
   return(0);    
  }

	.file	"tmpxft_002395ed_00000000-6_13_MatrixTranspose.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10initMatrixPiii.part.0, @function
_Z10initMatrixPiii.part.0:
.LFB2179:
	.cfi_startproc
	leal	-1(%rsi), %eax
	movslq	%esi, %r9
	movq	%rdi, %rdx
	xorl	%r8d, %r8d
	leaq	4(,%rax,4), %r10
	salq	$2, %r9
	xorl	%eax, %eax
.L3:
	cmpl	%esi, %r8d
	jge	.L6
	movq	%rdx, %rdi
	movq	%r10, %rcx
	incl	%r8d
	addq	%r9, %rdx
	rep stosb
	jmp	.L3
.L6:
	ret
	.cfi_endproc
.LFE2179:
	.size	_Z10initMatrixPiii.part.0, .-_Z10initMatrixPiii.part.0
	.globl	_Z10initMatrixPiii
	.type	_Z10initMatrixPiii, @function
_Z10initMatrixPiii:
.LFB2027:
	.cfi_startproc
	endbr64
	testl	%edx, %edx
	je	.L8
	movslq	%esi, %rcx
	xorl	%eax, %eax
	salq	$2, %rcx
	jmp	.L9
.L8:
	jmp	_Z10initMatrixPiii.part.0
.L10:
	movl	%eax, (%rdi,%rdx,4)
	incq	%rdx
	cmpl	%edx, %esi
	jg	.L10
	incl	%eax
	addq	%rcx, %rdi
.L9:
	cmpl	%esi, %eax
	jge	.L13
	xorl	%edx, %edx
	jmp	.L10
.L13:
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z10initMatrixPiii, .-_Z10initMatrixPiii
	.globl	_Z11checkResultPiS_i
	.type	_Z11checkResultPiS_i, @function
_Z11checkResultPiS_i:
.LFB2028:
	.cfi_startproc
	endbr64
	movslq	%edx, %r10
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	xorl	%r9d, %r9d
	movl	$1, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	salq	$2, %r10
	xorl	%r8d, %r8d
	xorl	%ebx, %ebx
.L15:
	cmpl	%edx, %r8d
	jge	.L22
	movq	%r9, %rax
	xorl	%r11d, %r11d
.L17:
	movl	(%rsi,%rax), %r14d
	cmpl	%r14d, (%rdi,%rax)
	cmovne	%ebx, %ecx
	incl	%r11d
	addq	$4, %rax
	cmpl	%r11d, %edx
	jne	.L17
	incl	%r8d
	addq	%r10, %r9
	jmp	.L15
.L22:
	movl	%ecx, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z11checkResultPiS_i, .-_Z11checkResultPiS_i
	.globl	_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i
	.type	_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i, @function
_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%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)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z29transposeKernel_single_threadPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L23:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i, .-_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i
	.globl	_Z29transposeKernel_single_threadPiS_i
	.type	_Z29transposeKernel_single_threadPiS_i, @function
_Z29transposeKernel_single_threadPiS_i:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i
	.cfi_endproc
.LFE2055:
	.size	_Z29transposeKernel_single_threadPiS_i, .-_Z29transposeKernel_single_threadPiS_i
	.globl	_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i
	.type	_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i, @function
_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i:
.LFB2056:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%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)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z30transposeKernel_thread_per_rowPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L28:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i, .-_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i
	.globl	_Z30transposeKernel_thread_per_rowPiS_i
	.type	_Z30transposeKernel_thread_per_rowPiS_i, @function
_Z30transposeKernel_thread_per_rowPiS_i:
.LFB2057:
	.cfi_startproc
	endbr64
	jmp	_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i
	.cfi_endproc
.LFE2057:
	.size	_Z30transposeKernel_thread_per_rowPiS_i, .-_Z30transposeKernel_thread_per_rowPiS_i
	.globl	_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i
	.type	_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i, @function
_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i:
.LFB2058:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%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)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z34transposeKernel_thread_per_elementPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	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 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L33:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L35:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2058:
	.size	_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i, .-_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i
	.globl	_Z34transposeKernel_thread_per_elementPiS_i
	.type	_Z34transposeKernel_thread_per_elementPiS_i, @function
_Z34transposeKernel_thread_per_elementPiS_i:
.LFB2059:
	.cfi_startproc
	endbr64
	jmp	_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i
	.cfi_endproc
.LFE2059:
	.size	_Z34transposeKernel_thread_per_elementPiS_i, .-_Z34transposeKernel_thread_per_elementPiS_i
	.globl	_Z32__device_stub__Z11init_kernelPiiPii
	.type	_Z32__device_stub__Z11init_kernelPiiPii, @function
_Z32__device_stub__Z11init_kernelPiiPii:
.LFB2060:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L38
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z11init_kernelPii(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L38:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L40
	call	__stack_chk_fail@PLT
.L40:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2060:
	.size	_Z32__device_stub__Z11init_kernelPiiPii, .-_Z32__device_stub__Z11init_kernelPiiPii
	.globl	_Z11init_kernelPii
	.type	_Z11init_kernelPii, @function
_Z11init_kernelPii:
.LFB2061:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z11init_kernelPiiPii
	.cfi_endproc
.LFE2061:
	.size	_Z11init_kernelPii, .-_Z11init_kernelPii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\n Error: Allocating memory to h_input!!!"
.LC1:
	.string	"\n Error: Allocating memory to h_output!!!"
.LC4:
	.string	"\n     Time required for sequential execution is %.3f ms\n"
.LC5:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/NileshchandraPikle/Introduction-to-CUDA-Programming/master/13_MatrixTranspose.cu"
.LC6:
	.string	"Error %s %d"
.LC7:
	.string	"\n Code %d Reason %s \n"
.LC9:
	.string	"\n 1. Parallel runtime Single thread %0.3f ms"
.LC10:
	.string	"\n    Parallel Speedup Single thread = (CPU time)/(GPU time) = %0.3f\n"
.LC11:
	.string	"\n Something went wrong => Results are not matching for 'transposeKernel_single_thread' !!!"
.LC12:
	.string	"\n 2. Parallel runtime thread per row is %0.3f ms"
.LC13:
	.string	"\n    Parallel Speedup thread per row = (CPU time)/(GPU time) = %0.3f\n"
.LC14:
	.string	"\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_row'!!!"
.LC15:
	.string	"\n 3. Parallel runtime thread per element is %0.3f ms"
.LC16:
	.string	"\n    Parallel Speedup thread per element = (CPU time)/(GPU time) = %0.3f\n"
.LC17:
	.string	"\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_element'!!!"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$16777216, %edi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movl	$16777216, %edi
	movq	%rax, %r12
	call	malloc@PLT
	movl	$16777216, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%rax, %rbx
	testq	%r12, %r12
	jne	.L44
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L44:
	testq	%rbp, %rbp
	jne	.L45
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L45:
	testq	%rbx, %rbx
	jne	.L46
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L46:
	movl	$2048, %esi
	movq	%rbp, %rdi
	call	_Z10initMatrixPiii.part.0
	movq	%rbx, %rdi
	call	_Z10initMatrixPiii.part.0
	movl	$1, %edx
	movq	%r12, %rdi
	call	_Z10initMatrixPiii
	call	clock@PLT
	xorl	%edx, %edx
	movq	%rax, %r13
.L47:
	movq	%rdx, %rax
	leaq	(%r12,%rdx,4), %rsi
	xorl	%ecx, %ecx
	salq	$13, %rax
	addq	%rbp, %rax
.L48:
	movl	(%rsi), %edi
	addq	$8192, %rsi
	movl	%edi, (%rax,%rcx,4)
	incq	%rcx
	cmpq	$2048, %rcx
	jne	.L48
	incq	%rdx
	cmpq	$2048, %rdx
	jne	.L47
	call	clock@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	subq	%r13, %rax
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC2(%rip), %xmm0
	movb	$1, %al
	mulsd	.LC3(%rip), %xmm0
	movq	%xmm0, %r14
	call	__printf_chk@PLT
	leaq	16(%rsp), %rdi
	movl	$16777216, %esi
	call	cudaMalloc@PLT
	movl	$191, %ecx
	movl	%eax, %r13d
	testl	%eax, %eax
	jne	.L68
	leaq	24(%rsp), %rdi
	movl	$16777216, %esi
	call	cudaMalloc@PLT
	movl	%eax, %r13d
	testl	%eax, %eax
	je	.L51
	movl	$192, %ecx
.L68:
	leaq	.LC5(%rip), %rdx
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%r13d, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	movl	%r13d, %edx
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L51:
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movl	$16777216, %edx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movl	$196, %ecx
	movl	%eax, %r13d
	testl	%eax, %eax
	jne	.L68
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	movl	$16777216, %edx
	movq	%rbp, %rsi
	call	cudaMemcpy@PLT
	movl	$197, %ecx
	movl	%eax, %r13d
	testl	%eax, %eax
	jne	.L68
	leaq	32(%rsp), %rdi
	movl	$0x00000000, 4(%rsp)
	call	cudaEventCreate@PLT
	leaq	40(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movl	$1, 100(%rsp)
	movq	%rdi, %rdx
	movq	%rdi, 92(%rsp)
	movq	%rdi, 80(%rsp)
	movl	$1, 88(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2048, %edx
	call	_Z52__device_stub__Z29transposeKernel_single_threadPiS_iPiS_i
.L54:
	movq	40(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	40(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	leaq	4(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	32(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	40(%rsp), %rdi
	call	cudaEventDestroy@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	4(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r14, %xmm0
	movb	$1, %al
	cvtss2sd	4(%rsp), %xmm1
	divsd	%xmm1, %xmm0
	movl	$2, %edi
	leaq	.LC10(%rip), %rsi
	call	__printf_chk@PLT
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$16777216, %edx
	call	cudaMemcpy@PLT
	movl	$223, %ecx
	movl	%eax, %r13d
	testl	%eax, %eax
	jne	.L68
	movl	$2048, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z11checkResultPiS_i
	testl	%eax, %eax
	jne	.L56
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
.L56:
	movl	$2048, %esi
	movq	%rbx, %rdi
	call	_Z10initMatrixPiii.part.0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movl	$1, 100(%rsp)
	movq	%rdi, %rdx
	movq	%rdi, 92(%rsp)
	movq	%rdi, 80(%rsp)
	movl	$1, 88(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L57
	movq	24(%rsp), %rdi
	movl	$2048, %esi
	call	_Z32__device_stub__Z11init_kernelPiiPii
.L57:
	leaq	48(%rsp), %rdi
	movl	$0x00000000, 8(%rsp)
	call	cudaEventCreate@PLT
	leaq	56(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$33554433, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$268435457, %edi
	salq	$7, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$4, %rdi
	movq	%rdx, 92(%rsp)
	movl	$1, 100(%rsp)
	movq	%rdi, 80(%rsp)
	movl	$1, 88(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L58
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2048, %edx
	call	_Z53__device_stub__Z30transposeKernel_thread_per_rowPiS_iPiS_i
.L58:
	movq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	56(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	leaq	8(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	48(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	56(%rsp), %rdi
	call	cudaEventDestroy@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	8(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r14, %xmm0
	movb	$1, %al
	cvtss2sd	8(%rsp), %xmm1
	divsd	%xmm1, %xmm0
	movl	$2, %edi
	leaq	.LC13(%rip), %rsi
	call	__printf_chk@PLT
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$16777216, %edx
	call	cudaMemcpy@PLT
	movl	$262, %ecx
	movl	%eax, %r13d
	testl	%eax, %eax
	jne	.L68
	movl	$2048, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z11checkResultPiS_i
	testl	%eax, %eax
	jne	.L60
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
.L60:
	movl	$2048, %esi
	movq	%rbx, %rdi
	call	_Z10initMatrixPiii.part.0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movl	$1, 100(%rsp)
	movq	%rdi, %rdx
	movq	%rdi, 92(%rsp)
	movq	%rdi, 80(%rsp)
	movl	$1, 88(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L61
	movq	24(%rsp), %rdi
	movl	$2048, %esi
	call	_Z32__device_stub__Z11init_kernelPiiPii
.L61:
	movabsq	$1099511628032, %rax
	leaq	64(%rsp), %rdi
	movl	$1, 100(%rsp)
	movq	%rax, 92(%rsp)
	movl	$0x00000000, 12(%rsp)
	call	cudaEventCreate@PLT
	leaq	72(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	64(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	92(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	100(%rsp), %esi
	movl	$1, %ecx
	movabsq	$34359738376, %rdx
	movl	$1, 88(%rsp)
	movq	%rdx, 80(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L62
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2048, %edx
	call	_Z57__device_stub__Z34transposeKernel_thread_per_elementPiS_iPiS_i
.L62:
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	72(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	72(%rsp), %rdx
	movq	64(%rsp), %rsi
	leaq	12(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	64(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	72(%rsp), %rdi
	call	cudaEventDestroy@PLT
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	12(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r14, %xmm0
	movb	$1, %al
	cvtss2sd	12(%rsp), %xmm1
	divsd	%xmm1, %xmm0
	movl	$2, %edi
	leaq	.LC16(%rip), %rsi
	call	__printf_chk@PLT
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$16777216, %edx
	call	cudaMemcpy@PLT
	movl	$302, %ecx
	movl	%eax, %r13d
	testl	%eax, %eax
	jne	.L68
	movl	$2048, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z11checkResultPiS_i
	testl	%eax, %eax
	jne	.L64
	leaq	.LC17(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
.L64:
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L65
	call	__stack_chk_fail@PLT
.L65:
	addq	$112, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	main, .-main
	.section	.rodata.str1.1
.LC18:
	.string	"_Z11init_kernelPii"
.LC19:
	.string	"_Z34transposeKernel_thread_per_elementPiS_i"
.LC20:
	.string	"_Z30transposeKernel_thread_per_rowPiS_i"
.LC21:
	.string	"_Z29transposeKernel_single_threadPiS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2063:
	.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	.LC18(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z11init_kernelPii(%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
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC19(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z34transposeKernel_thread_per_elementPiS_i(%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	.LC20(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z30transposeKernel_thread_per_rowPiS_i(%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	.LC21(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z29transposeKernel_single_threadPiS_i(%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
.LFE2063:
	.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
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	0
	.long	1093567616
	.align 8
.LC3:
	.long	0
	.long	1083129856
	.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:

/*
   Program name: MatrixTranspose.cu
    Author name: Dr. Nileshchandra Pikle
           Email: nilesh.pikle@gmail.com
  Contact Number: 7276834418  
         Webpage: https://piklenileshchandra.wixsite.com/personal
    
    Purpose: To perform Matrix Transpose using CUDA
             
    Description: Matrix transpose program is considered to demonstrate 2D thread block and 2D grid.

    Three versions of matrix transpose are created here
    1. Matrix transpose using single thread (sequential)
    2. Matrix transpose using N threads (1D thread block) where N is #rows in matrix
    3. Matrix transpose using N x N threads (2D thread block) where N is #rows in matrix
    
   *In first version a single thread on GPU responsible to perform entire transpose operation hence it 
    is a sequetial operation. This version performs worst than sequential CPU as GPU core is lightweight
   *In second version each thread is responsible to take a row of matrix and store into new matrix in 
    column. Consider this as coarse grained. 
   *In third version each thread is responsible to take a single element of matrix and store into new 
    matrix at transposed position. Consider this as coarse grained. 

    Note: 3rd version is NOT optimized. However it outperforms first and second version. Shared memory 
          level optimizations can be performed to accelerate execution even further. To see how to 
          optimize second version refer link below
          Link: https://devblogs.nvidia.com/efficient-matrix-transpose-cuda-cc/


*/


#include <hip/hip_runtime.h>
#include<stdio.h>
#include<stdio.h>
#include<math.h>
#include<time.h>

#define CHECK(call)                                                     \
{									\
  const hipError_t error = call;					\
  if(error != hipSuccess)						\
  {									\
    printf("Error %s %d", __FILE__, __LINE__);				\
    printf("\n Code %d Reason %s \n",error, hipGetErrorString(error));	\
    exit(1);								\
  }									\
}	
 __global__ void transposeKernel_single_thread(int *d_input, int *d_output, int MAT_SIZE)
 {
   int i,j;
   for(i = 0; i < MAT_SIZE; i++)
   {
     for(j = 0; j < MAT_SIZE; j++)
     {
       d_output[j * MAT_SIZE + i] = d_input[i * MAT_SIZE + j];
     }
   } 
 }

 __global__ void transposeKernel_thread_per_row(int *d_input, int *d_output, int MAT_SIZE)
 {
    int j;
    int gid = threadIdx.x + blockIdx.x *blockDim.x;
    if(gid < MAT_SIZE)
    {
      for(j = 0; j < MAT_SIZE; j++)
      {
        d_output[gid * MAT_SIZE + j] = d_input[j * MAT_SIZE + gid];
      }
    }
 }

  __global__ void transposeKernel_thread_per_element(int *d_input, int *d_output, int MAT_SIZE)
  {
    int row,col;
    col = threadIdx.x + blockIdx.x * blockDim.x;
    row = threadIdx.y + blockIdx.y * blockDim.y;

    if( (row < MAT_SIZE) && (col < MAT_SIZE) )
    { 
     d_output[row * MAT_SIZE + col] = d_input[col * MAT_SIZE + row];
    }
  }

  void initMatrix(int *matrix, int MAT_SIZE, int flag)
  {
    int i,j;
    if(flag == 0)
    {
      for(i = 0; i < MAT_SIZE; i++)
      {
        for(j = 0; j < MAT_SIZE; j++)
        {
          matrix[i * MAT_SIZE + j] = 0;
        }
      }
     }else{
            for(i = 0; i < MAT_SIZE; i++)
            {
             for(j = 0; j < MAT_SIZE; j++)
             {
               matrix[i * MAT_SIZE + j] = i;
               //printf(" %d ",matrix[i * MAT_SIZE + j]);
             }
            //printf("\n");
            }
           }
    }

    int checkResult(int *hg_output, int *h_output, int MAT_SIZE)
    {
       int i,j, flag = 1;
       for(i = 0; i < MAT_SIZE; i++)
       {
         for(j = 0; j < MAT_SIZE; j++)
         {
           if(hg_output[i*MAT_SIZE+j] != h_output[i*MAT_SIZE+j])
           {
             flag = 0;
           }
         }
       }
       return flag;
    }
    
    __global__ void init_kernel(int *d_output, int MAT_SIZE)
    {
       int i,j;
       for(i = 0; i < MAT_SIZE; i++)
       {
         for(j = 0; j < MAT_SIZE; j++)
         {
            d_output[i * MAT_SIZE + j] = 0;
         }
       }
    }

  int main()
  {
    int i,j; 
    int MAT_SIZE = 2048;   // Matrix size MAT_SIZE x MAT_SIZE
    int *h_input  = (int *)malloc(MAT_SIZE*MAT_SIZE*sizeof(int));
    int *h_output = (int *)malloc(MAT_SIZE*MAT_SIZE*sizeof(int));
    int *hg_output = (int *)malloc(MAT_SIZE*MAT_SIZE*sizeof(int));
    
    if(!h_input)
    {
     printf("\n Error: Allocating memory to h_input!!!");
    }
    if(!h_output)
    {
     printf("\n Error: Allocating memory to h_output!!!");
    }
    if(!hg_output)
    {
     printf("\n Error: Allocating memory to h_output!!!");
    }
    /*
      Data initialization for input matrix h_input and h_output matrix
       
    */

     initMatrix(h_output,MAT_SIZE,0);  // matrix initialization function 0 indicates initialize with 0
     initMatrix(hg_output,MAT_SIZE,0); // matrix initialization function 0 indicates initialize with 0
     initMatrix(h_input,MAT_SIZE,1);  // matrix initialization function 0 indicates initialize with row index (refer function)

    /*
      Matrix transpose code: h_output stores the transposed output of matrix of h_input       
    */
    clock_t t;     // clock function to calculate execution time of sequential program
    t = clock();   // record start time

    for(i = 0; i < MAT_SIZE; i++)
    {
      for(j = 0; j < MAT_SIZE; j++)
      {
       h_output[i * MAT_SIZE + j] = h_input[j * MAT_SIZE + i];
       //printf(" %d ", h_output[i * MAT_SIZE + j]);
      }
      //printf("\n");
    }
   t = clock() - t;
   double CPUtime_taken = ((double)t)/CLOCKS_PER_SEC;
   printf("\n     Time required for sequential execution is %.3f ms\n", CPUtime_taken*1000);

    /*
       Parallel CUDA programs start here
    */

    /* Allocate memory to d_input and d_output array on device */
    int *d_input, *d_output;
    CHECK(hipMalloc((void **)&d_input, MAT_SIZE*MAT_SIZE*sizeof(int)));
    CHECK(hipMalloc((void **)&d_output, MAT_SIZE*MAT_SIZE*sizeof(int)));
    
    /* Transfer data from host memory to device memory h_input to d_input and h_output to d_output array on device*/ 

    CHECK(hipMemcpy(d_input, h_input, MAT_SIZE*MAT_SIZE*sizeof(int), hipMemcpyHostToDevice));
    CHECK(hipMemcpy(d_output, h_output, MAT_SIZE*MAT_SIZE*sizeof(int), hipMemcpyHostToDevice));

/**********************************************************************************************/
    /* 1. Transpose kerene launch using single thread
          A thread block of single thread is configured to perform matrix transpose
          This is exactly like sequential version of matrix transpose as only one 
          thread is responsible to perform entire matrix transpose operation          
    */
     
  /* Declaring Time variables to measure GPU time*/
  float GPUelapsed = 0.0;  // To store final kernel execution time 
  hipEvent_t start, stop; // Variables to record start and stop of kernel
  hipEventCreate(&start); // Event create start
  hipEventCreate(&stop);  // Event create stop
  hipEventRecord(start, 0); // Record time at start variables

/**/transposeKernel_single_thread<<<1,1>>>(d_input,d_output,MAT_SIZE); // krenel for matrix transpose

  hipEventRecord(stop, 0);
  hipEventSynchronize(stop);
  hipEventElapsedTime(&GPUelapsed, start, stop);
  hipEventDestroy(start); // Event destroy start
  hipEventDestroy(stop);  // Event destroy stop
  printf("\n 1. Parallel runtime Single thread %0.3f ms",GPUelapsed);
  printf("\n    Parallel Speedup Single thread = (CPU time)/(GPU time) = %0.3f\n", (CPUtime_taken*1000)/GPUelapsed);
    
  CHECK(hipMemcpy(hg_output, d_output, MAT_SIZE*MAT_SIZE*sizeof(int), hipMemcpyDeviceToHost)); // copy results back on host

    int check = checkResult(hg_output, h_output, MAT_SIZE); // This function checks two matrices are same or not
    if(check == 0)printf("\n Something went wrong => Results are not matching for 'transposeKernel_single_thread' !!!");
    initMatrix(hg_output,MAT_SIZE,0); // Initialize hg_output matrix to 0
    init_kernel<<<1,1>>>(d_output, MAT_SIZE); // Initialize d_output matrix to 0
/**********************************************************************************************/

    /* 2. Transpose kerene launch using thread per row
          In this configuration, # threads launched are equalt to # rows of matrix i.e MAT_SIZE
          Where each thread is responsible to perform transpose operation.
          Each thread reads the matrix entries row major order from d_input matrix and stores in 
          d_output matrix in column major order.    

    # threads launched = # rows in matrix
    numT = # threads per thread block
    numB = # thread blocks in the grid
  */

   int numT, numB;
   numT = 128;
   numB = ceil(MAT_SIZE/(float)numT);
  /* Declaring Time variables to measure GPU time*/
  float GPUelapsed2 = 0.0;  // To store final kernel execution time 
  hipEvent_t start2, stop2; // Variables to record start and stop of kernel
  hipEventCreate(&start2); // Event create start
  hipEventCreate(&stop2);  // Event create stop
  hipEventRecord(start2, 0); // Record time at start variables

/**/ transposeKernel_thread_per_row<<<numB,numT>>>(d_input,d_output,MAT_SIZE); // krenel for matrix transpose

  hipEventRecord(stop2, 0);
  hipEventSynchronize(stop2);
  hipEventElapsedTime(&GPUelapsed2, start2, stop2);
  hipEventDestroy(start2); // Event destroy start
  hipEventDestroy(stop2);  // Event destroy stop
  printf("\n 2. Parallel runtime thread per row is %0.3f ms",GPUelapsed2);
  printf("\n    Parallel Speedup thread per row = (CPU time)/(GPU time) = %0.3f\n", (CPUtime_taken*1000)/GPUelapsed2);
    
  CHECK(hipMemcpy(hg_output, d_output, MAT_SIZE*MAT_SIZE*sizeof(int), hipMemcpyDeviceToHost)); // copy results back on host

    check = checkResult(hg_output, h_output, MAT_SIZE); // This function checks two matrices are same or not
    if(check == 0)printf("\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_row'!!!");
    initMatrix(hg_output,MAT_SIZE,0); // Initialize hg_output matrix to 0
    init_kernel<<<1,1>>>(d_output, MAT_SIZE); // Initialize d_output matrix to 0


/**********************************************************************************************/
    /* 3. Transpose kerene launch using thread per element of the matrix
          In this configuration, # threads launched are equalt to # of elements in the matrix i.e MAT_SIZE x MAT_SIZE
          Where each thread is responsible to perform transpose operation on a single element.
          Each thread reads a single corresponding matrix element from d_input matrix and stores in 
          d_output matrix in transposed indices.    

    # threads launched = # elements in matrix
    numT = # threads per thread block in 2D
    numB = # thread blocks in the grid in 2D
  */

   dim3 num2T(8, 8,1);
   dim3 num2B(ceil(MAT_SIZE/(float)num2T.x), ceil(MAT_SIZE/(float)num2T.y ),1 );
   //printf("\n num2B.x = %d num2B.y = %d ", num2B.x,num2B.x );
  /* Declaring Time variables to measure GPU time */
  float GPUelapsed3 = 0.0;  // To store final kernel execution time 
  hipEvent_t start3, stop3; // Variables to record start and stop of kernel
  hipEventCreate(&start3); // Event create start
  hipEventCreate(&stop3);  // Event create stop
  hipEventRecord(start3, 0); // Record time at start variables

/**/ transposeKernel_thread_per_element<<<num2B,num2T>>>(d_input,d_output,MAT_SIZE); // krenel for matrix transpose

  hipEventRecord(stop3, 0);
  hipEventSynchronize(stop3);
  hipEventElapsedTime(&GPUelapsed3, start3, stop3);
  hipEventDestroy(start3); // Event destroy start
  hipEventDestroy(stop3);  // Event destroy stop
  printf("\n 3. Parallel runtime thread per element is %0.3f ms",GPUelapsed3);
  printf("\n    Parallel Speedup thread per element = (CPU time)/(GPU time) = %0.3f\n", (CPUtime_taken*1000)/GPUelapsed3);
    
  CHECK(hipMemcpy(hg_output, d_output, MAT_SIZE*MAT_SIZE*sizeof(int), hipMemcpyDeviceToHost)); // copy results back on host

    check = checkResult(hg_output, h_output, MAT_SIZE); // This function checks two matrices are same or not
    if(check == 0)
   {
    printf("\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_element'!!!");
   }

   free(h_input);
   free(h_output);
   free(hg_output);
   hipFree(d_input);
   hipFree(d_output);
   return(0);    
  }

	.text
	.file	"13_MatrixTranspose.hip"
	.globl	_Z44__device_stub__transposeKernel_single_threadPiS_i # -- Begin function _Z44__device_stub__transposeKernel_single_threadPiS_i
	.type	_Z44__device_stub__transposeKernel_single_threadPiS_i,@function
_Z44__device_stub__transposeKernel_single_threadPiS_i: # @_Z44__device_stub__transposeKernel_single_threadPiS_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z29transposeKernel_single_threadPiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.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
	retq
.Lfunc_end0:
	.size	_Z44__device_stub__transposeKernel_single_threadPiS_i, .Lfunc_end0-_Z44__device_stub__transposeKernel_single_threadPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z45__device_stub__transposeKernel_thread_per_rowPiS_i # -- Begin function _Z45__device_stub__transposeKernel_thread_per_rowPiS_i
	.type	_Z45__device_stub__transposeKernel_thread_per_rowPiS_i,@function
_Z45__device_stub__transposeKernel_thread_per_rowPiS_i: # @_Z45__device_stub__transposeKernel_thread_per_rowPiS_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z30transposeKernel_thread_per_rowPiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.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
	retq
.Lfunc_end1:
	.size	_Z45__device_stub__transposeKernel_thread_per_rowPiS_i, .Lfunc_end1-_Z45__device_stub__transposeKernel_thread_per_rowPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z49__device_stub__transposeKernel_thread_per_elementPiS_i # -- Begin function _Z49__device_stub__transposeKernel_thread_per_elementPiS_i
	.type	_Z49__device_stub__transposeKernel_thread_per_elementPiS_i,@function
_Z49__device_stub__transposeKernel_thread_per_elementPiS_i: # @_Z49__device_stub__transposeKernel_thread_per_elementPiS_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z34transposeKernel_thread_per_elementPiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.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
	retq
.Lfunc_end2:
	.size	_Z49__device_stub__transposeKernel_thread_per_elementPiS_i, .Lfunc_end2-_Z49__device_stub__transposeKernel_thread_per_elementPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z10initMatrixPiii              # -- Begin function _Z10initMatrixPiii
	.type	_Z10initMatrixPiii,@function
_Z10initMatrixPiii:                     # @_Z10initMatrixPiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebx
	movq	%rdi, %r14
	testl	%edx, %edx
	je	.LBB3_6
# %bb.1:                                # %.preheader26
	testl	%ebx, %ebx
	jle	.LBB3_9
# %bb.2:                                # %.preheader25.lr.ph
	movl	%ebx, %eax
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB3_3:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_4 Depth 2
	movl	%ecx, %esi
	leaq	(%r14,%rsi,4), %rsi
	xorl	%edi, %edi
.LBB3_4:                                #   Parent Loop BB3_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%edx, (%rsi,%rdi,4)
	incq	%rdi
	cmpq	%rdi, %rax
	jne	.LBB3_4
# %bb.5:                                # %._crit_edge
                                        #   in Loop: Header=BB3_3 Depth=1
	incq	%rdx
	addl	%ebx, %ecx
	cmpq	%rax, %rdx
	jne	.LBB3_3
	jmp	.LBB3_9
.LBB3_6:                                # %.preheader24
	testl	%ebx, %ebx
	jle	.LBB3_9
# %bb.7:                                # %.preheader.lr.ph
	movl	%ebx, %r12d
	leaq	(,%r12,4), %r15
	xorl	%ebp, %ebp
.LBB3_8:                                # %._crit_edge32
                                        # =>This Inner Loop Header: Depth=1
	movl	%ebp, %eax
	leaq	(%r14,%rax,4), %rdi
	xorl	%esi, %esi
	movq	%r15, %rdx
	callq	memset@PLT
	addl	%ebx, %ebp
	decq	%r12
	jne	.LBB3_8
.LBB3_9:                                # %.loopexit
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z10initMatrixPiii, .Lfunc_end3-_Z10initMatrixPiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z11checkResultPiS_i            # -- Begin function _Z11checkResultPiS_i
	.type	_Z11checkResultPiS_i,@function
_Z11checkResultPiS_i:                   # @_Z11checkResultPiS_i
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB4_1
# %bb.3:                                # %.preheader.lr.ph
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%edx, %ecx
	movl	$1, %eax
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
.LBB4_4:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_5 Depth 2
	movl	%r9d, %ebx
	leaq	(%rsi,%rbx,4), %r11
	leaq	(%rdi,%rbx,4), %rbx
	xorl	%r14d, %r14d
.LBB4_5:                                #   Parent Loop BB4_4 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rbx,%r14,4), %ebp
	cmpl	(%r11,%r14,4), %ebp
	cmovnel	%r8d, %eax
	incq	%r14
	cmpq	%r14, %rcx
	jne	.LBB4_5
# %bb.6:                                # %._crit_edge
                                        #   in Loop: Header=BB4_4 Depth=1
	incq	%r10
	addl	%edx, %r9d
	cmpq	%rcx, %r10
	jne	.LBB4_4
# %bb.7:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %rbp
	retq
.LBB4_1:
	movl	$1, %eax
	retq
.Lfunc_end4:
	.size	_Z11checkResultPiS_i, .Lfunc_end4-_Z11checkResultPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__init_kernelPii # -- Begin function _Z26__device_stub__init_kernelPii
	.type	_Z26__device_stub__init_kernelPii,@function
_Z26__device_stub__init_kernelPii:      # @_Z26__device_stub__init_kernelPii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z11init_kernelPii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.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
	retq
.Lfunc_end5:
	.size	_Z26__device_stub__init_kernelPii, .Lfunc_end5-_Z26__device_stub__init_kernelPii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI6_0:
	.quad	0x412e848000000000              # double 1.0E+6
.LCPI6_1:
	.quad	0x408f400000000000              # double 1000
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$16777216, %edi                 # imm = 0x1000000
	callq	malloc
	movq	%rax, %r13
	movl	$16777216, %edi                 # imm = 0x1000000
	callq	malloc
	movq	%rax, %r14
	movl	$16777216, %edi                 # imm = 0x1000000
	callq	malloc
	movq	%rax, %r15
	testq	%r13, %r13
	jne	.LBB6_2
# %bb.1:
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
.LBB6_2:
	testq	%r14, %r14
	jne	.LBB6_4
# %bb.3:
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
.LBB6_4:
	testq	%r15, %r15
	jne	.LBB6_6
# %bb.5:
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
.LBB6_6:                                # %._crit_edge32.i148.preheader
	xorl	%ebx, %ebx
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movq	%r13, %rax
.LBB6_7:                                # %.lr.ph.i
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_8 Depth 2
	xorl	%ecx, %ecx
.LBB6_8:                                #   Parent Loop BB6_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%ebx, (%rax,%rcx,4)
	incq	%rcx
	cmpq	$2048, %rcx                     # imm = 0x800
	jne	.LBB6_8
# %bb.9:                                # %._crit_edge.i
                                        #   in Loop: Header=BB6_7 Depth=1
	incq	%rbx
	addq	$8192, %rax                     # imm = 0x2000
	cmpq	$2048, %rbx                     # imm = 0x800
	jne	.LBB6_7
# %bb.10:                               # %_Z10initMatrixPiii.exit154
	xorl	%ebx, %ebx
	callq	clock
	movq	%rax, %r12
	movq	%r13, %rax
	movq	%r14, %rcx
.LBB6_11:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_12 Depth 2
	movq	%rax, %rdx
	xorl	%esi, %esi
.LBB6_12:                               #   Parent Loop BB6_11 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rdx), %edi
	movl	%edi, (%rcx,%rsi,4)
	incq	%rsi
	addq	$8192, %rdx                     # imm = 0x2000
	cmpq	$2048, %rsi                     # imm = 0x800
	jne	.LBB6_12
# %bb.13:                               #   in Loop: Header=BB6_11 Depth=1
	incq	%rbx
	addq	$8192, %rcx                     # imm = 0x2000
	addq	$4, %rax
	cmpq	$2048, %rbx                     # imm = 0x800
	jne	.LBB6_11
# %bb.14:
	callq	clock
	subq	%r12, %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI6_0(%rip), %xmm0
	mulsd	.LCPI6_1(%rip), %xmm0
	movl	$.L.str.2, %edi
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	movb	$1, %al
	callq	printf
	leaq	16(%rsp), %rdi
	movl	$16777216, %esi                 # imm = 0x1000000
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB6_15
# %bb.17:
	leaq	8(%rsp), %rdi
	movl	$16777216, %esi                 # imm = 0x1000000
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB6_18
# %bb.19:
	movq	16(%rsp), %rdi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB6_20
# %bb.21:
	movq	%r13, 96(%rsp)                  # 8-byte Spill
	movq	8(%rsp), %rdi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB6_22
# %bb.23:
	movabsq	$4294967297, %r12               # imm = 0x100000001
	movl	$0, 68(%rsp)
	leaq	88(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventCreate
	leaq	40(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_25
# %bb.24:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$2048, %edx                     # imm = 0x800
	callq	_Z44__device_stub__transposeKernel_single_threadPiS_i
.LBB6_25:
	movq	40(%rsp), %rdi
	xorl	%ebx, %ebx
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	40(%rsp), %rdi
	callq	hipEventSynchronize
	movq	88(%rsp), %rsi
	movq	40(%rsp), %rdx
	leaq	68(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventElapsedTime
	movq	88(%rsp), %rdi
	callq	hipEventDestroy
	movq	40(%rsp), %rdi
	callq	hipEventDestroy
	cvtss2sd	(%r13), %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	cvtss2sd	(%r13), %xmm1
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	divsd	%xmm1, %xmm0
	movl	$.L.str.7, %edi
	movb	$1, %al
	callq	printf
	movq	8(%rsp), %rsi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB6_55
# %bb.26:                               # %.lr.ph.i155.preheader
	movl	$1, %eax
	movl	$8192, %ecx                     # imm = 0x2000
	movq	%r15, %rdx
	movq	%r14, %rsi
	xorl	%edi, %edi
.LBB6_27:                               # %.lr.ph.i155
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_28 Depth 2
	xorl	%r8d, %r8d
.LBB6_28:                               #   Parent Loop BB6_27 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rdx,%r8,4), %r9d
	cmpl	(%rsi,%r8,4), %r9d
	cmovnel	%ebx, %eax
	incq	%r8
	cmpq	$2048, %r8                      # imm = 0x800
	jne	.LBB6_28
# %bb.29:                               # %._crit_edge.i159
                                        #   in Loop: Header=BB6_27 Depth=1
	incq	%rdi
	addq	%rcx, %rsi
	addq	%rcx, %rdx
	cmpq	$2048, %rdi                     # imm = 0x800
	jne	.LBB6_27
# %bb.30:                               # %_Z11checkResultPiS_i.exit
	testl	%eax, %eax
	jne	.LBB6_32
# %bb.31:
	movl	$.L.str.8, %edi
	xorl	%eax, %eax
	callq	printf
.LBB6_32:                               # %_Z10initMatrixPiii.exit165
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_34
# %bb.33:
	movq	8(%rsp), %rdi
	movl	$2048, %esi                     # imm = 0x800
	callq	_Z26__device_stub__init_kernelPii
.LBB6_34:
	movl	$0, 64(%rsp)
	leaq	80(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventCreate
	leaq	32(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	leaq	15(%r12), %rdi
	leaq	127(%r12), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_36
# %bb.35:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$2048, %edx                     # imm = 0x800
	callq	_Z45__device_stub__transposeKernel_thread_per_rowPiS_i
.LBB6_36:
	movq	32(%rsp), %rdi
	xorl	%ebx, %ebx
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	32(%rsp), %rdi
	callq	hipEventSynchronize
	movq	80(%rsp), %rsi
	movq	32(%rsp), %rdx
	leaq	64(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventElapsedTime
	movq	80(%rsp), %rdi
	callq	hipEventDestroy
	movq	32(%rsp), %rdi
	callq	hipEventDestroy
	cvtss2sd	(%r13), %xmm0
	movl	$.L.str.9, %edi
	movb	$1, %al
	callq	printf
	cvtss2sd	(%r13), %xmm1
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	divsd	%xmm1, %xmm0
	movl	$.L.str.10, %edi
	movb	$1, %al
	callq	printf
	movq	8(%rsp), %rsi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB6_56
# %bb.37:                               # %.lr.ph.i166.preheader
	movl	$1, %eax
	movl	$8192, %ecx                     # imm = 0x2000
	movq	%r15, %rdx
	movq	%r14, %rsi
	xorl	%edi, %edi
	movq	96(%rsp), %r13                  # 8-byte Reload
.LBB6_38:                               # %.lr.ph.i166
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_39 Depth 2
	xorl	%r8d, %r8d
.LBB6_39:                               #   Parent Loop BB6_38 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rdx,%r8,4), %r9d
	cmpl	(%rsi,%r8,4), %r9d
	cmovnel	%ebx, %eax
	incq	%r8
	cmpq	$2048, %r8                      # imm = 0x800
	jne	.LBB6_39
# %bb.40:                               # %._crit_edge.i175
                                        #   in Loop: Header=BB6_38 Depth=1
	incq	%rdi
	addq	%rcx, %rsi
	addq	%rcx, %rdx
	cmpq	$2048, %rdi                     # imm = 0x800
	jne	.LBB6_38
# %bb.41:                               # %_Z11checkResultPiS_i.exit178
	testl	%eax, %eax
	jne	.LBB6_43
# %bb.42:
	movl	$.L.str.11, %edi
	xorl	%eax, %eax
	callq	printf
.LBB6_43:                               # %_Z10initMatrixPiii.exit184
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_45
# %bb.44:
	movq	8(%rsp), %rdi
	movl	$2048, %esi                     # imm = 0x800
	callq	_Z26__device_stub__init_kernelPii
.LBB6_45:
	movl	$0, 60(%rsp)
	leaq	72(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventCreate
	leaq	24(%rsp), %rdi
	callq	hipEventCreate
	movq	(%r12), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movabsq	$1099511628032, %rdi            # imm = 0x10000000100
	movabsq	$34359738376, %rdx              # imm = 0x800000008
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_47
# %bb.46:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$2048, %edx                     # imm = 0x800
	callq	_Z49__device_stub__transposeKernel_thread_per_elementPiS_i
.LBB6_47:
	movq	24(%rsp), %rdi
	xorl	%ebx, %ebx
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	24(%rsp), %rdi
	callq	hipEventSynchronize
	movq	72(%rsp), %rsi
	movq	24(%rsp), %rdx
	leaq	60(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventElapsedTime
	movq	72(%rsp), %rdi
	callq	hipEventDestroy
	movq	24(%rsp), %rdi
	callq	hipEventDestroy
	cvtss2sd	(%r12), %xmm0
	movl	$.L.str.12, %edi
	movb	$1, %al
	callq	printf
	cvtss2sd	(%r12), %xmm1
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	divsd	%xmm1, %xmm0
	movl	$.L.str.13, %edi
	movb	$1, %al
	callq	printf
	movq	8(%rsp), %rsi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB6_57
# %bb.48:                               # %.lr.ph.i185.preheader
	movl	$1, %eax
	movl	$8192, %ecx                     # imm = 0x2000
	movq	%r15, %rdx
	movq	%r14, %rsi
	xorl	%edi, %edi
.LBB6_49:                               # %.lr.ph.i185
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_50 Depth 2
	xorl	%r8d, %r8d
.LBB6_50:                               #   Parent Loop BB6_49 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rdx,%r8,4), %r9d
	cmpl	(%rsi,%r8,4), %r9d
	cmovnel	%ebx, %eax
	incq	%r8
	cmpq	$2048, %r8                      # imm = 0x800
	jne	.LBB6_50
# %bb.51:                               # %._crit_edge.i194
                                        #   in Loop: Header=BB6_49 Depth=1
	incq	%rdi
	addq	%rcx, %rsi
	addq	%rcx, %rdx
	cmpq	$2048, %rdi                     # imm = 0x800
	jne	.LBB6_49
# %bb.52:                               # %_Z11checkResultPiS_i.exit197
	testl	%eax, %eax
	jne	.LBB6_54
# %bb.53:
	movl	$.L.str.14, %edi
	xorl	%eax, %eax
	callq	printf
.LBB6_54:
	movq	%r13, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB6_15:
	.cfi_def_cfa_offset 160
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$193, %edx
	jmp	.LBB6_16
.LBB6_18:
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$194, %edx
	jmp	.LBB6_16
.LBB6_20:
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$198, %edx
	jmp	.LBB6_16
.LBB6_22:
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$199, %edx
	jmp	.LBB6_16
.LBB6_55:
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$225, %edx
	jmp	.LBB6_16
.LBB6_56:
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$264, %edx                      # imm = 0x108
	jmp	.LBB6_16
.LBB6_57:
	movl	%eax, %ebp
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	movl	$304, %edx                      # imm = 0x130
.LBB6_16:
	xorl	%eax, %eax
	callq	printf
	movl	%ebp, %edi
	callq	hipGetErrorString
	movl	$.L.str.5, %edi
	movl	%ebp, %esi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	exit
.Lfunc_end6:
	.size	main, .Lfunc_end6-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z29transposeKernel_single_threadPiS_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z30transposeKernel_thread_per_rowPiS_i, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z34transposeKernel_thread_per_elementPiS_i, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11init_kernelPii, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB8_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z29transposeKernel_single_threadPiS_i,@object # @_Z29transposeKernel_single_threadPiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z29transposeKernel_single_threadPiS_i
	.p2align	3, 0x0
_Z29transposeKernel_single_threadPiS_i:
	.quad	_Z44__device_stub__transposeKernel_single_threadPiS_i
	.size	_Z29transposeKernel_single_threadPiS_i, 8

	.type	_Z30transposeKernel_thread_per_rowPiS_i,@object # @_Z30transposeKernel_thread_per_rowPiS_i
	.globl	_Z30transposeKernel_thread_per_rowPiS_i
	.p2align	3, 0x0
_Z30transposeKernel_thread_per_rowPiS_i:
	.quad	_Z45__device_stub__transposeKernel_thread_per_rowPiS_i
	.size	_Z30transposeKernel_thread_per_rowPiS_i, 8

	.type	_Z34transposeKernel_thread_per_elementPiS_i,@object # @_Z34transposeKernel_thread_per_elementPiS_i
	.globl	_Z34transposeKernel_thread_per_elementPiS_i
	.p2align	3, 0x0
_Z34transposeKernel_thread_per_elementPiS_i:
	.quad	_Z49__device_stub__transposeKernel_thread_per_elementPiS_i
	.size	_Z34transposeKernel_thread_per_elementPiS_i, 8

	.type	_Z11init_kernelPii,@object      # @_Z11init_kernelPii
	.globl	_Z11init_kernelPii
	.p2align	3, 0x0
_Z11init_kernelPii:
	.quad	_Z26__device_stub__init_kernelPii
	.size	_Z11init_kernelPii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"\n Error: Allocating memory to h_input!!!"
	.size	.L.str, 41

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"\n Error: Allocating memory to h_output!!!"
	.size	.L.str.1, 42

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\n     Time required for sequential execution is %.3f ms\n"
	.size	.L.str.2, 57

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Error %s %d"
	.size	.L.str.3, 12

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/NileshchandraPikle/Introduction-to-CUDA-Programming/master/13_MatrixTranspose.hip"
	.size	.L.str.4, 144

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\n Code %d Reason %s \n"
	.size	.L.str.5, 22

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"\n 1. Parallel runtime Single thread %0.3f ms"
	.size	.L.str.6, 45

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"\n    Parallel Speedup Single thread = (CPU time)/(GPU time) = %0.3f\n"
	.size	.L.str.7, 69

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"\n Something went wrong => Results are not matching for 'transposeKernel_single_thread' !!!"
	.size	.L.str.8, 91

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"\n 2. Parallel runtime thread per row is %0.3f ms"
	.size	.L.str.9, 49

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"\n    Parallel Speedup thread per row = (CPU time)/(GPU time) = %0.3f\n"
	.size	.L.str.10, 70

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_row'!!!"
	.size	.L.str.11, 91

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"\n 3. Parallel runtime thread per element is %0.3f ms"
	.size	.L.str.12, 53

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"\n    Parallel Speedup thread per element = (CPU time)/(GPU time) = %0.3f\n"
	.size	.L.str.13, 74

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"\n Something went wrong => Results are not matching for 'transposeKernel_thread_per_element'!!!"
	.size	.L.str.14, 95

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z29transposeKernel_single_threadPiS_i"
	.size	.L__unnamed_1, 39

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z30transposeKernel_thread_per_rowPiS_i"
	.size	.L__unnamed_2, 40

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z34transposeKernel_thread_per_elementPiS_i"
	.size	.L__unnamed_3, 44

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z11init_kernelPii"
	.size	.L__unnamed_4, 19

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.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 _Z44__device_stub__transposeKernel_single_threadPiS_i
	.addrsig_sym _Z45__device_stub__transposeKernel_thread_per_rowPiS_i
	.addrsig_sym _Z49__device_stub__transposeKernel_thread_per_elementPiS_i
	.addrsig_sym _Z26__device_stub__init_kernelPii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z29transposeKernel_single_threadPiS_i
	.addrsig_sym _Z30transposeKernel_thread_per_rowPiS_i
	.addrsig_sym _Z34transposeKernel_thread_per_elementPiS_i
	.addrsig_sym _Z11init_kernelPii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z29transposeKernel_single_threadPiS_i ; -- Begin function _Z29transposeKernel_single_threadPiS_i
	.globl	_Z29transposeKernel_single_threadPiS_i
	.p2align	8
	.type	_Z29transposeKernel_single_threadPiS_i,@function
_Z29transposeKernel_single_threadPiS_i: ; @_Z29transposeKernel_single_threadPiS_i
; %bb.0:
	s_load_b32 s10, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s10, 1
	s_cbranch_scc1 .LBB0_5
; %bb.1:                                ; %.preheader.preheader
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mov_b32_e32 v0, 0
	s_mov_b32 s5, 0
	s_mov_b32 s11, 0
	s_mov_b32 s6, s5
.LBB0_2:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_3 Depth 2
	s_mov_b32 s7, s5
	s_mov_b32 s4, s11
	s_lshl_b64 s[8:9], s[6:7], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s8, s0, s8
	s_addc_u32 s9, s1, s9
.LBB0_3:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v1, v0, s[8:9]
	s_lshl_b64 s[12:13], s[4:5], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s12, s2, s12
	s_addc_u32 s13, s3, s13
	s_add_i32 s7, s7, 1
	s_add_i32 s4, s4, s10
	s_add_u32 s8, s8, 4
	s_addc_u32 s9, s9, 0
	s_cmp_lg_u32 s10, s7
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[12:13]
	s_cbranch_scc1 .LBB0_3
; %bb.4:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_add_i32 s11, s11, 1
	s_add_i32 s6, s6, s10
	s_cmp_lg_u32 s11, s10
	s_cbranch_scc1 .LBB0_2
.LBB0_5:                                ; %._crit_edge19
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z29transposeKernel_single_threadPiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 2
		.amdhsa_next_free_sgpr 14
		.amdhsa_reserve_vcc 0
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z29transposeKernel_single_threadPiS_i, .Lfunc_end0-_Z29transposeKernel_single_threadPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 156
; NumSgprs: 14
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 14
; NumVGPRsForWavesPerEU: 2
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z30transposeKernel_thread_per_rowPiS_i ; -- Begin function _Z30transposeKernel_thread_per_rowPiS_i
	.globl	_Z30transposeKernel_thread_per_rowPiS_i
	.p2align	8
	.type	_Z30transposeKernel_thread_per_rowPiS_i,@function
_Z30transposeKernel_thread_per_rowPiS_i: ; @_Z30transposeKernel_thread_per_rowPiS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b32 s4, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_max_i32_e32 v0, 0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s4, v0
	s_cbranch_execz .LBB1_3
; %bb.1:                                ; %.lr.ph
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mul_lo_u32 v2, v1, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, s2, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v4, vcc_lo
	s_mov_b32 s2, s4
.LBB1_2:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v2, 31, v1
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s2, 0
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	v_add_nc_u32_e32 v1, s4, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v5, vcc_lo, s0, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s1, v6, vcc_lo
	global_load_b32 v0, v[5:6], off
	s_waitcnt vmcnt(0)
	global_store_b32 v[3:4], v0, off
	v_add_co_u32 v3, vcc_lo, v3, 4
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
	s_cbranch_scc1 .LBB1_2
.LBB1_3:                                ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z30transposeKernel_thread_per_rowPiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 7
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z30transposeKernel_thread_per_rowPiS_i, .Lfunc_end1-_Z30transposeKernel_thread_per_rowPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 216
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z34transposeKernel_thread_per_elementPiS_i ; -- Begin function _Z34transposeKernel_thread_per_elementPiS_i
	.globl	_Z34transposeKernel_thread_per_elementPiS_i
	.p2align	8
	.type	_Z34transposeKernel_thread_per_elementPiS_i,@function
_Z34transposeKernel_thread_per_elementPiS_i: ; @_Z34transposeKernel_thread_per_elementPiS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x24
	s_load_b32 s2, s[0:1], 0x10
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[0:1], null, s14, s3, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, s4, v[3:4]
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_i32_e32 v2, v0, v1
	v_cmpx_gt_i32_e64 s2, v2
	s_cbranch_execz .LBB2_2
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x0
	v_mad_u64_u32 v[2:3], null, v0, s2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v3, vcc_lo
	global_load_b32 v4, v[2:3], off
	v_mad_u64_u32 v[2:3], null, v1, s2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v4, off
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z34transposeKernel_thread_per_elementPiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z34transposeKernel_thread_per_elementPiS_i, .Lfunc_end2-_Z34transposeKernel_thread_per_elementPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 220
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z11init_kernelPii      ; -- Begin function _Z11init_kernelPii
	.globl	_Z11init_kernelPii
	.p2align	8
	.type	_Z11init_kernelPii,@function
_Z11init_kernelPii:                     ; @_Z11init_kernelPii
; %bb.0:
	s_load_b32 s6, s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s6, 1
	s_cbranch_scc1 .LBB3_5
; %bb.1:                                ; %.preheader.preheader
	s_load_b64 s[0:1], s[0:1], 0x0
	v_mov_b32_e32 v0, 0
	s_mov_b32 s3, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s2, s3
	s_mov_b32 s7, s3
.LBB3_2:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB3_3 Depth 2
	s_lshl_b64 s[4:5], s[2:3], 2
	s_mov_b32 s8, s6
	s_waitcnt lgkmcnt(0)
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
.LBB3_3:                                ;   Parent Loop BB3_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_add_i32 s8, s8, -1
	global_store_b32 v0, v0, s[4:5]
	s_add_u32 s4, s4, 4
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s8, 0
	s_cbranch_scc1 .LBB3_3
; %bb.4:                                ; %._crit_edge
                                        ;   in Loop: Header=BB3_2 Depth=1
	s_add_i32 s7, s7, 1
	s_add_i32 s2, s2, s6
	s_cmp_lg_u32 s7, s6
	s_cbranch_scc1 .LBB3_2
.LBB3_5:                                ; %._crit_edge14
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11init_kernelPii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 12
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 1
		.amdhsa_next_free_sgpr 9
		.amdhsa_reserve_vcc 0
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end3:
	.size	_Z11init_kernelPii, .Lfunc_end3-_Z11init_kernelPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 124
; NumSgprs: 9
; NumVgprs: 1
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 9
; NumVGPRsForWavesPerEU: 1
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.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_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z29transposeKernel_single_threadPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     14
    .sgpr_spill_count: 0
    .symbol:         _Z29transposeKernel_single_threadPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z30transposeKernel_thread_per_rowPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z30transposeKernel_thread_per_rowPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z34transposeKernel_thread_per_elementPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z34transposeKernel_thread_per_elementPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 12
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11init_kernelPii
    .private_segment_fixed_size: 0
    .sgpr_count:     9
    .sgpr_spill_count: 0
    .symbol:         _Z11init_kernelPii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     1
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
