#include "cuda_MP7.cuh"

void cuda_MP7(int argc, char* argv[])
{
  /* Case of 0 arguments: Default seed is used */
  if (argc < 2) {
	srand(0);
  }
  /* Case of 1 argument: Seed is specified as first command line argument */
  else {
	int seed = atoi(argv[1]);
	srand(seed);
  }

  uint8_t *gold_bins = (uint8_t*)malloc(HISTO_HEIGHT*HISTO_WIDTH * sizeof(uint8_t));

  // Use kernel_bins for your final result
  uint8_t *kernel_bins = (uint8_t*)malloc(HISTO_HEIGHT*HISTO_WIDTH * sizeof(uint8_t));

  // A 2D array of histogram bin-ids.  One can think of each of these bins-ids as
  // being associated with a pixel in a 2D image.
  uint32_t **input = generate_histogram_bins();

  cout << "Input example: " << endl;
  for (int i = 0; i < 10; i++) {
	for (int j = 0; j < 14; j++) {
	  cout << input[i][j] << " ";
	}
	cout << endl;
  }

//  TIME_IT("ref_2dhisto",
//	  50,
//	  ref_2dhisto(input, INPUT_HEIGHT, INPUT_WIDTH, gold_bins);)

  ref_2dhisto(input, INPUT_HEIGHT, INPUT_WIDTH, gold_bins);
	/* Include your setup code below (temp variables, function calls, etc.) */

  uint32_t **d_input = NULL;
  cudaMalloc((void**)&d_input, INPUT_HEIGHT * INPUT_WIDTH * sizeof(uint32_t));

  uint8_t *d_obins = NULL;
  cudaMalloc((void**)&d_obins, HISTO_HEIGHT * HISTO_WIDTH * sizeof(uint8_t));
  
  uint32_t *temp_bins = NULL;
  cudaMalloc((void**)&temp_bins, HISTO_HEIGHT * HISTO_WIDTH * sizeof(uint32_t));

  cudaMemcpy(d_input, &(input[0][0]), INPUT_HEIGHT * INPUT_WIDTH * sizeof(uint32_t),
	cudaMemcpyHostToDevice);

	/* End of setup code */

	/* This is the call you will use to time your parallel implementation */
//	TIME_IT("opt_2dhisto",
//	  50,
//	  opt_2dhisto( /*Define your own function parameters*/);)

  opt_2dhisto(d_input, d_obins, temp_bins, INPUT_HEIGHT, INPUT_WIDTH);
	/* Include your teardown code below (temporary variables, function calls, etc.) */

  cudaMemcpy(kernel_bins, d_obins, HISTO_HEIGHT*HISTO_WIDTH * sizeof(uint8_t),
	cudaMemcpyDeviceToHost);

  cudaFree(temp_bins);
  cudaFree(d_input);
  cudaFree(d_obins);
	/* End of teardown code */

  int passed = 1;
  cout << "Gold_Bins vs. Kernal_Bins" << endl;
  for (int i = 0; i < HISTO_HEIGHT*HISTO_WIDTH; i++) {
	if (gold_bins[i] != kernel_bins[i]) {
	  cout << i << " " << gold_bins[i] << " " << kernel_bins[i] << endl;
	  passed = 0;
	  break;
	}
  }
  (passed) ? printf("\n    Test PASSED\n") : printf("\n    Test FAILED\n");

  free(gold_bins);
  free(kernel_bins);
}

int ref_2dhisto(uint32_t *input[], size_t height, size_t width, uint8_t bins[])
{

  // Zero out all the bins
  memset(bins, 0, HISTO_HEIGHT*HISTO_WIDTH * sizeof(bins[0]));

  for (size_t j = 0; j < height; ++j)
  {
	for (size_t i = 0; i < width; ++i)
	{
	  const uint32_t value = input[j][i];

	  uint8_t *p = (uint8_t*)bins;

	  // Increment the appropriate bin, but do not roll-over the max value
	  if (p[value] < UINT8_MAX)
		++p[value];
	}
  }

  return 0;
}

