#include <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <cuda.h>

//#define SHARED_SIZE_LIMIT 1024
#define NUM_THREADS 1024
#define NUM_BLOCKS 32768
#define NUM_VALS NUM_THREADS*NUM_BLOCKS
#define SHARED_SIZE_LIMIT 1024


int random_float() {
  return (int)rand()/(int)RAND_MAX;
}

void array_print(int *arr, int length)  {
  int i;
  for (i = 0; i < length; ++i) {
    printf("%d ",  arr[i]);
  }
  printf("\n");
}

void array_fill(int *v) {
   int i;
   for (i = 0; i < NUM_VALS; i++) {
     v[i] = rand();
   }
}

void test (int *v) {
  int i;
  int val = v[0];
  for (i = 1; i < NUM_VALS; ++i) {
    if (val < v[i]) {
        printf("val: %d, v[%d]: %d.\n", val, i, v[i]);
        printf("TEST FAIL\n\n");
        return;
    } else {
        printf("val: %d, v[%d]: %d.\n", val, i, v[i]);

        val = v[i];
    }
  }

  printf("TEST OK\n\n");

}
   

/*
void array_fill(int *arr, int length) {
  srand(time(NULL));
  int i;
  for (i = 0; i < length; ++i) {
    arr[i] = length-i;//random_float();
  }
}*/

void array_copy(int *dst, int *src, int length) {
  int i;
  for (i=0; i<length; ++i) {
    dst[i] = src[i];
  }
}

//Comparamos dos elementos y en caso de ser decrecientes, los swapeamos.
__device__ inline void comparator(int &A, int &B, uint dir) {
    int temp;
    if ((A <= B) == dir) {
        temp = A;
        A = B;
        B = temp;
    }
}

/*La cosa en este bitonicsort es que compartimos memoria.
Asíues pese que la idea principal es la misma, nosotros lo que hacemos es 
comparaciones entre elementos de las distintas memorias para hacer que,
finalmente el vector termine ordenado.
*/

__global__ void bitonicSortShared(int *dev_values) 
{
    int tx = threadIdx.x;
    int bx = blockIdx.x;
    int index = blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    __shared__ int sh_values[SHARED_SIZE_LIMIT];
    sh_values[tx] = dev_values[index];
    sh_values[tx + (SHARED_SIZE_LIMIT/2)] = dev_values[index + (SHARED_SIZE_LIMIT/2)];
    for (uint size = 2; size < SHARED_SIZE_LIMIT; size <<= 1) {
      uint ddd = (tx & (size / 2)) == 0;//direction: ascending or descending
      for (uint stride = size/2; stride > 0; stride >>= 1) {
          __syncthreads();
          uint pos = 2 * tx - (tx & (stride - 1));
          comparator(sh_values[pos], sh_values[pos + stride], ddd);
      }
    }
    uint ddd = ((bx&1) == 0); //    uint ddd = ((bx&1)==0);
    {
      for (uint stride = SHARED_SIZE_LIMIT/2; stride > 0; stride >>= 1) {
          __syncthreads();
          uint pos = 2 * tx - (tx & (stride - 1));
          comparator(sh_values[pos + 0], sh_values[pos + stride], ddd);
      }
    }
    __syncthreads();
    dev_values[index] = sh_values[tx];
    dev_values[index+(SHARED_SIZE_LIMIT/2)] = sh_values[tx+(SHARED_SIZE_LIMIT/2)];
}

void bitonic_sort(int *values) 
{
  int *dev_values;
  size_t size = NUM_VALS * sizeof(int);
  cudaMalloc((void**) &dev_values, size);
  cudaMemcpy(dev_values, values, size, cudaMemcpyHostToDevice);
  dim3 numBlocks(NUM_BLOCKS, 1);
  dim3 numThreads(NUM_THREADS, 1);
  cudaDeviceSynchronize();
  uint blockCount = NUM_VALS / SHARED_SIZE_LIMIT;
  uint threadCount = SHARED_SIZE_LIMIT / 2;
  printf("blockCount=%d, threadCount=%d, SHARED_SIZE_LIMIT=%d\n", blockCount, threadCount, SHARED_SIZE_LIMIT);
  bitonicSortShared<<<blockCount, threadCount>>>(dev_values);
  cudaDeviceSynchronize();
  cudaMemcpy(values, dev_values, size, cudaMemcpyDeviceToHost);
  cudaFree(dev_values);
}

