#include <stdio.h>


__global__ void kicache_test4_2 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_4 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_6 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_8 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_10 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_12 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_14 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_16 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_18 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_20 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_22 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_24 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_26 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_28 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_30 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_32 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_34 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_36 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_38 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_40 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_42 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_44 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_46 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_48 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_50 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_52 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_54 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_56 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_58 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_60 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_62 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_64 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_72 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_80 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_88 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_96 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_104 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_112 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_120 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_128 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_136 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_144 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_152 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_160 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_168 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_176 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_184 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_192 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_200 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_208 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_216 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_224 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_232 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
//__global__ void kicache_test4_240 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
//__global__ void kicache_test4_248 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
//__global__ void kicache_test4_256 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);




void measure_icache()
{
	
	unsigned int ts[1024];			// ts, output from kernel. Two elements used per thread.
	unsigned int *d_ts;
	unsigned int *d_out;			// Unused memory for storing output
	

    	dim3 Db = dim3(1);
    	dim3 Dg = dim3(1,1,1);
	
	// Allocate device array.
	if (cudaSuccess != cudaMalloc((void**)&d_ts, sizeof(ts)))
	{
		printf ("cudaMalloc failed %s:%d\n", __FILE__, __LINE__);
		return;
	}
	if (cudaSuccess != cudaMalloc((void**)&d_out, 4))
	{
		printf ("cudaMalloc failed %s:%d\n", __FILE__, __LINE__);
		return;
	}
	
    fprintf (stderr, "Running icache test...\n");
	
	

	// Measure instruction cache size
	printf ("Instruction cache size:\n");
	



	unsigned int sum_times[32];
	
	
	printf ("  0.5 KB steps: ");
	Db.x = 1;
	for (int p2 = 1; p2 <= 32; p2++)
	{
		unsigned int sum_time = 0;
		bool failed = false;
		for (int i=0;i<100 && !failed ;i++)
		{
			cudaGetLastError();		// Clear previous error code, if any
			switch (p2) {
				case 1: kicache_test4_2 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 2: kicache_test4_4 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 3: kicache_test4_6 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 4: kicache_test4_8 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 5: kicache_test4_10 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 6: kicache_test4_12 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 7: kicache_test4_14 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 8: kicache_test4_16 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 9: kicache_test4_18 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 10: kicache_test4_20 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 11: kicache_test4_22 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 12: kicache_test4_24 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 13: kicache_test4_26 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 14: kicache_test4_28 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 15: kicache_test4_30 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 16: kicache_test4_32 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 17: kicache_test4_34 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 18: kicache_test4_36 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 19: kicache_test4_38 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 20: kicache_test4_40 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 21: kicache_test4_42 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 22: kicache_test4_44 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 23: kicache_test4_46 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 24: kicache_test4_48 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 25: kicache_test4_50 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 26: kicache_test4_52 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 27: kicache_test4_54 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 28: kicache_test4_56 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 29: kicache_test4_58 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 30: kicache_test4_60 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 31: kicache_test4_62 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 32: kicache_test4_64 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
			}
				
			if (cudaGetLastError() != cudaSuccess)
			{
				failed = true;
				break;
			}
				
			cudaThreadSynchronize();
			cudaMemcpy(ts, d_ts, sizeof(ts), cudaMemcpyDeviceToHost);
			
			sum_time += (ts[1] - ts[0]);
			
		}
		
		if (failed)
		{
			printf ("xxxx ");
		}
		else
		{
			// Compute average latency over the lifetime of each warp (sum_time), and average throughput of the kernel (sum_max_time).
			printf ("%.5f ", sum_time/100.0/(p2*64)); fflush(stdout);
		}
		
		sum_times[p2-1] = sum_time;
	}
	printf (" (icache = ");
	for (int last_i=1, i=1;i<32;i++)
	{
		//printf("sum time %d, %d \n",sum_times[i]/(i+1),sum_times[last_i]/(last_i+1)*1.33);
		if (sum_times[i]/(i+1) > sum_times[last_i]/(last_i+1) /**1.33*/)
		{
			printf ("%.1fKB ", i*0.5);
			last_i = i;
		}	
	}
	printf (")\n");
	
	
	printf ("  2 KB steps: ");
	Db.x = 1;
	for (int p2 = 1; p2 <= 29/*32*/; p2++)
	{
		unsigned int sum_time = 0;
		bool failed = false;
		for (int i=0;i<100 && !failed ;i++)
		{
			cudaGetLastError();		// Clear previous error code, if any
			switch (p2) {
				case 1: kicache_test4_8 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 2: kicache_test4_16 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 3: kicache_test4_24 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 4: kicache_test4_32 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 5: kicache_test4_40 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 6: kicache_test4_48 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 7: kicache_test4_56 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 8: kicache_test4_64 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 9: kicache_test4_72 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 10: kicache_test4_80 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 11: kicache_test4_88 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 12: kicache_test4_96 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 13: kicache_test4_104 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 14: kicache_test4_112 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 15: kicache_test4_120 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 16: kicache_test4_128 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 17: kicache_test4_136 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 18: kicache_test4_144 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 19: kicache_test4_152 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 20: kicache_test4_160 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 21: kicache_test4_168 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 22: kicache_test4_176 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 23: kicache_test4_184 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 24: kicache_test4_192 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 25: kicache_test4_200 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 26: kicache_test4_208 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 27: kicache_test4_216 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 28: kicache_test4_224 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 29: kicache_test4_232 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				//case 30: kicache_test4_240 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				//case 31: kicache_test4_248 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				//case 32: kicache_test4_256 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;

			}			

				
			if (cudaGetLastError() != cudaSuccess)
			{
				failed = true;
				break;
			}
				
			cudaThreadSynchronize();
			cudaMemcpy(ts, d_ts, sizeof(ts), cudaMemcpyDeviceToHost);
			
			sum_time += (ts[1] - ts[0]);
		}
		
		if (failed)
		{
			printf ("xxxx ");
		}
		else
		{
			// Compute average latency over the lifetime of each warp (sum_time), and average throughput of the kernel (sum_max_time).
			printf ("%.1f ", sum_time/100.0/(p2*256)); fflush(stdout);
		}
		sum_times[p2-1] = sum_time;
	}
	printf (" (icache = ");
	for (int last_i=1, i=1;i<28/*32*/;i++)
	{
		if (sum_times[i]/(i+1) > sum_times[last_i]/(last_i+1) /**1.33*/)
		{
			printf ("%.1fKB ", i*2.0);
			last_i = i;
		}	
	}
	printf (")\n");
	
	
	printf ("\n Test instruction cache sharing by running two thread blocks concurently.\n");
	
	Db.x = 1;
	dim3 Dg2 = dim3(31,1,1);
	
	for (int blk2 = 1; blk2 <31; blk2+= (blk2 == 1? 9 : 10))
	{
		printf ("  TPC 0,%d (2 KB steps): ", blk2);
		int mask = (1<<blk2) | 1;	// Enable two blocks for execution
		for (int p2 = 1; p2 <= 28/*32*/; p2++)
		{
			unsigned int sum_time = 0;
			bool failed = false;
			for (int i=0;i<50 && !failed ;i++)
			{
				cudaGetLastError();		// Clear previous error code, if any
				switch (p2) {
					case 1: kicache_test4_8 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 2: kicache_test4_16 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 3: kicache_test4_24 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 4: kicache_test4_32 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 5: kicache_test4_40 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 6: kicache_test4_48 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 7: kicache_test4_56 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 8: kicache_test4_64 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 9: kicache_test4_72 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 10: kicache_test4_80 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 11: kicache_test4_88 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 12: kicache_test4_96 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 13: kicache_test4_104 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 14: kicache_test4_112 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 15: kicache_test4_120 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 16: kicache_test4_128 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 17: kicache_test4_136 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 18: kicache_test4_144 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 19: kicache_test4_152 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 20: kicache_test4_160 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 21: kicache_test4_168 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 22: kicache_test4_176 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 23: kicache_test4_184 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 24: kicache_test4_192 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 25: kicache_test4_200 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 26: kicache_test4_208 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 27: kicache_test4_216 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 28: kicache_test4_224 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 29: kicache_test4_232 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 30: kicache_test4_240 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 31: kicache_test4_248 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 32: kicache_test4_256 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;

				}

				if (cudaGetLastError() != cudaSuccess)
				{
					failed = true;
					break;
				}
					
				cudaThreadSynchronize();
				cudaMemcpy(ts, d_ts, sizeof(ts), cudaMemcpyDeviceToHost);
				
				sum_time += (ts[1] - ts[0]);
			}
			
			if (failed)
			{
				printf ("xxxx ");
			}
			else
			{
				printf ("%.1f ", sum_time/50.0/(p2*256)); fflush(stdout);
			}
			sum_times[p2-1] = sum_time;
		}
		printf (" (apparent icache = ");
		for (int last_i=1, i=1;i<32;i++)
		{
			if (sum_times[i]/(i+1) > sum_times[last_i]/(last_i+1) *1.25)
			{
				printf ("%.1fKB ", i*2.0);
				last_i = i;
			}	
		}
		printf (")\n");	
	}
	
	
	printf ("\n");
	
	cudaFree(d_ts);
	cudaFree(d_out);

	
	
}