void** alloc_2d(size_t y_size, size_t x_size, size_t element_size)
{
  const size_t x_size_padded = (x_size + 128) & 0xFFFFFF80;

  uint8_t *data = (uint8_t*)calloc(x_size_padded * y_size, element_size);
  void   **res = (void**)calloc(y_size, sizeof(void*));

  if (data == 0 || res == 0)
  {
	free(data);
	free(res);
	res = 0;
	goto exit;
  }

  for (size_t i = 0; i < y_size; ++i)
	res[i] = data + (i * x_size_padded * element_size);

exit:
  return res;
}

// Generate another bin for the histogram.  The bins are created as a random walk ...
static uint32_t next_bin(uint32_t pix)
{
  const uint16_t bottom = pix & ((1 << HISTO_LOG) - 1);
  const uint16_t top = (uint16_t)(pix >> HISTO_LOG);

  int new_bottom = NEXT(bottom, SPREAD_BOTTOM)
	CLAMP(new_bottom, 0, HISTO_WIDTH - 1)

  int new_top = NEXT(top, SPREAD_TOP)
	CLAMP(new_top, 0, HISTO_HEIGHT - 1)

	const uint32_t result = (new_bottom | (new_top << HISTO_LOG));

  return result;
}

// Return a 2D array of histogram bin-ids.  This function generates
// bin-ids with correlation characteristics similar to some actual images.
// The key point here is that the pixels (and thus the bin-ids) are *NOT*
// randomly distributed ... a given pixel tends to be similar to the
// pixels near it.
static uint32_t **generate_histogram_bins()
{
  uint32_t **input = (uint32_t**)alloc_2d(INPUT_HEIGHT, INPUT_WIDTH, sizeof(uint32_t));

  input[0][0] = HISTO_WIDTH / 2 | ((HISTO_HEIGHT / 2) << HISTO_LOG);
  for (int i = 1; i < INPUT_WIDTH; ++i)
	input[0][i] = next_bin(input[0][i - 1]);
  for (int j = 1; j < INPUT_HEIGHT; ++j)
  {
	input[j][0] = next_bin(input[j - 1][0]);
	for (int i = 1; i < INPUT_WIDTH; ++i)
	  input[j][i] = next_bin(input[j][i - 1]);
  }

  return input;
}

/*
int gettimeofday(struct timeval * tp, struct timezone * tzp)
{
  // Note: some broken versions only have 8 trailing zero's, the correct epoch has 9 trailing zero's
  // This magic number is the number of 100 nanosecond intervals since January 1, 1601 (UTC)
  // until 00:00:00 January 1, 1970 
  static const uint64_t EPOCH = ((uint64_t)116444736000000000ULL);

  SYSTEMTIME  system_time;
  FILETIME    file_time;
  uint64_t    time;

  GetSystemTime(&system_time);
  SystemTimeToFileTime(&system_time, &file_time);
  time = ((uint64_t)file_time.dwLowDateTime);
  time += ((uint64_t)file_time.dwHighDateTime) << 32;

  tp->tv_sec = (long)((time - EPOCH) / 10000000L);
  tp->tv_usec = (long)(system_time.wMilliseconds * 1000);
  return 0;
}
*/

void opt_2dhisto(uint32_t **d_input, uint8_t *d_bins, uint32_t *d_temp_bins, 
  size_t height, size_t width)
{
  /* This function should only contain a call to the GPU
  histogramming kernel. Any memory allocations and
  transfers must be done outside this function */

  unsigned int num_blocks = ceil((float)width / TILE_SIZE_MP7);

  histo_kernel << <num_blocks, 1024 >> > (d_input, d_temp_bins);

  histo_32to8_kernel << <1, 1024 >> >(d_bins, d_temp_bins, 1024);

  cudaThreadSynchronize();

}