int main(void) 
{
  //int *values = (int*) malloc( NUM_VALS * sizeof(int));
  //int *ref = (int*) malloc( NUM_VALS * sizeof(int));
  int *host_values;
  cudaMallocHost( &host_values, NUM_VALS * sizeof(int));
//  cudaMallocHost( &original_values, numBytes);

  float TiempoKernel; 

  cudaEvent_t E1, E2;

  cudaEventCreate(&E1);
  cudaEventCreate(&E2);

  array_fill(host_values);	 

  cudaEventRecord(E1, 0);
  cudaEventSynchronize(E1);

  cudaFuncSetCacheConfig(bitonicSortShared, cudaFuncCachePreferL1);

  bitonic_sort(host_values);
  test(host_values);
  cudaEventRecord(E2, 0);
  cudaEventSynchronize(E2);

  cudaEventElapsedTime(&TiempoKernel, E1, E2);

  printf("Tiempo Kernel: %4.6f milseg\n", TiempoKernel);

  cudaFree(host_values);
}

	.file	"tmpxft_002486bf_00000000-6_multiGPU.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2037:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2037:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z12random_floatv
	.type	_Z12random_floatv, @function
_Z12random_floatv:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	rand@PLT
	movl	$2147483647, %ecx
	cltd
	idivl	%ecx
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z12random_floatv, .-_Z12random_floatv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%d "
.LC1:
	.string	"\n"
	.text
	.globl	_Z11array_printPii
	.type	_Z11array_printPii, @function
