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

typedef struct node {
  int data;
  struct node *parent;
  struct node *left;
  struct node *right;
  int height;
  int sema;
} node;

__device__ node* global_root = NULL;

__device__ volatile int MASTER_LOCK = 0;

__device__ int lock(node* n) {
  int status = atomicExch(&n->sema, 1);
  return (!status && !MASTER_LOCK);
}

__device__ void unlock(node* n) {
  atomicExch(&n->sema, 0);
}

__device__ node* new_node(int val, node* parent) {
  node *tmp = (node *) malloc(sizeof(node));
  tmp->data = val;
  tmp->parent = parent;
  tmp->left = tmp->right = NULL;
  tmp->height = 1;
  tmp->sema = 0;
  return tmp;
}

__device__ node* find(node* root, int key) {
  if (root == NULL) return NULL;

  if (root->data == key) return root;
  else if (root->data > key) return find(root->left, key);
  else return find(root->right, key);
}

__device__ int height(node *root)
{
  if (root == NULL)
      return 0;
  return root->height;
}

__device__ int get_balance(node *root)
{
  if (root == NULL)
      return 0;
  return height(root->left) - height(root->right);
}


__device__ node* left_rotate(node* root, node* parent)
{
  node* temp1 = root->right;
  node* temp2 = temp1->left;

  temp1->left = root;
  root->parent = temp1;
  root->right = temp2;

  if (temp2)
    temp2->parent = root;

  root->height = max(height(root->left), height(root->right))+1;
  temp1->height = max(height(temp1->left), height(temp1->right))+1;

  temp1->parent = parent;
  return temp1;
}

__device__ node* right_rotate(node* root, node* parent)
{
  node* temp1 = root->left;
  node* temp2 = temp1->right;

  temp1->right = root;
  root->parent = temp1;
  root->left = temp2;

  if(temp2)
    temp2->parent = root;

  root->height = max(height(root->left), height(root->right))+1;
  temp1->height = max(height(temp1->left), height(temp1->right))+1;

  temp1->parent = parent;
  return temp1;
}


__device__ void rebalance(node* root, int key) {
  root->height = max(height(root->left),height(root->right)) + 1;
  int balance = get_balance(root);

  // Left Left Case
  node* p = root->parent;
  if (balance > 1 && key < root->left->data) {
    if (p) {
      if (root->data < p->data)
        p->left = right_rotate(root, p);
      else
        p->right = right_rotate(root, p);
    }
    else 
      global_root = right_rotate(root, global_root);
  }

  // Right Right Case
  else if (balance < -1 && key > root->right->data) {
    if (p) {
      if (root->data < p->data) 
        p->left = left_rotate(root, p);
      else
        p->right = left_rotate(root, p);
    }
    else
      global_root = left_rotate(root, global_root);
  }

  // Left Right Case
  else if (balance > 1 && key > root->left->data) {
    root->left =  left_rotate(root->left, root);
    if (p) {
      if (root->data < p->data)
        p->left = right_rotate(root, p);
      else 
        p->right = right_rotate(root, p);
    }
    else 
      global_root = right_rotate(root, global_root);
  }

  // Right Left Case
  else if (balance < -1 && key < root->right->data)
  {
    root->right = right_rotate(root->right, root);
    if (p) {
      if (root->data < p->data)
        p->left = left_rotate(root, p);
      else
        p->right = left_rotate(root, p);
    }
    else
      global_root = left_rotate(root, global_root);
  }
  else {
    if (root->parent)
      rebalance(root->parent, key);
  }
  return;
}


__device__ void insert(node* root, int key) {

  if (root == NULL) {             // Empty Tree
    root = new_node(key, NULL); 
    return;
  }
  
  int acquired = lock(root);

  if (acquired) {
    if (key < root->data) {
      if (root->left == NULL) {     // Can be inserted to the immediate left
        root->left = new_node(key, root);
        unlock(root);
        while (!atomicExch((int*)&MASTER_LOCK, 1));
        rebalance(root, key);
        atomicExch((int*)&MASTER_LOCK, 0);
        return;
      } else {                      // Release this Node and proceed
        unlock(root);
        insert(root->left, key);
      }
    } else {
      if (root->right == NULL) {    // Can be inserted to the immediate right
        root->right = new_node(key, root);
        unlock(root);
        while (!atomicExch((int*)&MASTER_LOCK, 1));
        rebalance(root, key);
        atomicExch((int*)&MASTER_LOCK, 0);
        return;
      } else {
        unlock(root);               // Release this Node and proceed
        insert(root->right, key);
      }
    }
    

  } else {
    insert(root, key);
  }
}


__device__ void pre_order(node* root)
{
    if(root != NULL)
    {
      printf("%d ", root->data);
      pre_order(root->left);
      pre_order(root->right);
    }
    return;
}

__device__ void in_order(node* root)
{
    if(root != NULL)
    {
      in_order(root->left);
      printf("%d ", root->data);
      in_order(root->right);
    }
    return;
}

	.file	"tmpxft_002d9edb_00000000-6_avl-fine.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2041:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2041:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z4lockP4node
	.type	_Z4lockP4node, @function