__global__ void histo_kernel(uint32_t **d_input, uint32_t *d_ouput)
{
  __shared__ uint32_t private_bins[HISTO_HEIGHT];
  unsigned int idx = threadIdx.x + blockIdx.x * blockDim.x;
  unsigned int ti = threadIdx.x;
  private_bins[ti] = 0;

  __syncthreads();

  int start_col = TILE_SIZE_MP7 * blockIdx.x;

  for (int i = 0; i < TILE_SIZE_MP7; i++) {
	if (start_col + i < INPUT_WIDTH)
	  atomicAdd(&private_bins[d_input[ti][start_col + i]], 1);
  }

  __syncthreads();

  atomicAdd(&(d_ouput[ti]), private_bins[ti]);
  atomicAdd(&(d_ouput[ti + 512]), private_bins[ti + 512]);
}

__global__ void histo_cuda_kernel(int *d_input, int *d_output)
{
  __shared__ int private_bins[HISTO_MAX];
//  unsigned int idx = threadIdx.x + blockIdx.x * blockDim.x;
  unsigned int ti = threadIdx.x;

  if (ti < HISTO_MAX)
	private_bins[ti] = 0;

  __syncthreads();

  int start_col = TILE_SIZE_MP7 * blockIdx.x;

  for (int i = 0; i < TILE_SIZE_MP7; i++) {
	if (start_col + i < INPUT_WIDTH)
	  atomicAdd(&(private_bins[d_input[ti * INPUT_WIDTH + (start_col + i)]]), 1);
  }

  __syncthreads();

  if (ti < HISTO_MAX)
	atomicAdd(&(d_output[ti]), private_bins[ti]);

}


__global__ void histo_32to8_kernel(uint8_t *d_ouput, uint32_t *d_temp, const int sz)
{
  int idx = blockDim.x * blockIdx.x + threadIdx.x;
  if (idx < sz)
  {
	if (d_temp[idx]< UINT8_MAX)
	  d_ouput[idx] = (uint8_t)d_temp[idx];
	if (d_temp[idx] >= UINT8_MAX)
	  d_ouput[idx] = (uint8_t)UINT_MAX;
  }
  __syncthreads();
}


void histo_cuda()
{
  cout << "Starting CUDA histo test..." << endl;

  int *input_array = (int*)malloc(INPUT_HEIGHT * INPUT_WIDTH * sizeof(int));;
  int gold_histo[HISTO_MAX] = { 0 };
  int kernel_histo[HISTO_MAX] = { 0 };

  for (int i = 0; i < INPUT_HEIGHT; i++)
	for (int j = 0; j < INPUT_WIDTH; j++)
	  input_array[i * INPUT_WIDTH+ j] = rand() % (HISTO_MAX - 1);

  for (int i = 0; i < 10; i++) {
	for (int j = 0; j < 14; j++) {
	  cout << input_array[i * INPUT_WIDTH + j] << " ";
	}
	cout << endl;
  }

  for (int i = 0; i < INPUT_HEIGHT; i++)
	for (int j = 0; j < INPUT_WIDTH; j++)
	  gold_histo[input_array[i * INPUT_WIDTH + j]]++;

  cout << "Golden result..." << endl;
  for (int i = 0; i < 20; i++)
	cout << gold_histo[i] << " ";
  cout << endl;

  int *d_input = NULL;
  cudaMalloc((void**)&d_input, INPUT_HEIGHT * INPUT_WIDTH * sizeof(int));

  cudaMemcpy(d_input, input_array, INPUT_HEIGHT * INPUT_WIDTH * 
	sizeof(int), cudaMemcpyHostToDevice);

  int *d_output = NULL;
  cudaMalloc((void**)&d_output, HISTO_MAX * sizeof(int));

  unsigned int num_blocks = ceil((float)INPUT_WIDTH / TILE_SIZE_MP7);

  histo_cuda_kernel << <num_blocks, 1024 >> > (d_input, d_output);

  cudaThreadSynchronize();

  cudaMemcpy(kernel_histo, d_output, HISTO_MAX * sizeof(int),
	cudaMemcpyDeviceToHost);

  cudaFree(d_input);
  cudaFree(d_output);

  cout << "Kernel result..." << endl;
  for (int i = 0; i < 20; i++)
	cout << kernel_histo[i] << " ";
  cout << endl;

  bool passed = true;
  cout << "Gold_Bins vs. Kernal_Bins" << endl;
  for (int i = 0; i < HISTO_MAX; i++) {
	if (gold_histo[i] != kernel_histo[i]) {
	  cout << i << ": " << gold_histo[i] << " vs. " << kernel_histo[i] << endl;
	  passed = false;
	  break;
	}
  }
  (passed) ? printf("\n    Test PASSED\n") : printf("\n    Test FAILED\n");

  cout << "Ending CUDA histo test..." << endl;
}
	.file	"tmpxft_0027212b_00000000-6_cuda_MP7.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL8next_binj, @function