int main()
{
	measure_icache();
}




	.file	"tmpxft_002be6ae_00000000-6_icache.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/mirsoleimani/SAMMicrobenchmark/master/cudabmk/icache.cu"
.LC1:
	.string	"cudaMalloc failed %s:%d\n"
.LC2:
	.string	"Running icache test...\n"
.LC3:
	.string	"Instruction cache size:\n"
.LC4:
	.string	"  0.5 KB steps: "
.LC5:
	.string	" (icache = "
.LC7:
	.string	"%.1fKB "
.LC8:
	.string	")\n"
.LC9:
	.string	"  2 KB steps: "
.LC10:
	.string	"\n Test instruction cache sharing by running two thread blocks concurently.\n"
.LC11:
	.string	"  TPC 0,%d (2 KB steps): "
.LC12:
	.string	" (apparent icache = "
.LC14:
	.string	"\n"
.LC15:
	.string	"xxxx "
.LC17:
	.string	"%.1f "
.LC19:
	.string	"%.5f "
	.text
	.globl	_Z14measure_icachev
	.type	_Z14measure_icachev, @function
_Z14measure_icachev:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$4096, %rsp
	.cfi_def_cfa_offset 4152
	orq	$0, (%rsp)
	subq	$216, %rsp
	.cfi_def_cfa_offset 4368
	movl	$4096, %esi
	movq	%fs:40, %rax
	movq	%rax, 4296(%rsp)
	movabsq	$4294967297, %rax
	leaq	16(%rsp), %rdi
	movl	$1, 44(%rsp)
	movl	$1, 56(%rsp)
	movq	%rax, 36(%rsp)
	movq	%rax, 48(%rsp)
	call	cudaMalloc@PLT
	movl	$78, %ecx
	testl	%eax, %eax
	jne	.L157
	leaq	24(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L5
	movl	$83, %ecx
.L157:
	leaq	.LC0(%rip), %rdx
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L2
.L5:
	movq	stderr(%rip), %rdi
	leaq	.LC2(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	movl	$1, %ebp
	leaq	.L9(%rip), %r13
	call	__fprintf_chk@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L6:
	movl	%ebp, %r14d
	xorl	%ebx, %ebx
	movl	$100, %r12d
.L44:
	call	cudaGetLastError@PLT
	leal	-2(%r14), %eax
	movq	36(%rsp), %rdx
	movl	44(%rsp), %ecx
	cmpl	$30, %eax
	movq	48(%rsp), %rdi
	movl	56(%rsp), %esi
	ja	.L7
	movslq	0(%r13,%rax,4), %rax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addq	%r13, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L9:
	.long	.L39-.L9
	.long	.L38-.L9
	.long	.L37-.L9
	.long	.L36-.L9
	.long	.L35-.L9
	.long	.L34-.L9
	.long	.L33-.L9
	.long	.L32-.L9
	.long	.L31-.L9
	.long	.L30-.L9
	.long	.L29-.L9
	.long	.L28-.L9
	.long	.L27-.L9
	.long	.L26-.L9
	.long	.L25-.L9
	.long	.L24-.L9
	.long	.L23-.L9
	.long	.L22-.L9
	.long	.L21-.L9
	.long	.L20-.L9
	.long	.L19-.L9
	.long	.L18-.L9
	.long	.L17-.L9
	.long	.L16-.L9
	.long	.L15-.L9
	.long	.L14-.L9
	.long	.L13-.L9
	.long	.L12-.L9
	.long	.L11-.L9
	.long	.L10-.L9
	.long	.L8-.L9
	.text
.L7:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$1, %ecx
	movl	$1, %edx
	call	_Z15kicache_test4_2PjS_iii@PLT
	jmp	.L41
.L39:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$2, %ecx
	movl	$1, %edx
	call	_Z15kicache_test4_4PjS_iii@PLT
	jmp	.L41
.L38:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$3, %ecx
	movl	$1, %edx
	call	_Z15kicache_test4_6PjS_iii@PLT
	jmp	.L41
.L37:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$4, %ecx
	movl	$1, %edx
	call	_Z15kicache_test4_8PjS_iii@PLT
	jmp	.L41
.L36:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$5, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_10PjS_iii@PLT
	jmp	.L41
.L35:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$6, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_12PjS_iii@PLT
	jmp	.L41
.L34:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$7, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_14PjS_iii@PLT
	jmp	.L41
.L33:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$8, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_16PjS_iii@PLT
	jmp	.L41
.L32:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$9, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_18PjS_iii@PLT
	jmp	.L41
.L31:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$10, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_20PjS_iii@PLT
	jmp	.L41
.L30:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$11, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_22PjS_iii@PLT
	jmp	.L41
.L29:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$12, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_24PjS_iii@PLT
	jmp	.L41
.L28:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$13, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_26PjS_iii@PLT
	jmp	.L41
.L27:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$14, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_28PjS_iii@PLT
	jmp	.L41
.L26:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$15, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_30PjS_iii@PLT
	jmp	.L41
.L25:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$16, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_32PjS_iii@PLT
	jmp	.L41
.L24:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$17, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_34PjS_iii@PLT
	jmp	.L41
.L23:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$18, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_36PjS_iii@PLT
	jmp	.L41
.L22:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$19, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_38PjS_iii@PLT
	jmp	.L41
.L21:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$20, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_40PjS_iii@PLT
	jmp	.L41
.L20:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$21, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_42PjS_iii@PLT
	jmp	.L41
.L19:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$22, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_44PjS_iii@PLT
	jmp	.L41
.L18:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$23, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_46PjS_iii@PLT
	jmp	.L41
.L17:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$24, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_48PjS_iii@PLT
	jmp	.L41
.L16:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$25, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_50PjS_iii@PLT
	jmp	.L41
.L15:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$26, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_52PjS_iii@PLT
	jmp	.L41
.L14:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$27, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_54PjS_iii@PLT
	jmp	.L41
.L13:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$28, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_56PjS_iii@PLT
	jmp	.L41
.L12:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$29, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_58PjS_iii@PLT
	jmp	.L41
.L11:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$30, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_60PjS_iii@PLT
	jmp	.L41
.L10:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$31, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_62PjS_iii@PLT
	jmp	.L41
.L8:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$32, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_64PjS_iii@PLT
.L41:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	jne	.L43
	call	cudaThreadSynchronize@PLT
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	leaq	200(%rsp), %rdi
	movl	$4096, %edx
	call	cudaMemcpy@PLT
	addl	204(%rsp), %ebx
	subl	200(%rsp), %ebx
	decl	%r12d
	jne	.L44
	movl	%ebx, %eax
	movl	$2, %edi
	leaq	.LC19(%rip), %rsi
	cvtsi2sdq	%rax, %xmm0
	movl	%ebp, %eax
	divsd	.LC18(%rip), %xmm0
	sall	$6, %eax
	cvtsi2sdl	%eax, %xmm1
	movb	$1, %al
	divsd	%xmm1, %xmm0
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
.L142:
	movl	%ebx, 68(%rsp,%rbp,4)
	incq	%rbp
	leaq	72(%rsp), %r15
	cmpq	$33, %rbp
	jne	.L6
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1, %ebx
	call	__printf_chk@PLT
	movl	$1, %ebp
.L48:
	movl	(%r15,%rbx,4), %eax
	leal	1(%rbx), %ecx
	xorl	%edx, %edx
	leal	1(%rbp), %esi
	divl	%ecx
	xorl	%edx, %edx
	movl	%eax, %ecx
	movslq	%ebp, %rax
	movl	72(%rsp,%rax,4), %eax
	divl	%esi
	cmpl	%ecx, %eax
	jnb	.L47
	cvtsi2sdl	%ebx, %xmm0
	mulsd	.LC6(%rip), %xmm0
	movb	$1, %al
	movl	%ebx, %ebp
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
.L47:
	incq	%rbx
	cmpq	$32, %rbx
	jne	.L48
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1, %ebx
	call	__printf_chk@PLT
	leaq	.LC9(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.L52(%rip), %r14
.L49:
	movl	%ebx, %r13d
	xorl	%ebp, %ebp
	movl	$100, %r12d
.L84:
	call	cudaGetLastError@PLT
	leal	-2(%r13), %eax
	movq	36(%rsp), %rdx
	movl	44(%rsp), %ecx
	cmpl	$27, %eax
	movq	48(%rsp), %rdi
	movl	56(%rsp), %esi
	ja	.L50
	movslq	(%r14,%rax,4), %rax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addq	%r14, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L52:
	.long	.L79-.L52
	.long	.L78-.L52
	.long	.L77-.L52
	.long	.L76-.L52
	.long	.L75-.L52
	.long	.L74-.L52
	.long	.L73-.L52
	.long	.L72-.L52
	.long	.L71-.L52
	.long	.L70-.L52
	.long	.L69-.L52
	.long	.L68-.L52
	.long	.L67-.L52
	.long	.L66-.L52
	.long	.L65-.L52
	.long	.L64-.L52
	.long	.L63-.L52
	.long	.L62-.L52
	.long	.L61-.L52
	.long	.L60-.L52
	.long	.L59-.L52
	.long	.L58-.L52
	.long	.L57-.L52
	.long	.L56-.L52
	.long	.L55-.L52
	.long	.L54-.L52
	.long	.L53-.L52
	.long	.L51-.L52
	.text
.L50:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$1, %ecx
	movl	$1, %edx
	call	_Z15kicache_test4_8PjS_iii@PLT
	jmp	.L81
.L79:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$2, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_16PjS_iii@PLT
	jmp	.L81
.L78:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$3, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_24PjS_iii@PLT
	jmp	.L81
.L77:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$4, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_32PjS_iii@PLT
	jmp	.L81
.L76:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$5, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_40PjS_iii@PLT
	jmp	.L81
.L75:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$6, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_48PjS_iii@PLT
	jmp	.L81
.L74:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$7, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_56PjS_iii@PLT
	jmp	.L81
.L73:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$8, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_64PjS_iii@PLT
	jmp	.L81
.L72:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$9, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_72PjS_iii@PLT
	jmp	.L81
.L71:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$10, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_80PjS_iii@PLT
	jmp	.L81
.L70:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$11, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_88PjS_iii@PLT
	jmp	.L81
.L69:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$12, %ecx
	movl	$1, %edx
	call	_Z16kicache_test4_96PjS_iii@PLT
	jmp	.L81
.L68:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$13, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_104PjS_iii@PLT
	jmp	.L81
.L67:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$14, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_112PjS_iii@PLT
	jmp	.L81
.L66:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$15, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_120PjS_iii@PLT
	jmp	.L81
.L65:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$16, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_128PjS_iii@PLT
	jmp	.L81
.L64:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$17, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_136PjS_iii@PLT
	jmp	.L81
.L63:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$18, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_144PjS_iii@PLT
	jmp	.L81
.L62:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$19, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_152PjS_iii@PLT
	jmp	.L81
.L61:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$20, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_160PjS_iii@PLT
	jmp	.L81
.L60:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$21, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_168PjS_iii@PLT
	jmp	.L81
.L59:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$22, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_176PjS_iii@PLT
	jmp	.L81
.L58:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$23, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_184PjS_iii@PLT
	jmp	.L81
.L57:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$24, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_192PjS_iii@PLT
	jmp	.L81
.L56:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$25, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_200PjS_iii@PLT
	jmp	.L81
.L55:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$26, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_208PjS_iii@PLT
	jmp	.L81
.L54:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$27, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_216PjS_iii@PLT
	jmp	.L81
.L53:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$28, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_224PjS_iii@PLT
	jmp	.L81
.L51:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L81
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %r8d
	movl	$29, %ecx
	movl	$1, %edx
	call	_Z17kicache_test4_232PjS_iii@PLT
.L81:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	jne	.L83
	call	cudaThreadSynchronize@PLT
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	leaq	200(%rsp), %rdi
	movl	$4096, %edx
	call	cudaMemcpy@PLT
	addl	204(%rsp), %ebp
	subl	200(%rsp), %ebp
	decl	%r12d
	jne	.L84
	movl	%ebp, %eax
	movl	$2, %edi
	leaq	.LC17(%rip), %rsi
	cvtsi2sdq	%rax, %xmm0
	movl	%ebx, %eax
	divsd	.LC18(%rip), %xmm0
	sall	$8, %eax
	cvtsi2sdl	%eax, %xmm1
	movb	$1, %al
	divsd	%xmm1, %xmm0
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
.L137:
	movl	%ebp, -4(%r15,%rbx,4)
	incq	%rbx
	cmpq	$30, %rbx
	jne	.L49
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1, %ebx
	call	__printf_chk@PLT
	movl	$1, %r12d
	leaq	.LC7(%rip), %rbp
.L88:
	movl	(%r15,%rbx,4), %eax
	leal	1(%rbx), %ecx
	xorl	%edx, %edx
	leal	1(%r12), %esi
	divl	%ecx
	xorl	%edx, %edx
	movl	%eax, %ecx
	movslq	%r12d, %rax
	movl	72(%rsp,%rax,4), %eax
	divl	%esi
	cmpl	%ecx, %eax
	jnb	.L87
	cvtsi2sdl	%ebx, %xmm0
	movq	%rbp, %rsi
	movl	$2, %edi
	movb	$1, %al
	movl	%ebx, %r12d
	addsd	%xmm0, %xmm0
	call	__printf_chk@PLT
.L87:
	incq	%rbx
	cmpq	$28, %rbx
	jne	.L88
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1, %r12d
	call	__printf_chk@PLT
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, 68(%rsp)
	movabsq	$4294967327, %rax
	movq	%rax, 60(%rsp)
.L132:
	movl	%r12d, %edx
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1, %ebx
	leaq	.LC11(%rip), %rsi
	leaq	.L92(%rip), %r13
	call	__printf_chk@PLT
	movl	$1, %eax
	movl	%r12d, %ecx
	sall	%cl, %eax
	orl	$1, %eax
	movl	%eax, 8(%rsp)
.L89:
	movl	$50, 12(%rsp)
	movl	%ebx, %r14d
	xorl	%ebp, %ebp
.L123:
	call	cudaGetLastError@PLT
	leal	-2(%r14), %eax
	movq	36(%rsp), %rdx
	movl	44(%rsp), %ecx
	cmpl	$26, %eax
	movq	60(%rsp), %rdi
	movl	68(%rsp), %esi
	ja	.L90
	movslq	0(%r13,%rax,4), %rax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addq	%r13, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L92:
	.long	.L118-.L92
	.long	.L117-.L92
	.long	.L116-.L92
	.long	.L115-.L92
	.long	.L114-.L92
	.long	.L113-.L92
	.long	.L112-.L92
	.long	.L111-.L92
	.long	.L110-.L92
	.long	.L109-.L92
	.long	.L108-.L92
	.long	.L107-.L92
	.long	.L106-.L92
	.long	.L105-.L92
	.long	.L104-.L92
	.long	.L103-.L92
	.long	.L102-.L92
	.long	.L101-.L92
	.long	.L100-.L92
	.long	.L99-.L92
	.long	.L98-.L92
	.long	.L97-.L92
	.long	.L96-.L92
	.long	.L95-.L92
	.long	.L94-.L92
	.long	.L93-.L92
	.long	.L91-.L92
	.text
.L90:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z15kicache_test4_8PjS_iii@PLT
	jmp	.L120
.L118:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_16PjS_iii@PLT
	jmp	.L120
.L117:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$3, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_24PjS_iii@PLT
	jmp	.L120
.L116:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$4, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_32PjS_iii@PLT
	jmp	.L120
.L115:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$5, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_40PjS_iii@PLT
	jmp	.L120
.L114:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$6, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_48PjS_iii@PLT
	jmp	.L120
.L113:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$7, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_56PjS_iii@PLT
	jmp	.L120
.L112:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$8, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_64PjS_iii@PLT
	jmp	.L120
.L111:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$9, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_72PjS_iii@PLT
	jmp	.L120
.L110:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$10, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_80PjS_iii@PLT
	jmp	.L120
.L109:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$11, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_88PjS_iii@PLT
	jmp	.L120
.L108:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$12, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z16kicache_test4_96PjS_iii@PLT
	jmp	.L120
.L107:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$13, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_104PjS_iii@PLT
	jmp	.L120
.L106:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$14, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_112PjS_iii@PLT
	jmp	.L120
.L105:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$15, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_120PjS_iii@PLT
	jmp	.L120
.L104:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$16, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_128PjS_iii@PLT
	jmp	.L120
.L103:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$17, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_136PjS_iii@PLT
	jmp	.L120
.L102:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$18, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_144PjS_iii@PLT
	jmp	.L120
.L101:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$19, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_152PjS_iii@PLT
	jmp	.L120
.L100:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$20, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_160PjS_iii@PLT
	jmp	.L120
.L99:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$21, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_168PjS_iii@PLT
	jmp	.L120
.L98:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$22, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_176PjS_iii@PLT
	jmp	.L120
.L97:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$23, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_184PjS_iii@PLT
	jmp	.L120
.L96:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$24, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_192PjS_iii@PLT
	jmp	.L120
.L95:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$25, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_200PjS_iii@PLT
	jmp	.L120
.L94:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$26, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_208PjS_iii@PLT
	jmp	.L120
.L93:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$27, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_216PjS_iii@PLT
	jmp	.L120
.L91:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movl	8(%rsp), %edx
	movq	24(%rsp), %rsi
	movl	$28, %ecx
	movl	$2, %r8d
	movq	16(%rsp), %rdi
	call	_Z17kicache_test4_224PjS_iii@PLT
.L120:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	jne	.L122
	call	cudaThreadSynchronize@PLT
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	leaq	200(%rsp), %rdi
	movl	$4096, %edx
	call	cudaMemcpy@PLT
	addl	204(%rsp), %ebp
	subl	200(%rsp), %ebp
	decl	12(%rsp)
	jne	.L123
	movl	%ebp, %eax
	movl	$2, %edi
	leaq	.LC17(%rip), %rsi
	cvtsi2sdq	%rax, %xmm0
	movl	%ebx, %eax
	divsd	.LC16(%rip), %xmm0
	sall	$8, %eax
	cvtsi2sdl	%eax, %xmm1
	movb	$1, %al
	divsd	%xmm1, %xmm0
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
.L134:
	movl	%ebp, -4(%r15,%rbx,4)
	incq	%rbx
	cmpq	$29, %rbx
	jne	.L89
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1, %ebx
	call	__printf_chk@PLT
	movl	$1, %ebp
.L130:
	movl	(%r15,%rbx,4), %eax
	leal	1(%rbx), %ecx
	xorl	%edx, %edx
	divl	%ecx
	xorl	%edx, %edx
	leal	1(%rbp), %ecx
	cvtsi2sdl	%eax, %xmm1
	movslq	%ebp, %rax
	movl	72(%rsp,%rax,4), %eax
	divl	%ecx
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC13(%rip), %xmm0
	comisd	%xmm0, %xmm1
	jbe	.L126
	cvtsi2sdl	%ebx, %xmm0
	leaq	.LC7(%rip), %rsi
	movb	$1, %al
	movl	%ebx, %ebp
	movl	$2, %edi
	addsd	%xmm0, %xmm0
	call	__printf_chk@PLT
.L126:
	incq	%rbx
	cmpq	$32, %rbx
	jne	.L130
	xorl	%eax, %eax
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	xorl	%eax, %eax
	cmpl	$1, %r12d
	setne	%al
	leal	9(%rax,%r12), %r12d
	cmpl	$30, %r12d
	jle	.L132
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	jmp	.L2
.L122:
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L134
.L83:
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L137
.L43:
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L142
.L2:
	movq	4296(%rsp), %rax
	subq	%fs:40, %rax
	je	.L143
	call	__stack_chk_fail@PLT
.L143:
	addq	$4312, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z14measure_icachev, .-_Z14measure_icachev
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z14measure_icachev
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2054:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2054:
	.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
.LC6:
	.long	0
	.long	1071644672
	.align 8
.LC13:
	.long	0
	.long	1072955392
	.align 8
.LC16:
	.long	0
	.long	1078525952
	.align 8
.LC18:
	.long	0
	.long	1079574528
	.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:


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


__global__ void kicache_test4_2 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_4 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_6 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_8 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_10 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_12 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_14 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_16 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_18 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_20 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_22 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_24 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_26 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_28 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_30 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_32 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_34 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_36 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_38 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_40 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_42 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_44 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_46 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_48 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_50 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_52 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_54 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_56 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_58 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_60 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_62 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_64 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_72 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_80 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_88 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_96 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_104 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_112 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_120 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_128 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_136 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_144 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_152 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_160 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_168 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_176 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_184 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_192 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_200 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_208 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_216 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_224 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
__global__ void kicache_test4_232 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
//__global__ void kicache_test4_240 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
//__global__ void kicache_test4_248 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);
//__global__ void kicache_test4_256 (unsigned int *ts, unsigned int* out, int p1, int p2, int its);




void measure_icache()
{
	
	unsigned int ts[1024];			// ts, output from kernel. Two elements used per thread.
	unsigned int *d_ts;
	unsigned int *d_out;			// Unused memory for storing output
	

    	dim3 Db = dim3(1);
    	dim3 Dg = dim3(1,1,1);
	
	// Allocate device array.
	if (hipSuccess != hipMalloc((void**)&d_ts, sizeof(ts)))
	{
		printf ("hipMalloc failed %s:%d\n", __FILE__, __LINE__);
		return;
	}
	if (hipSuccess != hipMalloc((void**)&d_out, 4))
	{
		printf ("hipMalloc failed %s:%d\n", __FILE__, __LINE__);
		return;
	}
	
    fprintf (stderr, "Running icache test...\n");
	
	

	// Measure instruction cache size
	printf ("Instruction cache size:\n");
	



	unsigned int sum_times[32];
	
	
	printf ("  0.5 KB steps: ");
	Db.x = 1;
	for (int p2 = 1; p2 <= 32; p2++)
	{
		unsigned int sum_time = 0;
		bool failed = false;
		for (int i=0;i<100 && !failed ;i++)
		{
			hipGetLastError();		// Clear previous error code, if any
			switch (p2) {
				case 1: kicache_test4_2 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 2: kicache_test4_4 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 3: kicache_test4_6 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 4: kicache_test4_8 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 5: kicache_test4_10 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 6: kicache_test4_12 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 7: kicache_test4_14 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 8: kicache_test4_16 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 9: kicache_test4_18 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 10: kicache_test4_20 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 11: kicache_test4_22 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 12: kicache_test4_24 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 13: kicache_test4_26 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 14: kicache_test4_28 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 15: kicache_test4_30 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 16: kicache_test4_32 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 17: kicache_test4_34 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 18: kicache_test4_36 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 19: kicache_test4_38 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 20: kicache_test4_40 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 21: kicache_test4_42 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 22: kicache_test4_44 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 23: kicache_test4_46 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 24: kicache_test4_48 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 25: kicache_test4_50 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 26: kicache_test4_52 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 27: kicache_test4_54 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 28: kicache_test4_56 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 29: kicache_test4_58 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 30: kicache_test4_60 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 31: kicache_test4_62 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 32: kicache_test4_64 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
			}
				
			if (hipGetLastError() != hipSuccess)
			{
				failed = true;
				break;
			}
				
			hipDeviceSynchronize();
			hipMemcpy(ts, d_ts, sizeof(ts), hipMemcpyDeviceToHost);
			
			sum_time += (ts[1] - ts[0]);
			
		}
		
		if (failed)
		{
			printf ("xxxx ");
		}
		else
		{
			// Compute average latency over the lifetime of each warp (sum_time), and average throughput of the kernel (sum_max_time).
			printf ("%.5f ", sum_time/100.0/(p2*64)); fflush(stdout);
		}
		
		sum_times[p2-1] = sum_time;
	}
	printf (" (icache = ");
	for (int last_i=1, i=1;i<32;i++)
	{
		//printf("sum time %d, %d \n",sum_times[i]/(i+1),sum_times[last_i]/(last_i+1)*1.33);
		if (sum_times[i]/(i+1) > sum_times[last_i]/(last_i+1) /**1.33*/)
		{
			printf ("%.1fKB ", i*0.5);
			last_i = i;
		}	
	}
	printf (")\n");
	
	
	printf ("  2 KB steps: ");
	Db.x = 1;
	for (int p2 = 1; p2 <= 29/*32*/; p2++)
	{
		unsigned int sum_time = 0;
		bool failed = false;
		for (int i=0;i<100 && !failed ;i++)
		{
			hipGetLastError();		// Clear previous error code, if any
			switch (p2) {
				case 1: kicache_test4_8 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 2: kicache_test4_16 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 3: kicache_test4_24 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 4: kicache_test4_32 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 5: kicache_test4_40 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 6: kicache_test4_48 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 7: kicache_test4_56 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 8: kicache_test4_64 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 9: kicache_test4_72 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 10: kicache_test4_80 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 11: kicache_test4_88 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 12: kicache_test4_96 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 13: kicache_test4_104 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 14: kicache_test4_112 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 15: kicache_test4_120 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 16: kicache_test4_128 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 17: kicache_test4_136 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 18: kicache_test4_144 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 19: kicache_test4_152 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 20: kicache_test4_160 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 21: kicache_test4_168 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 22: kicache_test4_176 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 23: kicache_test4_184 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 24: kicache_test4_192 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 25: kicache_test4_200 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 26: kicache_test4_208 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 27: kicache_test4_216 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 28: kicache_test4_224 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				case 29: kicache_test4_232 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				//case 30: kicache_test4_240 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				//case 31: kicache_test4_248 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;
				//case 32: kicache_test4_256 <<<Dg, Db>>>(d_ts, d_out, 1, p2, 2); break;

			}			

				
			if (hipGetLastError() != hipSuccess)
			{
				failed = true;
				break;
			}
				
			hipDeviceSynchronize();
			hipMemcpy(ts, d_ts, sizeof(ts), hipMemcpyDeviceToHost);
			
			sum_time += (ts[1] - ts[0]);
		}
		
		if (failed)
		{
			printf ("xxxx ");
		}
		else
		{
			// Compute average latency over the lifetime of each warp (sum_time), and average throughput of the kernel (sum_max_time).
			printf ("%.1f ", sum_time/100.0/(p2*256)); fflush(stdout);
		}
		sum_times[p2-1] = sum_time;
	}
	printf (" (icache = ");
	for (int last_i=1, i=1;i<28/*32*/;i++)
	{
		if (sum_times[i]/(i+1) > sum_times[last_i]/(last_i+1) /**1.33*/)
		{
			printf ("%.1fKB ", i*2.0);
			last_i = i;
		}	
	}
	printf (")\n");
	
	
	printf ("\n Test instruction cache sharing by running two thread blocks concurently.\n");
	
	Db.x = 1;
	dim3 Dg2 = dim3(31,1,1);
	
	for (int blk2 = 1; blk2 <31; blk2+= (blk2 == 1? 9 : 10))
	{
		printf ("  TPC 0,%d (2 KB steps): ", blk2);
		int mask = (1<<blk2) | 1;	// Enable two blocks for execution
		for (int p2 = 1; p2 <= 28/*32*/; p2++)
		{
			unsigned int sum_time = 0;
			bool failed = false;
			for (int i=0;i<50 && !failed ;i++)
			{
				hipGetLastError();		// Clear previous error code, if any
				switch (p2) {
					case 1: kicache_test4_8 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 2: kicache_test4_16 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 3: kicache_test4_24 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 4: kicache_test4_32 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 5: kicache_test4_40 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 6: kicache_test4_48 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 7: kicache_test4_56 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 8: kicache_test4_64 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 9: kicache_test4_72 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 10: kicache_test4_80 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 11: kicache_test4_88 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 12: kicache_test4_96 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 13: kicache_test4_104 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 14: kicache_test4_112 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 15: kicache_test4_120 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 16: kicache_test4_128 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 17: kicache_test4_136 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 18: kicache_test4_144 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 19: kicache_test4_152 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 20: kicache_test4_160 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 21: kicache_test4_168 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 22: kicache_test4_176 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 23: kicache_test4_184 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 24: kicache_test4_192 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 25: kicache_test4_200 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 26: kicache_test4_208 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 27: kicache_test4_216 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					case 28: kicache_test4_224 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 29: kicache_test4_232 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 30: kicache_test4_240 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 31: kicache_test4_248 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;
					//case 32: kicache_test4_256 <<<Dg2, Db>>>(d_ts, d_out, mask, p2, 2); break;

				}

				if (hipGetLastError() != hipSuccess)
				{
					failed = true;
					break;
				}
					
				hipDeviceSynchronize();
				hipMemcpy(ts, d_ts, sizeof(ts), hipMemcpyDeviceToHost);
				
				sum_time += (ts[1] - ts[0]);
			}
			
			if (failed)
			{
				printf ("xxxx ");
			}
			else
			{
				printf ("%.1f ", sum_time/50.0/(p2*256)); fflush(stdout);
			}
			sum_times[p2-1] = sum_time;
		}
		printf (" (apparent icache = ");
		for (int last_i=1, i=1;i<32;i++)
		{
			if (sum_times[i]/(i+1) > sum_times[last_i]/(last_i+1) *1.25)
			{
				printf ("%.1fKB ", i*2.0);
				last_i = i;
			}	
		}
		printf (")\n");	
	}
	
	
	printf ("\n");
	
	hipFree(d_ts);
	hipFree(d_out);

	
	
}

int main()
{
	measure_icache();
}




	.text
	.file	"icache.hip"
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z14measure_icachev
.LCPI0_0:
	.quad	0x4059000000000000              # double 100
.LCPI0_1:
	.quad	0x3fe0000000000000              # double 0.5
.LCPI0_2:
	.quad	0x4049000000000000              # double 50
.LCPI0_3:
	.quad	0x3ff4000000000000              # double 1.25
	.text
	.globl	_Z14measure_icachev
	.type	_Z14measure_icachev,@function
_Z14measure_icachev:                    # @_Z14measure_icachev
	.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	$4264, %rsp                     # imm = 0x10A8
	.cfi_def_cfa_offset 4320
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	8(%rsp), %rdi
	movl	$4096, %esi                     # imm = 0x1000
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB0_3
# %bb.1:
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	movl	$80, %edx
	jmp	.LBB0_2
.LBB0_3:
	leaq	16(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB0_5
# %bb.4:
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	movl	$85, %edx
.LBB0_2:
	xorl	%eax, %eax
	callq	printf
.LBB0_222:
	addq	$4264, %rsp                     # imm = 0x10A8
	.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
.LBB0_5:
	.cfi_def_cfa_offset 4320
	movq	stderr(%rip), %rcx
	movl	$1, %r15d
	movl	$.L.str.2, %edi
	movl	$23, %esi
	movl	$1, %edx
	callq	fwrite@PLT
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$.L.str.4, %edi
	xorl	%eax, %eax
	callq	printf
	leaq	160(%rsp), %r14
.LBB0_6:                                # %.preheader1377
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_7 Depth 2
	movl	$100, %r12d
	leal	-1(%r15), %r13d
	xorl	%ebp, %ebp
.LBB0_7:                                #   Parent Loop BB0_6 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	hipGetLastError
	cmpl	$31, %r13d
	ja	.LBB0_73
# %bb.8:                                #   in Loop: Header=BB0_7 Depth=2
	jmpq	*.LJTI0_0(,%r13,8)
.LBB0_9:                                #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.10:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$1, %ecx
	movl	$2, %r8d
	callq	_Z30__device_stub__kicache_test4_2PjS_iii
	jmp	.LBB0_73
.LBB0_63:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.64:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$28, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_56PjS_iii
	jmp	.LBB0_73
.LBB0_41:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.42:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$17, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_34PjS_iii
	jmp	.LBB0_73
.LBB0_37:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.38:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$15, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_30PjS_iii
	jmp	.LBB0_73
.LBB0_31:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.32:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$12, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_24PjS_iii
	jmp	.LBB0_73
.LBB0_17:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.18:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$5, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_10PjS_iii
	jmp	.LBB0_73
.LBB0_33:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.34:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$13, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_26PjS_iii
	jmp	.LBB0_73
.LBB0_27:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.28:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$10, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_20PjS_iii
	jmp	.LBB0_73
.LBB0_59:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.60:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$26, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_52PjS_iii
	jmp	.LBB0_73
.LBB0_57:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.58:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$25, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_50PjS_iii
	jmp	.LBB0_73
.LBB0_13:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.14:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$3, %ecx
	movl	$2, %r8d
	callq	_Z30__device_stub__kicache_test4_6PjS_iii
	jmp	.LBB0_73
.LBB0_39:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.40:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$16, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_32PjS_iii
	jmp	.LBB0_73
.LBB0_15:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.16:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$4, %ecx
	movl	$2, %r8d
	callq	_Z30__device_stub__kicache_test4_8PjS_iii
	jmp	.LBB0_73
.LBB0_23:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.24:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$8, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_16PjS_iii
	jmp	.LBB0_73
.LBB0_11:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.12:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$2, %ecx
	movl	$2, %r8d
	callq	_Z30__device_stub__kicache_test4_4PjS_iii
	jmp	.LBB0_73
.LBB0_71:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.72:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$32, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_64PjS_iii
	jmp	.LBB0_73
.LBB0_43:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.44:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$18, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_36PjS_iii
	jmp	.LBB0_73
.LBB0_49:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.50:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$21, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_42PjS_iii
	jmp	.LBB0_73
.LBB0_19:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.20:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$6, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_12PjS_iii
	jmp	.LBB0_73
.LBB0_51:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.52:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$22, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_44PjS_iii
	jmp	.LBB0_73
.LBB0_35:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.36:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$14, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_28PjS_iii
	jmp	.LBB0_73
.LBB0_21:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.22:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$7, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_14PjS_iii
	jmp	.LBB0_73
.LBB0_45:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.46:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$19, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_38PjS_iii
	jmp	.LBB0_73
.LBB0_29:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.30:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$11, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_22PjS_iii
	jmp	.LBB0_73
.LBB0_25:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.26:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$9, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_18PjS_iii
	jmp	.LBB0_73
.LBB0_69:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.70:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$31, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_62PjS_iii
	jmp	.LBB0_73
.LBB0_53:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.54:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$23, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_46PjS_iii
	jmp	.LBB0_73
.LBB0_47:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.48:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$20, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_40PjS_iii
	jmp	.LBB0_73
.LBB0_55:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.56:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$24, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_48PjS_iii
	jmp	.LBB0_73
.LBB0_65:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.66:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$29, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_58PjS_iii
	jmp	.LBB0_73
.LBB0_67:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.68:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$30, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_60PjS_iii
	jmp	.LBB0_73
.LBB0_61:                               #   in Loop: Header=BB0_7 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_73
# %bb.62:                               #   in Loop: Header=BB0_7 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$27, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_54PjS_iii
.LBB0_73:                               #   in Loop: Header=BB0_7 Depth=2
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_223
# %bb.74:                               #   in Loop: Header=BB0_7 Depth=2
	callq	hipDeviceSynchronize
	movq	8(%rsp), %rsi
	movl	$4096, %edx                     # imm = 0x1000
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	addl	164(%rsp), %ebp
	subl	160(%rsp), %ebp
	decl	%r12d
	jne	.LBB0_7
# %bb.75:                               # %.critedge
                                        #   in Loop: Header=BB0_6 Depth=1
	movl	%ebp, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	movl	%r15d, %eax
	shll	$6, %eax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%eax, %xmm1
	divsd	%xmm1, %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	jmp	.LBB0_76
.LBB0_223:                              #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
.LBB0_76:                               #   in Loop: Header=BB0_6 Depth=1
	movl	%ebp, 28(%rsp,%r15,4)
	incq	%r15
	cmpq	$33, %r15
	jne	.LBB0_6
# %bb.77:
	movl	$.L.str.7, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %esi
	movl	$1, %ebx
	movl	$2, %ebp
.LBB0_78:                               # =>This Inner Loop Header: Depth=1
	movl	32(%rsp,%rbx,4), %eax
	xorl	%edx, %edx
	divl	%ebp
	movl	%eax, %ecx
	movslq	%esi, %rdi
	movl	32(%rsp,%rdi,4), %eax
	incl	%edi
	xorl	%edx, %edx
	divl	%edi
	cmpl	%eax, %ecx
	jbe	.LBB0_80
# %bb.79:                               #   in Loop: Header=BB0_78 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2sd	%ebx, %xmm0
	mulsd	.LCPI0_1(%rip), %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	movl	%ebx, %esi
.LBB0_80:                               #   in Loop: Header=BB0_78 Depth=1
	incq	%rbx
	incl	%ebp
	cmpq	$32, %rbx
	jne	.LBB0_78
# %bb.81:
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$.L.str.10, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %r15d
	leaq	160(%rsp), %r14
.LBB0_82:                               # %.preheader1376
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_83 Depth 2
	movl	$100, %r12d
	leal	-1(%r15), %r13d
	xorl	%ebp, %ebp
.LBB0_83:                               #   Parent Loop BB0_82 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	hipGetLastError
	cmpl	$28, %r13d
	ja	.LBB0_143
# %bb.84:                               #   in Loop: Header=BB0_83 Depth=2
	jmpq	*.LJTI0_1(,%r13,8)
.LBB0_85:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.86:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$1, %ecx
	movl	$2, %r8d
	callq	_Z30__device_stub__kicache_test4_8PjS_iii
	jmp	.LBB0_143
.LBB0_139:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.140:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$28, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_224PjS_iii
	jmp	.LBB0_143
.LBB0_117:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.118:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$17, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_136PjS_iii
	jmp	.LBB0_143
.LBB0_113:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.114:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$15, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_120PjS_iii
	jmp	.LBB0_143
.LBB0_107:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.108:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$12, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_96PjS_iii
	jmp	.LBB0_143
.LBB0_93:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.94:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$5, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_40PjS_iii
	jmp	.LBB0_143
.LBB0_109:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.110:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$13, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_104PjS_iii
	jmp	.LBB0_143
.LBB0_103:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.104:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$10, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_80PjS_iii
	jmp	.LBB0_143
.LBB0_135:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.136:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$26, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_208PjS_iii
	jmp	.LBB0_143
.LBB0_133:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.134:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$25, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_200PjS_iii
	jmp	.LBB0_143
.LBB0_89:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.90:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$3, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_24PjS_iii
	jmp	.LBB0_143
.LBB0_115:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.116:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$16, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_128PjS_iii
	jmp	.LBB0_143
.LBB0_91:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.92:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$4, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_32PjS_iii
	jmp	.LBB0_143
.LBB0_99:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.100:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$8, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_64PjS_iii
	jmp	.LBB0_143
.LBB0_87:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.88:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$2, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_16PjS_iii
	jmp	.LBB0_143
.LBB0_119:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.120:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$18, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_144PjS_iii
	jmp	.LBB0_143
.LBB0_125:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.126:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$21, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_168PjS_iii
	jmp	.LBB0_143
.LBB0_95:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.96:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$6, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_48PjS_iii
	jmp	.LBB0_143
.LBB0_127:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.128:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$22, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_176PjS_iii
	jmp	.LBB0_143
.LBB0_111:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.112:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$14, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_112PjS_iii
	jmp	.LBB0_143
.LBB0_97:                               #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.98:                               #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$7, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_56PjS_iii
	jmp	.LBB0_143
.LBB0_121:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.122:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$19, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_152PjS_iii
	jmp	.LBB0_143
.LBB0_105:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.106:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$11, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_88PjS_iii
	jmp	.LBB0_143
.LBB0_101:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.102:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$9, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_72PjS_iii
	jmp	.LBB0_143
.LBB0_129:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.130:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$23, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_184PjS_iii
	jmp	.LBB0_143
.LBB0_123:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.124:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$20, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_160PjS_iii
	jmp	.LBB0_143
.LBB0_131:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.132:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$24, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_192PjS_iii
	jmp	.LBB0_143
.LBB0_141:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.142:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$29, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_232PjS_iii
	jmp	.LBB0_143
.LBB0_137:                              #   in Loop: Header=BB0_83 Depth=2
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_143
# %bb.138:                              #   in Loop: Header=BB0_83 Depth=2
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$1, %edx
	movl	$27, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_216PjS_iii
.LBB0_143:                              #   in Loop: Header=BB0_83 Depth=2
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_224
# %bb.144:                              #   in Loop: Header=BB0_83 Depth=2
	callq	hipDeviceSynchronize
	movq	8(%rsp), %rsi
	movl	$4096, %edx                     # imm = 0x1000
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	addl	164(%rsp), %ebp
	subl	160(%rsp), %ebp
	decl	%r12d
	jne	.LBB0_83
# %bb.145:                              # %.critedge1010
                                        #   in Loop: Header=BB0_82 Depth=1
	movl	%ebp, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	movl	%r15d, %eax
	shll	$8, %eax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%eax, %xmm1
	divsd	%xmm1, %xmm0
	movl	$.L.str.11, %edi
	movb	$1, %al
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	jmp	.LBB0_146
.LBB0_224:                              #   in Loop: Header=BB0_82 Depth=1
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
.LBB0_146:                              #   in Loop: Header=BB0_82 Depth=1
	movl	%ebp, 28(%rsp,%r15,4)
	incq	%r15
	cmpq	$30, %r15
	jne	.LBB0_82
# %bb.147:
	movl	$.L.str.7, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %esi
	movl	$1, %ebx
	movl	$2, %ebp
.LBB0_148:                              # =>This Inner Loop Header: Depth=1
	movl	32(%rsp,%rbx,4), %eax
	xorl	%edx, %edx
	divl	%ebp
	movl	%eax, %ecx
	movslq	%esi, %rdi
	movl	32(%rsp,%rdi,4), %eax
	incl	%edi
	xorl	%edx, %edx
	divl	%edi
	cmpl	%eax, %ecx
	jbe	.LBB0_150
# %bb.149:                              #   in Loop: Header=BB0_148 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2sd	%ebx, %xmm0
	addsd	%xmm0, %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	movl	%ebx, %esi
.LBB0_150:                              #   in Loop: Header=BB0_148 Depth=1
	incq	%rbx
	incl	%ebp
	cmpq	$28, %rbx
	jne	.LBB0_148
# %bb.151:
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$1, %ebx
	leaq	160(%rsp), %r15
.LBB0_152:                              # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_153 Depth 2
                                        #       Child Loop BB0_154 Depth 3
                                        #     Child Loop BB0_217 Depth 2
	movl	$.L.str.13, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %ebp
	movq	%rbx, 24(%rsp)                  # 8-byte Spill
	movl	%ebx, %ecx
	shll	%cl, %ebp
	incl	%ebp
	movl	$1, %r12d
.LBB0_153:                              # %.preheader
                                        #   Parent Loop BB0_152 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB0_154 Depth 3
	xorl	%r13d, %r13d
	movl	$50, %r14d
	leal	-1(%r12), %ebx
.LBB0_154:                              #   Parent Loop BB0_152 Depth=1
                                        #     Parent Loop BB0_153 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	callq	hipGetLastError
	cmpl	$27, %ebx
	ja	.LBB0_212
# %bb.155:                              #   in Loop: Header=BB0_154 Depth=3
	jmpq	*.LJTI0_2(,%rbx,8)
.LBB0_156:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.157:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$1, %ecx
	movl	$2, %r8d
	callq	_Z30__device_stub__kicache_test4_8PjS_iii
	jmp	.LBB0_212
.LBB0_210:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.211:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$28, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_224PjS_iii
	jmp	.LBB0_212
.LBB0_188:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.189:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$17, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_136PjS_iii
	jmp	.LBB0_212
.LBB0_184:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.185:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$15, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_120PjS_iii
	jmp	.LBB0_212
.LBB0_178:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.179:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$12, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_96PjS_iii
	jmp	.LBB0_212
.LBB0_164:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.165:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$5, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_40PjS_iii
	jmp	.LBB0_212
.LBB0_180:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.181:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$13, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_104PjS_iii
	jmp	.LBB0_212
.LBB0_174:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.175:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$10, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_80PjS_iii
	jmp	.LBB0_212
.LBB0_206:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.207:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$26, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_208PjS_iii
	jmp	.LBB0_212
.LBB0_204:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.205:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$25, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_200PjS_iii
	jmp	.LBB0_212
.LBB0_160:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.161:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$3, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_24PjS_iii
	jmp	.LBB0_212
.LBB0_186:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.187:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$16, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_128PjS_iii
	jmp	.LBB0_212
.LBB0_162:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.163:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$4, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_32PjS_iii
	jmp	.LBB0_212
.LBB0_170:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.171:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$8, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_64PjS_iii
	jmp	.LBB0_212
.LBB0_158:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.159:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$2, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_16PjS_iii
	jmp	.LBB0_212
.LBB0_190:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.191:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$18, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_144PjS_iii
	jmp	.LBB0_212
.LBB0_196:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.197:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$21, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_168PjS_iii
	jmp	.LBB0_212
.LBB0_166:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.167:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$6, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_48PjS_iii
	jmp	.LBB0_212
.LBB0_198:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.199:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$22, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_176PjS_iii
	jmp	.LBB0_212
.LBB0_182:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.183:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$14, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_112PjS_iii
	jmp	.LBB0_212
.LBB0_168:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.169:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$7, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_56PjS_iii
	jmp	.LBB0_212
.LBB0_192:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.193:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$19, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_152PjS_iii
	jmp	.LBB0_212
.LBB0_176:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.177:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$11, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_88PjS_iii
	jmp	.LBB0_212
.LBB0_172:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.173:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$9, %ecx
	movl	$2, %r8d
	callq	_Z31__device_stub__kicache_test4_72PjS_iii
	jmp	.LBB0_212
.LBB0_200:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.201:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$23, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_184PjS_iii
	jmp	.LBB0_212
.LBB0_194:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.195:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$20, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_160PjS_iii
	jmp	.LBB0_212
.LBB0_202:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.203:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$24, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_192PjS_iii
	jmp	.LBB0_212
.LBB0_208:                              #   in Loop: Header=BB0_154 Depth=3
	movabsq	$4294967297, %rdx               # imm = 0x100000001
	leaq	30(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_212
# %bb.209:                              #   in Loop: Header=BB0_154 Depth=3
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%ebp, %edx
	movl	$27, %ecx
	movl	$2, %r8d
	callq	_Z32__device_stub__kicache_test4_216PjS_iii
.LBB0_212:                              #   in Loop: Header=BB0_154 Depth=3
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_225
# %bb.213:                              #   in Loop: Header=BB0_154 Depth=3
	callq	hipDeviceSynchronize
	movq	8(%rsp), %rsi
	movl	$4096, %edx                     # imm = 0x1000
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	addl	164(%rsp), %r13d
	subl	160(%rsp), %r13d
	decl	%r14d
	jne	.LBB0_154
# %bb.214:                              # %.critedge1012
                                        #   in Loop: Header=BB0_153 Depth=2
	movl	%r13d, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_2(%rip), %xmm0
	movl	%r12d, %eax
	shll	$8, %eax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%eax, %xmm1
	divsd	%xmm1, %xmm0
	movl	$.L.str.11, %edi
	movb	$1, %al
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	jmp	.LBB0_215
.LBB0_225:                              #   in Loop: Header=BB0_153 Depth=2
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
.LBB0_215:                              #   in Loop: Header=BB0_153 Depth=2
	movl	%r13d, 28(%rsp,%r12,4)
	incq	%r12
	cmpq	$29, %r12
	jne	.LBB0_153
# %bb.216:                              #   in Loop: Header=BB0_152 Depth=1
	movl	$.L.str.14, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %ecx
	movl	$1, %ebx
	movl	$2, %ebp
.LBB0_217:                              #   Parent Loop BB0_152 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	32(%rsp,%rbx,4), %eax
	xorl	%edx, %edx
	divl	%ebp
                                        # kill: def $eax killed $eax def $rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	movslq	%ecx, %rsi
	movl	32(%rsp,%rsi,4), %eax
	incl	%esi
	xorl	%edx, %edx
	divl	%esi
                                        # kill: def $eax killed $eax def $rax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%rax, %xmm1
	mulsd	.LCPI0_3(%rip), %xmm1
	ucomisd	%xmm1, %xmm0
	jbe	.LBB0_219
# %bb.218:                              #   in Loop: Header=BB0_217 Depth=2
	xorps	%xmm0, %xmm0
	cvtsi2sd	%ebx, %xmm0
	addsd	%xmm0, %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	movl	%ebx, %ecx
.LBB0_219:                              #   in Loop: Header=BB0_217 Depth=2
	incq	%rbx
	incl	%ebp
	cmpq	$32, %rbx
	jne	.LBB0_217
# %bb.220:                              #   in Loop: Header=BB0_152 Depth=1
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	xorl	%eax, %eax
	movq	24(%rsp), %rbx                  # 8-byte Reload
	cmpl	$1, %ebx
	sete	%al
	negl	%eax
	addl	%eax, %ebx
	addl	$10, %ebx
	cmpl	$31, %ebx
	jb	.LBB0_152
# %bb.221:
	movl	$10, %edi
	callq	putchar@PLT
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	jmp	.LBB0_222
.Lfunc_end0:
	.size	_Z14measure_icachev, .Lfunc_end0-_Z14measure_icachev
	.cfi_endproc
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
.LJTI0_0:
	.quad	.LBB0_9
	.quad	.LBB0_11
	.quad	.LBB0_13
	.quad	.LBB0_15
	.quad	.LBB0_17
	.quad	.LBB0_19
	.quad	.LBB0_21
	.quad	.LBB0_23
	.quad	.LBB0_25
	.quad	.LBB0_27
	.quad	.LBB0_29
	.quad	.LBB0_31
	.quad	.LBB0_33
	.quad	.LBB0_35
	.quad	.LBB0_37
	.quad	.LBB0_39
	.quad	.LBB0_41
	.quad	.LBB0_43
	.quad	.LBB0_45
	.quad	.LBB0_47
	.quad	.LBB0_49
	.quad	.LBB0_51
	.quad	.LBB0_53
	.quad	.LBB0_55
	.quad	.LBB0_57
	.quad	.LBB0_59
	.quad	.LBB0_61
	.quad	.LBB0_63
	.quad	.LBB0_65
	.quad	.LBB0_67
	.quad	.LBB0_69
	.quad	.LBB0_71
.LJTI0_1:
	.quad	.LBB0_85
	.quad	.LBB0_87
	.quad	.LBB0_89
	.quad	.LBB0_91
	.quad	.LBB0_93
	.quad	.LBB0_95
	.quad	.LBB0_97
	.quad	.LBB0_99
	.quad	.LBB0_101
	.quad	.LBB0_103
	.quad	.LBB0_105
	.quad	.LBB0_107
	.quad	.LBB0_109
	.quad	.LBB0_111
	.quad	.LBB0_113
	.quad	.LBB0_115
	.quad	.LBB0_117
	.quad	.LBB0_119
	.quad	.LBB0_121
	.quad	.LBB0_123
	.quad	.LBB0_125
	.quad	.LBB0_127
	.quad	.LBB0_129
	.quad	.LBB0_131
	.quad	.LBB0_133
	.quad	.LBB0_135
	.quad	.LBB0_137
	.quad	.LBB0_139
	.quad	.LBB0_141
.LJTI0_2:
	.quad	.LBB0_156
	.quad	.LBB0_158
	.quad	.LBB0_160
	.quad	.LBB0_162
	.quad	.LBB0_164
	.quad	.LBB0_166
	.quad	.LBB0_168
	.quad	.LBB0_170
	.quad	.LBB0_172
	.quad	.LBB0_174
	.quad	.LBB0_176
	.quad	.LBB0_178
	.quad	.LBB0_180
	.quad	.LBB0_182
	.quad	.LBB0_184
	.quad	.LBB0_186
	.quad	.LBB0_188
	.quad	.LBB0_190
	.quad	.LBB0_192
	.quad	.LBB0_194
	.quad	.LBB0_196
	.quad	.LBB0_198
	.quad	.LBB0_200
	.quad	.LBB0_202
	.quad	.LBB0_204
	.quad	.LBB0_206
	.quad	.LBB0_208
	.quad	.LBB0_210
                                        # -- End function
	.text
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	_Z14measure_icachev
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"hipMalloc failed %s:%d\n"
	.size	.L.str, 24

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/mirsoleimani/SAMMicrobenchmark/master/cudabmk/icache.hip"
	.size	.L.str.1, 119

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Running icache test...\n"
	.size	.L.str.2, 24

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"  0.5 KB steps: "
	.size	.L.str.4, 17

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"xxxx "
	.size	.L.str.5, 6

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%.5f "
	.size	.L.str.6, 6

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	" (icache = "
	.size	.L.str.7, 12

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"%.1fKB "
	.size	.L.str.8, 8

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"  2 KB steps: "
	.size	.L.str.10, 15

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"%.1f "
	.size	.L.str.11, 6

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"  TPC 0,%d (2 KB steps): "
	.size	.L.str.13, 26

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	" (apparent icache = "
	.size	.L.str.14, 21

	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Instruction cache size:"
	.size	.Lstr, 24

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"\n Test instruction cache sharing by running two thread blocks concurently."
	.size	.Lstr.3, 75

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	")"
	.size	.Lstr.4, 2

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

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

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

	.end_amdgpu_metadata