_Z4lockP4node:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z4lockP4node, .-_Z4lockP4node
	.globl	_Z6unlockP4node
	.type	_Z6unlockP4node, @function
_Z6unlockP4node:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z6unlockP4node, .-_Z6unlockP4node
	.globl	_Z8new_nodeiP4node
	.type	_Z8new_nodeiP4node, @function
_Z8new_nodeiP4node:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z8new_nodeiP4node, .-_Z8new_nodeiP4node
	.globl	_Z4findP4nodei
	.type	_Z4findP4nodei, @function
_Z4findP4nodei:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z4findP4nodei, .-_Z4findP4nodei
	.globl	_Z6heightP4node
	.type	_Z6heightP4node, @function
_Z6heightP4node:
.LFB2180:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z4lockP4node
	.cfi_endproc
.LFE2180:
	.size	_Z6heightP4node, .-_Z6heightP4node
	.globl	_Z11get_balanceP4node
	.type	_Z11get_balanceP4node, @function
_Z11get_balanceP4node:
.LFB2182:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z4lockP4node
	.cfi_endproc
.LFE2182:
	.size	_Z11get_balanceP4node, .-_Z11get_balanceP4node
	.globl	_Z11left_rotateP4nodeS0_
	.type	_Z11left_rotateP4nodeS0_, @function
_Z11left_rotateP4nodeS0_:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2033:
	.size	_Z11left_rotateP4nodeS0_, .-_Z11left_rotateP4nodeS0_
	.globl	_Z12right_rotateP4nodeS0_
	.type	_Z12right_rotateP4nodeS0_, @function
_Z12right_rotateP4nodeS0_:
.LFB2188:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z11left_rotateP4nodeS0_
	.cfi_endproc
.LFE2188:
	.size	_Z12right_rotateP4nodeS0_, .-_Z12right_rotateP4nodeS0_
	.globl	_Z9rebalanceP4nodei
	.type	_Z9rebalanceP4nodei, @function
_Z9rebalanceP4nodei:
.LFB2035:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2035:
	.size	_Z9rebalanceP4nodei, .-_Z9rebalanceP4nodei
	.globl	_Z6insertP4nodei
	.type	_Z6insertP4nodei, @function
_Z6insertP4nodei:
.LFB2190:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z9rebalanceP4nodei
	.cfi_endproc
.LFE2190:
	.size	_Z6insertP4nodei, .-_Z6insertP4nodei
	.globl	_Z9pre_orderP4node
	.type	_Z9pre_orderP4node, @function
_Z9pre_orderP4node:
.LFB2184:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z6unlockP4node
	.cfi_endproc
.LFE2184:
	.size	_Z9pre_orderP4node, .-_Z9pre_orderP4node
	.globl	_Z8in_orderP4node
	.type	_Z8in_orderP4node, @function
_Z8in_orderP4node:
.LFB2186:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z6unlockP4node
	.cfi_endproc