_ZL8next_binj:
.LFB4132:
	.cfi_startproc
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	andl	$1023, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	call	rand@PLT
	cvtsi2sdl	%eax, %xmm0
	subsd	.LC0(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)
	call	rand@PLT
	movl	%eax, %ebp
	call	rand@PLT
	call	rand@PLT
	cvtsi2sdl	%ebp, %xmm0
	movl	$1023, %edx
	subsd	.LC0(%rip), %xmm0
	mulsd	8(%rsp), %xmm0
	movsd	.LC2(%rip), %xmm1
	mulsd	.LC1(%rip), %xmm0
	mulsd	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	addsd	%xmm0, %xmm0
	cvttsd2sil	%xmm0, %eax
	addl	%ebx, %eax
	cmpl	%edx, %eax
	cmovg	%edx, %eax
	xorl	%edx, %edx
	testl	%eax, %eax
	cmovs	%edx, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4132:
	.size	_ZL8next_binj, .-_ZL8next_binj
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4138:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4138:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4817:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4817:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z11ref_2dhistoPPjmmPh
	.type	_Z11ref_2dhistoPPjmmPh, @function
_Z11ref_2dhistoPPjmmPh:
.LFB4130:
	.cfi_startproc
	endbr64
	movq	%rdx, %r10
	movq	%rcx, %rdx
	xorl	%eax, %eax
	movq	%rdi, %r9
	movl	$256, %ecx
	movq	%rdx, %rdi
	movq	%rsi, %r8
	rep stosl
	xorl	%eax, %eax
.L6:
	cmpq	%r8, %rax
	je	.L15
	xorl	%ecx, %ecx
.L9:
	cmpq	%r10, %rcx
	je	.L16
	movq	(%r9,%rax,8), %rsi
	movl	(%rsi,%rcx,4), %esi
	addq	%rdx, %rsi
	movb	(%rsi), %dil
	cmpb	$-1, %dil
	je	.L7
	incl	%edi
	movb	%dil, (%rsi)
.L7:
	incq	%rcx
	jmp	.L9
.L16:
	incq	%rax
	jmp	.L6
.L15:
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE4130:
	.size	_Z11ref_2dhistoPPjmmPh, .-_Z11ref_2dhistoPPjmmPh
	.globl	_Z8alloc_2dmmm
	.type	_Z8alloc_2dmmm, @function
_Z8alloc_2dmmm:
.LFB4131:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	leaq	128(%rsi), %rbp
	movq	%rdx, %rsi
	andl	$4294967168, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rbp, %rdi
	imulq	%r12, %rdi
	call	calloc@PLT
	movl	$8, %esi
	movq	%r12, %rdi
	movq	%rax, %rbx
	call	calloc@PLT
	movq	%rax, %r13
	testq	%rbx, %rbx
	je	.L18
	testq	%rax, %rax
	je	.L18
	imulq	%r14, %rbp
	xorl	%eax, %eax
	jmp	.L19
.L18:
	movq	%rbx, %rdi
	call	free@PLT
	movq	%r13, %rdi
	xorl	%r13d, %r13d
	call	free@PLT
	jmp	.L17
.L19:
	cmpq	%r12, %rax
	je	.L17
	movq	%rbx, 0(%r13,%rax,8)
	incq	%rax
	addq	%rbp, %rbx
	jmp	.L19
.L20:
.L17:
	popq	%rbx
	.cfi_def_cfa_offset 40
	movq	%r13, %rax
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4131:
	.size	_Z8alloc_2dmmm, .-_Z8alloc_2dmmm
	.globl	_Z35__device_stub__Z12histo_kernelPPjS_PPjS_
	.type	_Z35__device_stub__Z12histo_kernelPPjS_PPjS_, @function