_Z11array_printPii:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC0(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 48
.L5:
	cmpl	%ebx, %ebp
	jle	.L9
	movl	(%r12,%rbx,4), %edx
	movq	%r13, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%rbx
	call	__printf_chk@PLT
	jmp	.L5
.L9:
	popq	%rdx
	.cfi_def_cfa_offset 40
	leaq	.LC1(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z11array_printPii, .-_Z11array_printPii
	.globl	_Z10array_fillPi
	.type	_Z10array_fillPi, @function
_Z10array_fillPi:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	xorl	%ebx, %ebx
	pushq	%rdx
	.cfi_def_cfa_offset 32
.L11:
	call	rand@PLT
	movl	%eax, 0(%rbp,%rbx,4)
	incq	%rbx
	cmpq	$33554432, %rbx
	jne	.L11
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z10array_fillPi, .-_Z10array_fillPi
	.section	.rodata.str1.1
.LC2:
	.string	"val: %d, v[%d]: %d.\n"
.LC3:
	.string	"TEST FAIL\n\n"
.LC4:
	.string	"TEST OK\n\n"
	.text
	.globl	_Z4testPi
	.type	_Z4testPi, @function
_Z4testPi:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	.LC2(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	(%rdi), %edx
	movl	$1, %ebx
.L16:
	movl	0(%rbp,%rbx,4), %r8d
	movl	%ebx, %ecx
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	cmpl	%edx, %r8d
	jle	.L15
	call	__printf_chk@PLT
	leaq	.LC3(%rip), %rsi
	jmp	.L19
.L15:
	call	__printf_chk@PLT
	movl	0(%rbp,%rbx,4), %edx
	incq	%rbx
	cmpq	$33554432, %rbx
	jne	.L16
	leaq	.LC4(%rip), %rsi
.L19:
	popq	%rbx
	.cfi_def_cfa_offset 24
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z4testPi, .-_Z4testPi
	.globl	_Z10array_copyPiS_i
	.type	_Z10array_copyPiS_i, @function
_Z10array_copyPiS_i:
.LFB2031:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L21:
	cmpl	%eax, %edx
	jle	.L23
	movl	(%rsi,%rax,4), %ecx
	movl	%ecx, (%rdi,%rax,4)
	incq	%rax
	jmp	.L21
.L23:
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z10array_copyPiS_i, .-_Z10array_copyPiS_i
	.globl	_Z37__device_stub__Z17bitonicSortSharedPiPi
	.type	_Z37__device_stub__Z17bitonicSortSharedPiPi, @function
_Z37__device_stub__Z17bitonicSortSharedPiPi:
.LFB2059:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z17bitonicSortSharedPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L24:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z37__device_stub__Z17bitonicSortSharedPiPi, .-_Z37__device_stub__Z17bitonicSortSharedPiPi
	.globl	_Z17bitonicSortSharedPi
	.type	_Z17bitonicSortSharedPi, @function
_Z17bitonicSortSharedPi:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z17bitonicSortSharedPiPi
	.cfi_endproc
.LFE2060:
	.size	_Z17bitonicSortSharedPi, .-_Z17bitonicSortSharedPi
	.section	.rodata.str1.1
.LC5:
	.string	"blockCount=%d, threadCount=%d, SHARED_SIZE_LIMIT=%d\n"
	.text
	.globl	_Z12bitonic_sortPi
	.type	_Z12bitonic_sortPi, @function
_Z12bitonic_sortPi:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$134217728, %esi
	movq	%rdi, %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$134217728, %edx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movl	$1024, %r8d
	xorl	%eax, %eax
	movl	$512, %ecx
	movl	$32768, %edx
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$8388609, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$131073, %edi
	salq	$9, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$15, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	movq	8(%rsp), %rdi
	call	_Z37__device_stub__Z17bitonicSortSharedPiPi
.L30:
	call	cudaDeviceSynchronize@PLT
	movq	8(%rsp), %rsi
	movq	%rbx, %rdi
	movl	$2, %ecx
	movl	$134217728, %edx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$48, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z12bitonic_sortPi, .-_Z12bitonic_sortPi
	.section	.rodata.str1.1
.LC6:
	.string	"Tiempo Kernel: %4.6f milseg\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	xorl	%edx, %edx
	movl	$134217728, %esi
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	16(%rsp), %rdi
	call	cudaHostAlloc@PLT
	leaq	24(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	32(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	16(%rsp), %rdi
	call	_Z10array_fillPi
	movq	24(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	24(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movl	$2, %esi
	leaq	_Z17bitonicSortSharedPi(%rip), %rdi
	call	cudaFuncSetCacheConfig@PLT
	movq	16(%rsp), %rdi
	call	_Z12bitonic_sortPi
	movq	16(%rsp), %rdi
	call	_Z4testPi
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	32(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rsi
	leaq	12(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movl	$2, %edi
	movb	$1, %al
	leaq	.LC6(%rip), %rsi
	cvtss2sd	12(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L34
	call	__stack_chk_fail@PLT
.L34:
	xorl	%eax, %eax
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z17bitonicSortSharedPi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2062:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z17bitonicSortSharedPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2062:
	.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
	.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 <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <hip/hip_runtime.h>

//#define SHARED_SIZE_LIMIT 1024
#define NUM_THREADS 1024
#define NUM_BLOCKS 32768
#define NUM_VALS NUM_THREADS*NUM_BLOCKS
#define SHARED_SIZE_LIMIT 1024


int random_float() {
  return (int)rand()/(int)RAND_MAX;
}

void array_print(int *arr, int length)  {
  int i;
  for (i = 0; i < length; ++i) {
    printf("%d ",  arr[i]);
  }
  printf("\n");
}

void array_fill(int *v) {
   int i;
   for (i = 0; i < NUM_VALS; i++) {
     v[i] = rand();
   }
}

void test (int *v) {
  int i;
  int val = v[0];
  for (i = 1; i < NUM_VALS; ++i) {
    if (val < v[i]) {
        printf("val: %d, v[%d]: %d.\n", val, i, v[i]);
        printf("TEST FAIL\n\n");
        return;
    } else {
        printf("val: %d, v[%d]: %d.\n", val, i, v[i]);

        val = v[i];
    }
  }

  printf("TEST OK\n\n");

}
   

/*
void array_fill(int *arr, int length) {
  srand(time(NULL));
  int i;
  for (i = 0; i < length; ++i) {
    arr[i] = length-i;//random_float();
  }
}*/

void array_copy(int *dst, int *src, int length) {
  int i;
  for (i=0; i<length; ++i) {
    dst[i] = src[i];
  }
}

//Comparamos dos elementos y en caso de ser decrecientes, los swapeamos.
__device__ inline void comparator(int &A, int &B, uint dir) {
    int temp;
    if ((A <= B) == dir) {
        temp = A;
        A = B;
        B = temp;
    }
}

/*La cosa en este bitonicsort es que compartimos memoria.
Asíues pese que la idea principal es la misma, nosotros lo que hacemos es 
comparaciones entre elementos de las distintas memorias para hacer que,
finalmente el vector termine ordenado.
*/

__global__ void bitonicSortShared(int *dev_values) 
{
    int tx = threadIdx.x;
    int bx = blockIdx.x;
    int index = blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    __shared__ int sh_values[SHARED_SIZE_LIMIT];
    sh_values[tx] = dev_values[index];
    sh_values[tx + (SHARED_SIZE_LIMIT/2)] = dev_values[index + (SHARED_SIZE_LIMIT/2)];
    for (uint size = 2; size < SHARED_SIZE_LIMIT; size <<= 1) {
      uint ddd = (tx & (size / 2)) == 0;//direction: ascending or descending
      for (uint stride = size/2; stride > 0; stride >>= 1) {
          __syncthreads();
          uint pos = 2 * tx - (tx & (stride - 1));
          comparator(sh_values[pos], sh_values[pos + stride], ddd);
      }
    }
    uint ddd = ((bx&1) == 0); //    uint ddd = ((bx&1)==0);
    {
      for (uint stride = SHARED_SIZE_LIMIT/2; stride > 0; stride >>= 1) {
          __syncthreads();
          uint pos = 2 * tx - (tx & (stride - 1));
          comparator(sh_values[pos + 0], sh_values[pos + stride], ddd);
      }
    }
    __syncthreads();
    dev_values[index] = sh_values[tx];
    dev_values[index+(SHARED_SIZE_LIMIT/2)] = sh_values[tx+(SHARED_SIZE_LIMIT/2)];
}

void bitonic_sort(int *values) 
{
  int *dev_values;
  size_t size = NUM_VALS * sizeof(int);
  hipMalloc((void**) &dev_values, size);
  hipMemcpy(dev_values, values, size, hipMemcpyHostToDevice);
  dim3 numBlocks(NUM_BLOCKS, 1);
  dim3 numThreads(NUM_THREADS, 1);
  hipDeviceSynchronize();
  uint blockCount = NUM_VALS / SHARED_SIZE_LIMIT;
  uint threadCount = SHARED_SIZE_LIMIT / 2;
  printf("blockCount=%d, threadCount=%d, SHARED_SIZE_LIMIT=%d\n", blockCount, threadCount, SHARED_SIZE_LIMIT);
  bitonicSortShared<<<blockCount, threadCount>>>(dev_values);
  hipDeviceSynchronize();
  hipMemcpy(values, dev_values, size, hipMemcpyDeviceToHost);
  hipFree(dev_values);
}

int main(void) 
{
  //int *values = (int*) malloc( NUM_VALS * sizeof(int));
  //int *ref = (int*) malloc( NUM_VALS * sizeof(int));
  int *host_values;
  hipHostMalloc( &host_values, NUM_VALS * sizeof(int), hipHostMallocDefault);
//  cudaMallocHost( &original_values, numBytes);

  float TiempoKernel; 

  hipEvent_t E1, E2;

  hipEventCreate(&E1);
  hipEventCreate(&E2);

  array_fill(host_values);	 

  hipEventRecord(E1, 0);
  hipEventSynchronize(E1);

  hipFuncSetCacheConfig(reinterpret_cast<const void*>(bitonicSortShared), hipFuncCachePreferL1);

  bitonic_sort(host_values);
  test(host_values);
  hipEventRecord(E2, 0);
  hipEventSynchronize(E2);

  hipEventElapsedTime(&TiempoKernel, E1, E2);

  printf("Tiempo Kernel: %4.6f milseg\n", TiempoKernel);

  hipFree(host_values);
}

	.text
	.file	"multiGPU.hip"
	.globl	_Z12random_floatv               # -- Begin function _Z12random_floatv
	.type	_Z12random_floatv,@function
_Z12random_floatv:                      # @_Z12random_floatv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	rand
	movslq	%eax, %rcx
	movq	%rcx, %rax
	shlq	$30, %rax
	addq	%rcx, %rax
	movq	%rax, %rcx
	shrq	$63, %rcx
	sarq	$61, %rax
	addl	%ecx, %eax
                                        # kill: def $eax killed $eax killed $rax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z12random_floatv, .Lfunc_end0-_Z12random_floatv
	.cfi_endproc
                                        # -- End function
	.globl	_Z11array_printPii              # -- Begin function _Z11array_printPii
	.type	_Z11array_printPii,@function
_Z11array_printPii:                     # @_Z11array_printPii
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB1_4
# %bb.1:                                # %.lr.ph.preheader
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdi, %rbx
	movl	%esi, %r14d
	xorl	%r15d, %r15d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r15,4), %esi
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB1_2
# %bb.3:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB1_4:                                # %._crit_edge
	movl	$10, %edi
	jmp	putchar@PLT                     # TAILCALL
.Lfunc_end1:
	.size	_Z11array_printPii, .Lfunc_end1-_Z11array_printPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z10array_fillPi                # -- Begin function _Z10array_fillPi
	.type	_Z10array_fillPi,@function
_Z10array_fillPi:                       # @_Z10array_fillPi
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	xorl	%r14d, %r14d
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	movl	%eax, (%rbx,%r14,4)
	incq	%r14
	cmpq	$33554432, %r14                 # imm = 0x2000000
	jne	.LBB2_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z10array_fillPi, .Lfunc_end2-_Z10array_fillPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z4testPi                       # -- Begin function _Z4testPi
	.type	_Z4testPi,@function
_Z4testPi:                              # @_Z4testPi
	.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	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %rbx
	movl	$1, %r14d
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	movl	-4(%rbx,%r14,4), %ebp
	movl	(%rbx,%r14,4), %r15d
	movl	$.L.str.2, %edi
	movl	%ebp, %esi
	movl	%r14d, %edx
	movl	%r15d, %ecx
	xorl	%eax, %eax
	callq	printf
	cmpl	%r15d, %ebp
	jl	.LBB3_2
# %bb.3:                                #   in Loop: Header=BB3_1 Depth=1
	incq	%r14
	cmpq	$33554432, %r14                 # imm = 0x2000000
	jne	.LBB3_1
# %bb.4:
	movl	$.Lstr, %edi
	jmp	.LBB3_5
.LBB3_2:
	movl	$.Lstr.1, %edi
.LBB3_5:                                # %.loopexit
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end3:
	.size	_Z4testPi, .Lfunc_end3-_Z4testPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z10array_copyPiS_i             # -- Begin function _Z10array_copyPiS_i
	.type	_Z10array_copyPiS_i,@function
_Z10array_copyPiS_i:                    # @_Z10array_copyPiS_i
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB4_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%edx, %eax
	xorl	%ecx, %ecx
.LBB4_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rsi,%rcx,4), %edx
	movl	%edx, (%rdi,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB4_2
.LBB4_3:                                # %._crit_edge
	retq
.Lfunc_end4:
	.size	_Z10array_copyPiS_i, .Lfunc_end4-_Z10array_copyPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__bitonicSortSharedPi # -- Begin function _Z32__device_stub__bitonicSortSharedPi
	.type	_Z32__device_stub__bitonicSortSharedPi,@function
_Z32__device_stub__bitonicSortSharedPi: # @_Z32__device_stub__bitonicSortSharedPi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z17bitonicSortSharedPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z32__device_stub__bitonicSortSharedPi, .Lfunc_end5-_Z32__device_stub__bitonicSortSharedPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z12bitonic_sortPi              # -- Begin function _Z12bitonic_sortPi
	.type	_Z12bitonic_sortPi,@function
_Z12bitonic_sortPi:                     # @_Z12bitonic_sortPi
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	movq	%rsp, %r14
	movl	$134217728, %esi                # imm = 0x8000000
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$134217728, %edx                # imm = 0x8000000
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	callq	hipDeviceSynchronize
	movl	$.L.str.5, %edi
	movl	$32768, %esi                    # imm = 0x8000
	movl	$512, %edx                      # imm = 0x200
	movl	$1024, %ecx                     # imm = 0x400
	xorl	%eax, %eax
	callq	printf
	movabsq	$4294967808, %rdx               # imm = 0x100000200
	leaq	32256(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_2
# %bb.1:
	movq	(%rsp), %rdi
	callq	_Z32__device_stub__bitonicSortSharedPi
.LBB6_2:
	callq	hipDeviceSynchronize
	movq	(%rsp), %rsi
	movl	$134217728, %edx                # imm = 0x8000000
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z12bitonic_sortPi, .Lfunc_end6-_Z12bitonic_sortPi
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	24(%rsp), %rbx
	xorl	%r14d, %r14d
	movl	$134217728, %esi                # imm = 0x8000000
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	leaq	16(%rsp), %rdi
	callq	hipEventCreate
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	movq	(%rbx), %rbx
.LBB7_1:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	movl	%eax, (%rbx,%r14,4)
	incq	%r14
	cmpq	$33554432, %r14                 # imm = 0x2000000
	jne	.LBB7_1
# %bb.2:                                # %_Z10array_fillPi.exit
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rdi
	callq	hipEventSynchronize
	movl	$_Z17bitonicSortSharedPi, %edi
	movl	$2, %esi
	callq	hipFuncSetCacheConfig
	movq	24(%rsp), %rdi
	callq	_Z12bitonic_sortPi
	movq	24(%rsp), %rdi
	callq	_Z4testPi
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	8(%rsp), %rdi
	callq	hipEventSynchronize
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	leaq	36(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventElapsedTime
	cvtss2sd	(%rbx), %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	movq	24(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	main, .Lfunc_end7-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17bitonicSortSharedPi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB9_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%d "
	.size	.L.str, 4

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"val: %d, v[%d]: %d.\n"
	.size	.L.str.2, 21

	.type	_Z17bitonicSortSharedPi,@object # @_Z17bitonicSortSharedPi
	.section	.rodata,"a",@progbits
	.globl	_Z17bitonicSortSharedPi
	.p2align	3, 0x0
_Z17bitonicSortSharedPi:
	.quad	_Z32__device_stub__bitonicSortSharedPi
	.size	_Z17bitonicSortSharedPi, 8

	.type	.L.str.5,@object                # @.str.5
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.5:
	.asciz	"blockCount=%d, threadCount=%d, SHARED_SIZE_LIMIT=%d\n"
	.size	.L.str.5, 53

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Tiempo Kernel: %4.6f milseg\n"
	.size	.L.str.6, 29

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z17bitonicSortSharedPi"
	.size	.L__unnamed_1, 24

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

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"TEST OK\n"
	.size	.Lstr, 9

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"TEST FAIL\n"
	.size	.Lstr.1, 11

	.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 _Z32__device_stub__bitonicSortSharedPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z17bitonicSortSharedPi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17bitonicSortSharedPi ; -- Begin function _Z17bitonicSortSharedPi
	.globl	_Z17bitonicSortSharedPi
	.p2align	8
	.type	_Z17bitonicSortSharedPi,@function
_Z17bitonicSortSharedPi:                ; @_Z17bitonicSortSharedPi
; %bb.0:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_lshl_or_b32 v1, s15, 10, v0
	v_lshlrev_b32_e32 v3, 2, v0
	v_lshlrev_b32_e32 v4, 1, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v2, 31, v1
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, s0, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s1, v2, vcc_lo
	s_mov_b32 s1, 2
	s_clause 0x1
	global_load_b32 v5, v[1:2], off
	global_load_b32 v6, v[1:2], off offset:2048
	s_waitcnt vmcnt(0)
	ds_store_2addr_stride64_b32 v3, v5, v6 offset1:8
.LBB0_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_2 Depth 2
	s_lshr_b32 s2, s1, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v5, s2, v0
	v_cmp_ne_u32_e32 vcc_lo, 0, v5
.LBB0_2:                                ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_add_i32 s0, s2, -1
	s_waitcnt lgkmcnt(0)
	v_and_b32_e32 v5, s0, v0
	s_barrier
	buffer_gl0_inv
	v_sub_nc_u32_e32 v5, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v5, 2, v5
	v_lshl_add_u32 v6, s2, 2, v5
	ds_load_b32 v7, v5
	ds_load_b32 v8, v6
	s_waitcnt lgkmcnt(0)
	v_cmp_le_i32_e64 s0, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s3, vcc_lo, s0
	s_and_saveexec_b32 s0, s3
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=2
	ds_store_b32 v5, v8
	ds_store_b32 v6, v7
.LBB0_4:                                ; %_Z10comparatorRiS_j.exit
                                        ;   in Loop: Header=BB0_2 Depth=2
	s_or_b32 exec_lo, exec_lo, s0
	s_lshr_b32 s0, s2, 1
	s_cmp_lt_u32 s2, 2
	s_mov_b32 s2, s0
	s_cbranch_scc0 .LBB0_2
; %bb.5:                                ;   in Loop: Header=BB0_1 Depth=1
	s_lshl_b32 s0, s1, 1
	s_cmpk_gt_u32 s1, 0x1ff
	s_mov_b32 s1, s0
	s_cbranch_scc0 .LBB0_1
; %bb.6:                                ; %.preheader
	s_bitcmp0_b32 s15, 0
	s_movk_i32 s1, 0x200
	s_cselect_b32 s0, -1, 0
.LBB0_7:                                ; =>This Inner Loop Header: Depth=1
	s_add_i32 s2, s1, -1
	s_waitcnt lgkmcnt(0)
	v_and_b32_e32 v5, s2, v0
	s_barrier
	buffer_gl0_inv
	v_sub_nc_u32_e32 v5, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v5, 2, v5
	v_lshl_add_u32 v6, s1, 2, v5
	ds_load_b32 v7, v5
	ds_load_b32 v8, v6
	s_waitcnt lgkmcnt(0)
	v_cmp_le_i32_e32 vcc_lo, v7, v8
	s_xor_b32 s2, s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s3, s2, -1
	s_and_saveexec_b32 s2, s3
	s_cbranch_execz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_7 Depth=1
	ds_store_b32 v5, v8
	ds_store_b32 v6, v7
.LBB0_9:                                ; %_Z10comparatorRiS_j.exit39
                                        ;   in Loop: Header=BB0_7 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_lshr_b32 s2, s1, 1
	s_cmp_lt_u32 s1, 2
	s_mov_b32 s1, s2
	s_cbranch_scc0 .LBB0_7
; %bb.10:
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	ds_load_2addr_stride64_b32 v[3:4], v3 offset1:8
	s_waitcnt lgkmcnt(0)
	s_clause 0x1
	global_store_b32 v[1:2], v3, off
	global_store_b32 v[1:2], v4, off offset:2048
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17bitonicSortSharedPi
		.amdhsa_group_segment_fixed_size 4096
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 8
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 9
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z17bitonicSortSharedPi, .Lfunc_end0-_Z17bitonicSortSharedPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 448
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4096 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 4096
    .kernarg_segment_align: 8
    .kernarg_segment_size: 8
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17bitonicSortSharedPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17bitonicSortSharedPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