.LFE2186:
	.size	_Z8in_orderP4node, .-_Z8in_orderP4node
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"global_root"
.LC1:
	.string	"MASTER_LOCK"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2064:
	.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
	movl	$8, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	movq	%rdx, %rcx
	leaq	_ZL11global_root(%rip), %rsi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL11MASTER_LOCK(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.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
.LFE2064:
	.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
	.local	_ZL11MASTER_LOCK
	.comm	_ZL11MASTER_LOCK,4,4
	.local	_ZL11global_root
	.comm	_ZL11global_root,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 <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>

typedef struct node {
  int data;
  struct node *parent;
  struct node *left;
  struct node *right;
  int height;
  int sema;
} node;

__device__ node* global_root = NULL;

__device__ volatile int MASTER_LOCK = 0;

__device__ int lock(node* n) {
  int status = atomicExch(&n->sema, 1);
  return (!status && !MASTER_LOCK);
}

__device__ void unlock(node* n) {
  atomicExch(&n->sema, 0);
}

__device__ node* new_node(int val, node* parent) {
  node *tmp = (node *) malloc(sizeof(node));
  tmp->data = val;
  tmp->parent = parent;
  tmp->left = tmp->right = NULL;
  tmp->height = 1;
  tmp->sema = 0;
  return tmp;
}

__device__ node* find(node* root, int key) {
  if (root == NULL) return NULL;

  if (root->data == key) return root;
  else if (root->data > key) return find(root->left, key);
  else return find(root->right, key);
}

__device__ int height(node *root)
{
  if (root == NULL)
      return 0;
  return root->height;
}

__device__ int get_balance(node *root)
{
  if (root == NULL)
      return 0;
  return height(root->left) - height(root->right);
}


__device__ node* left_rotate(node* root, node* parent)
{
  node* temp1 = root->right;
  node* temp2 = temp1->left;

  temp1->left = root;
  root->parent = temp1;
  root->right = temp2;

  if (temp2)
    temp2->parent = root;

  root->height = max(height(root->left), height(root->right))+1;
  temp1->height = max(height(temp1->left), height(temp1->right))+1;

  temp1->parent = parent;
  return temp1;
}

__device__ node* right_rotate(node* root, node* parent)
{
  node* temp1 = root->left;
  node* temp2 = temp1->right;

  temp1->right = root;
  root->parent = temp1;
  root->left = temp2;

  if(temp2)
    temp2->parent = root;

  root->height = max(height(root->left), height(root->right))+1;
  temp1->height = max(height(temp1->left), height(temp1->right))+1;

  temp1->parent = parent;
  return temp1;
}


__device__ void rebalance(node* root, int key) {
  root->height = max(height(root->left),height(root->right)) + 1;
  int balance = get_balance(root);

  // Left Left Case
  node* p = root->parent;
  if (balance > 1 && key < root->left->data) {
    if (p) {
      if (root->data < p->data)
        p->left = right_rotate(root, p);
      else
        p->right = right_rotate(root, p);
    }
    else 
      global_root = right_rotate(root, global_root);
  }

  // Right Right Case
  else if (balance < -1 && key > root->right->data) {
    if (p) {
      if (root->data < p->data) 
        p->left = left_rotate(root, p);
      else
        p->right = left_rotate(root, p);
    }
    else
      global_root = left_rotate(root, global_root);
  }

  // Left Right Case
  else if (balance > 1 && key > root->left->data) {
    root->left =  left_rotate(root->left, root);
    if (p) {
      if (root->data < p->data)
        p->left = right_rotate(root, p);
      else 
        p->right = right_rotate(root, p);
    }
    else 
      global_root = right_rotate(root, global_root);
  }

  // Right Left Case
  else if (balance < -1 && key < root->right->data)
  {
    root->right = right_rotate(root->right, root);
    if (p) {
      if (root->data < p->data)
        p->left = left_rotate(root, p);
      else
        p->right = left_rotate(root, p);
    }
    else
      global_root = left_rotate(root, global_root);
  }
  else {
    if (root->parent)
      rebalance(root->parent, key);
  }
  return;
}


__device__ void insert(node* root, int key) {

  if (root == NULL) {             // Empty Tree
    root = new_node(key, NULL); 
    return;
  }
  
  int acquired = lock(root);

  if (acquired) {
    if (key < root->data) {
      if (root->left == NULL) {     // Can be inserted to the immediate left
        root->left = new_node(key, root);
        unlock(root);
        while (!atomicExch((int*)&MASTER_LOCK, 1));
        rebalance(root, key);
        atomicExch((int*)&MASTER_LOCK, 0);
        return;
      } else {                      // Release this Node and proceed
        unlock(root);
        insert(root->left, key);
      }
    } else {
      if (root->right == NULL) {    // Can be inserted to the immediate right
        root->right = new_node(key, root);
        unlock(root);
        while (!atomicExch((int*)&MASTER_LOCK, 1));
        rebalance(root, key);
        atomicExch((int*)&MASTER_LOCK, 0);
        return;
      } else {
        unlock(root);               // Release this Node and proceed
        insert(root->right, key);
      }
    }
    

  } else {
    insert(root, key);
  }
}


__device__ void pre_order(node* root)
{
    if(root != NULL)
    {
      printf("%d ", root->data);
      pre_order(root->left);
      pre_order(root->right);
    }
    return;
}

__device__ void in_order(node* root)
{
    if(root != NULL)
    {
      in_order(root->left);
      printf("%d ", root->data);
      in_order(root->right);
    }
    return;
}

	.text
	.file	"avl-fine.hip"
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB0_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB0_2:
	movl	$global_root, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$MASTER_LOCK, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$4, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end0:
	.size	__hip_module_ctor, .Lfunc_end0-__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	.LBB1_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
.LBB1_2:
	retq
.Lfunc_end1:
	.size	__hip_module_dtor, .Lfunc_end1-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	global_root,@object             # @global_root
	.local	global_root
	.comm	global_root,8,8
	.type	MASTER_LOCK,@object             # @MASTER_LOCK
	.local	MASTER_LOCK
	.comm	MASTER_LOCK,4,4
	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"global_root"
	.size	.L__unnamed_1, 12

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"MASTER_LOCK"
	.size	.L__unnamed_2, 12

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

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

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

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	global_root             ; @global_root
	.type	global_root,@object
	.section	.bss,"aw",@nobits
	.globl	global_root
	.p2align	3, 0x0
global_root:
	.quad	0
	.size	global_root, 8

	.protected	MASTER_LOCK             ; @MASTER_LOCK
	.type	MASTER_LOCK,@object
	.globl	MASTER_LOCK
	.p2align	2, 0x0
MASTER_LOCK:
	.long	0                               ; 0x0
	.size	MASTER_LOCK, 4

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 global_root
	.addrsig_sym MASTER_LOCK
	.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