_Z35__device_stub__Z12histo_kernelPPjS_PPjS_:
.LFB4160:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z12histo_kernelPPjS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L23:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4160:
	.size	_Z35__device_stub__Z12histo_kernelPPjS_PPjS_, .-_Z35__device_stub__Z12histo_kernelPPjS_PPjS_
	.globl	_Z12histo_kernelPPjS_
	.type	_Z12histo_kernelPPjS_, @function
_Z12histo_kernelPPjS_:
.LFB4161:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z12histo_kernelPPjS_PPjS_
	.cfi_endproc
.LFE4161:
	.size	_Z12histo_kernelPPjS_, .-_Z12histo_kernelPPjS_
	.globl	_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_
	.type	_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_, @function
_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_:
.LFB4162:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17histo_cuda_kernelPiS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L28:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4162:
	.size	_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_, .-_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_
	.globl	_Z17histo_cuda_kernelPiS_
	.type	_Z17histo_cuda_kernelPiS_, @function
_Z17histo_cuda_kernelPiS_:
.LFB4163:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_
	.cfi_endproc
.LFE4163:
	.size	_Z17histo_cuda_kernelPiS_, .-_Z17histo_cuda_kernelPiS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"Starting CUDA histo test..."
.LC4:
	.string	" "
.LC5:
	.string	"Golden result..."
.LC6:
	.string	"Kernel result..."
.LC7:
	.string	"Gold_Bins vs. Kernal_Bins"
.LC8:
	.string	": "
.LC9:
	.string	" vs. "
.LC10:
	.string	"\n    Test FAILED\n"
.LC11:
	.string	"\n    Test PASSED\n"
.LC12:
	.string	"Ending CUDA histo test..."
	.text
	.globl	_Z10histo_cudav
	.type	_Z10histo_cudav, @function
_Z10histo_cudav:
.LFB4135:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC3(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	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	$2120, %rsp
	.cfi_def_cfa_offset 2176
	movq	%fs:40, %rax
	movq	%rax, 2104(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %r15
	leaq	1080(%rsp), %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$4079616, %edi
	call	malloc@PLT
	leaq	56(%rsp), %rdi
	movl	$256, %ecx
	xorl	%esi, %esi
	movq	%rax, %rbx
	xorl	%eax, %eax
	rep stosl
	leaq	1080(%rsp), %rdi
	movl	$256, %ecx
	rep stosl
.L34:
	imulq	$3984, %rsi, %rbp
	movq	%rbx, %r14
	xorl	%r13d, %r13d
	addq	%rbx, %rbp
.L35:
	movq	%rsi, 8(%rsp)
	call	rand@PLT
	movl	$255, %ecx
	movq	8(%rsp), %rsi
	cltd
	idivl	%ecx
	movl	%edx, 0(%rbp,%r13,4)
	incq	%r13
	cmpq	$996, %r13
	jne	.L35
	incq	%rsi
	cmpq	$1024, %rsi
	jne	.L34
	xorl	%ebp, %ebp
.L36:
	imulq	$3984, %rbp, %rcx
	xorl	%r13d, %r13d
	addq	%rbx, %rcx
.L37:
	movl	(%rcx,%r13,4), %esi
	leaq	_ZSt4cout(%rip), %rdi
	movq	%rcx, 8(%rsp)
	incq	%r13
	call	_ZNSolsEi@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$14, %r13
	movq	8(%rsp), %rcx
	jne	.L37
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%rbp
	call	_ZNSolsEPFRSoS_E.isra.0
	cmpq	$10, %rbp
	jne	.L36
	xorl	%edx, %edx
.L38:
	xorl	%eax, %eax
.L39:
	movslq	(%r14,%rax,4), %rcx
	incq	%rax
	incl	56(%rsp,%rcx,4)
	cmpq	$996, %rax
	jne	.L39
	addl	$996, %edx
	addq	$3984, %r14
	cmpl	$1019904, %edx
	jne	.L38
	leaq	.LC5(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%ebp, %ebp
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L41:
	movl	(%r15,%rbp,4), %esi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%rbp
	call	_ZNSolsEi@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$20, %rbp
	jne	.L41
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%eax, %eax
	leaq	16(%rsp), %rdi
	movl	$4079616, %esi
	movq	%rax, 16(%rsp)
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$4079616, %edx
	call	cudaMemcpy@PLT
	xorl	%edx, %edx
	leaq	24(%rsp), %rdi
	movl	$1024, %esi
	movq	%rdx, 24(%rsp)
	call	cudaMalloc@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$10, %rdx
	movl	$1, %ecx
	movabsq	$4294967421, %rdi
	movl	$1, %esi
	movq	%rdx, 44(%rsp)
	movl	$1, 52(%rsp)
	movq	%rdi, 32(%rsp)
	movl	$1, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L42
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	call	_Z39__device_stub__Z17histo_cuda_kernelPiS_PiS_
.L42:
	call	cudaThreadSynchronize@PLT
	movq	24(%rsp), %rsi
	movq	%r12, %rdi
	xorl	%ebx, %ebx
	movl	$2, %ecx
	movl	$1024, %edx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L43:
	movl	(%r12,%rbx,4), %esi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%rbx
	call	_ZNSolsEi@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$20, %rbx
	jne	.L43
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%ebx, %ebx
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC7(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L46:
	movl	(%r15,%rbx,4), %ebp
	cmpl	(%r12,%rbx,4), %ebp
	je	.L44
	movl	%ebx, %esi
	leaq	_ZSt4cout(%rip), %rdi
	movslq	%ebx, %rbx
	call	_ZNSolsEi@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC9(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	1080(%rsp,%rbx,4), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC10(%rip), %rsi
	jmp	.L57
.L44:
	incq	%rbx
	cmpq	$256, %rbx
	jne	.L46
	leaq	.LC11(%rip), %rsi
.L57:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC12(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	2104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L47
	call	__stack_chk_fail@PLT
.L47:
	addq	$2120, %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
.LFE4135:
	.size	_Z10histo_cudav, .-_Z10histo_cudav
	.globl	_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji
	.type	_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji, @function
_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji:
.LFB4164:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L58
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z18histo_32to8_kernelPhPji(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L58:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L60
	call	__stack_chk_fail@PLT
.L60:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4164:
	.size	_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji, .-_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji
	.globl	_Z18histo_32to8_kernelPhPji
	.type	_Z18histo_32to8_kernelPhPji, @function
_Z18histo_32to8_kernelPhPji:
.LFB4165:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji
	.cfi_endproc
.LFE4165:
	.size	_Z18histo_32to8_kernelPhPji, .-_Z18histo_32to8_kernelPhPji
	.globl	_Z11opt_2dhistoPPjPhS_mm
	.type	_Z11opt_2dhistoPPjPhS_mm, @function
_Z11opt_2dhistoPPjPhS_mm:
.LFB4134:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdx, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	testq	%r8, %r8
	js	.L64
	cvtsi2ssq	%r8, %xmm0
	jmp	.L65
.L64:
	movq	%r8, %rax
	andl	$1, %r8d
	shrq	%rax
	orq	%r8, %rax
	cvtsi2ssq	%rax, %xmm0
	addss	%xmm0, %xmm0
.L65:
	mulss	.LC13(%rip), %xmm0
	call	ceilf@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	cvttss2siq	%xmm0, %rax
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, 28(%rsp)
	movq	%rdx, 20(%rsp)
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L66
	movq	%rbx, %rsi
	movq	%r12, %rdi
	call	_Z35__device_stub__Z12histo_kernelPPjS_PPjS_
.L66:
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	salq	$10, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 28(%rsp)
	movq	%rdx, 20(%rsp)
	movq	%rdi, 8(%rsp)
	movl	$1, 16(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L67
	movl	$1024, %edx
	movq	%rbx, %rsi
	movq	%rbp, %rdi
	call	_Z41__device_stub__Z18histo_32to8_kernelPhPjiPhPji
.L67:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	cudaThreadSynchronize@PLT
	.cfi_endproc
.LFE4134:
	.size	_Z11opt_2dhistoPPjPhS_mm, .-_Z11opt_2dhistoPPjPhS_mm
	.section	.rodata.str1.1
.LC14:
	.string	"Input example: "
	.text
	.globl	_Z8cuda_MP7iPPc
	.type	_Z8cuda_MP7iPPc, @function
_Z8cuda_MP7iPPc:
.LFB4129:
	.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	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jg	.L70
	xorl	%edi, %edi
	jmp	.L89
.L70:
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movl	%eax, %edi
.L89:
	call	srand@PLT
	movl	$1024, %edi
	xorl	%r12d, %r12d
	call	malloc@PLT
	movl	$1024, %edi
	movq	%rax, %r13
	call	malloc@PLT
	movl	$4, %edx
	movl	$996, %esi
	movl	$1024, %edi
	movq	%rax, %rbp
	call	_Z8alloc_2dmmm
	movq	%rax, %rbx
	movq	(%rax), %rax
	movl	$512, (%rax)
.L72:
	movq	(%rbx), %rax
	movq	%r12, %rdx
	addq	$4, %r12
	movl	(%rax,%rdx), %edi
	leaq	(%rax,%r12), %r14
	call	_ZL8next_binj
	movl	%eax, (%r14)
	cmpq	$3980, %r12
	jne	.L72
	xorl	%r12d, %r12d
.L74:
	movq	(%rbx,%r12), %rax
	movq	8(%rbx,%r12), %r14
	movl	(%rax), %edi
	call	_ZL8next_binj
	movl	%eax, (%r14)
	xorl	%r14d, %r14d
.L73:
	movq	8(%rbx,%r12), %rax
	movq	%r14, %rdx
	addq	$4, %r14
	movl	(%rax,%rdx), %edi
	leaq	(%rax,%r14), %r15
	call	_ZL8next_binj
	movl	%eax, (%r15)
	cmpq	$3980, %r14
	jne	.L73
	addq	$8, %r12
	cmpq	$8184, %r12
	jne	.L74
	leaq	.LC14(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%r12d, %r12d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	.LC4(%rip), %r14
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L75:
	xorl	%r15d, %r15d
.L76:
	movq	(%rbx,%r12,8), %rax
	leaq	_ZSt4cout(%rip), %rdi
	movl	(%rax,%r15), %esi
	addq	$4, %r15
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$56, %r15
	jne	.L76
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%r12
	call	_ZNSolsEPFRSoS_E.isra.0
	cmpq	$10, %r12
	jne	.L75
	movq	%r13, %rcx
	movl	$996, %edx
	movq	%rbx, %rdi
	movl	$1024, %esi
	call	_Z11ref_2dhistoPPjmmPh
	xorl	%eax, %eax
	movq	%rsp, %rdi
	movl	$4079616, %esi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	xorl	%edx, %edx
	leaq	8(%rsp), %rdi
	movl	$1024, %esi
	movq	%rdx, 8(%rsp)
	call	cudaMalloc@PLT
	xorl	%ecx, %ecx
	leaq	16(%rsp), %rdi
	movl	$4096, %esi
	movq	%rcx, 16(%rsp)
	call	cudaMalloc@PLT
	movq	(%rbx), %rsi
	movq	(%rsp), %rdi
	xorl	%ebx, %ebx
	movl	$1, %ecx
	movl	$4079616, %edx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$996, %r8d
	movq	(%rsp), %rdi
	movl	$1024, %ecx
	call	_Z11opt_2dhistoPPjPhS_mm
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbp, %rdi
	movl	$1024, %edx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	leaq	.LC7(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L80:
	movb	0(%r13,%rbx), %r15b
	cmpb	0(%rbp,%rbx), %r15b
	je	.L78
	movl	%ebx, %esi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEi@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsbl	%r15b, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsbl	0(%rbp,%rbx), %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC10(%rip), %rsi
	jmp	.L90
.L78:
	incq	%rbx
	cmpq	$1024, %rbx
	jne	.L80
	leaq	.LC11(%rip), %rsi
.L90:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L81
	call	__stack_chk_fail@PLT
.L81:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rdi
	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
	jmp	free@PLT
	.cfi_endproc
.LFE4129:
	.size	_Z8cuda_MP7iPPc, .-_Z8cuda_MP7iPPc
	.section	.rodata.str1.1
.LC15:
	.string	"_Z18histo_32to8_kernelPhPji"
.LC16:
	.string	"_Z17histo_cuda_kernelPiS_"
.LC17:
	.string	"_Z12histo_kernelPPjS_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4167:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC15(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z18histo_32to8_kernelPhPji(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC16(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z17histo_cuda_kernelPiS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC17(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z12histo_kernelPPjS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4167:
	.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
.LC0:
	.long	0
	.long	1071644672
	.align 8
.LC1:
	.long	0
	.long	1074790400
	.align 8
.LC2:
	.long	1719614413
	.long	1073127582
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC13:
	.long	1040187392
	.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:

#ifndef CUDA_MP7_CUH
#define CUDA_MP7_CUH

#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#include <string>
#include <vector>
#include <iostream>
#include <fstream>
#include <time.h>
//#include <sys/time.h>
//#include <Windows.h>

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



using namespace std;

//#define __USE_BSD 1
#define INPUT_WIDTH  996
#define INPUT_HEIGHT 1024

#define HISTO_WIDTH  1024
#define HISTO_HEIGHT 1
#define HISTO_LOG 10

#define UINT8_MAX 255
#define HISTO_MAX 256

#define TILE_SIZE_MP7 8 

void cuda_MP7(int argc, char* argv[]);
int ref_2dhisto(uint32_t *input[], size_t height, size_t width, uint8_t bins[]);
void** alloc_2d(size_t y_size, size_t x_size, size_t element_size);
static uint32_t next_bin(uint32_t pix);
//int gettimeofday(struct timeval * tp, struct timezone * tzp);
static uint32_t **generate_histogram_bins();
void opt_2dhisto(uint32_t **d_input, uint8_t *d_bins, uint32_t *d_temp_bins,
  size_t height, size_t width);

__global__ void histo_kernel(uint32_t **d_input, uint32_t *d_ouput);
__global__ void histo_cuda_kernel(int *d_input, int *d_output);
__global__ void histo_32to8_kernel(uint8_t *d_ouput, uint32_t *d_temp, const int sz);

void histo_cuda();

/*
#define TIME_IT(ROUTINE_NAME__, LOOPS__, ACTION__)\
{\
    printf("    Timing '%s' started\n", ROUTINE_NAME__);\
    struct timeval tv;\
    struct timezone tz;\
    const clock_t startTime = clock();\
    gettimeofday(&tv, &tz); long GTODStartTime =  tv.tv_sec * 1000 + tv.tv_usec / 1000 ;\
    for (int loops = 0; loops < (LOOPS__); ++loops)\
    {\
        ACTION__;\
    }\
    gettimeofday(&tv, &tz); long GTODEndTime =  tv.tv_sec * 1000 + tv.tv_usec / 1000 ;\
    const clock_t endTime = clock();\
    const clock_t elapsedTime = endTime - startTime;\
    const double timeInSeconds = (elapsedTime/(double)CLOCKS_PER_SEC);\
    printf("        GetTimeOfDay Time (for %d iterations) = %g\n", LOOPS__, (double)(GTODEndTime - GTODStartTime) / 1000. );\
    printf("        Clock Time        (for %d iterations) = %g\n", LOOPS__, timeInSeconds );\
    printf("    Timing '%s' ended\n", ROUTINE_NAME__);\
}
*/

#define SQRT_2    1.4142135623730950488
#define SPREAD_BOTTOM   (2)
#define SPREAD_TOP      (6)

#define NEXT(init_, spread_)\
    (init_ + (int)((rand() - 0.5) * (rand() - 0.5) * 4.0 * SQRT_2 * SQRT_2 * spread_));

#define CLAMP(value_, min_, max_)\
    if (value_ < 0)\
        value_ = (min_);\
    else if (value_ > (max_))\
        value_ = (max_);


#endif // !CUDA_MP7_CUH
	.text
	.file	"cuda_MP7.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

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

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

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

	.end_amdgpu_metadata
