/*
Copyright 2017 the arraydiff authors

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/

#include "common.cuh"
#include <cuda_runtime_api.h>
#include <stdint.h>
#include <stdlib.h>

__global__ void symm_unit_clip_fwd_f32_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    float *y)
{
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float a = c / max(fabs(c), 1.0f);*/
    float a = clip[0];
    float x_i = x[idx];
    y[idx] = x_i * ((x_i > 0.0f) + a * (x_i < 0.0f));
  }
}

extern "C" void arraydiff_cuda_kernel_symm_unit_clip_fwd_f32(
    size_t dim,
    const float *clip,
    const float *x,
    float *y,
    cudaStream_t stream)
{
  symm_unit_clip_fwd_f32_kernel<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, clip, x, y);
}

__global__ void symm_unit_clip_param_bwd_f32_atomic_naive_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *grad)
{
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float u = max(fabs(c), 1.0f);
    float du = 1.0f * (c > 1.0f) - 1.0f * (c < -1.0f);
    float x_i = x[idx];
    atomicAdd(&grad[0], (1.0f / u) * (1.0f - du * c / u) * dy[idx] * x_i * (x_i < 0.0f));*/
    float x_i = x[idx];
    atomicAdd(&grad[0], dy[idx] * x_i * (x_i < 0.0f));
  }
}

__global__ void symm_unit_clip_param_bwd_f32_atomic_fast_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *grad)
{
  __shared__ float cache[1024];
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float u = max(fabs(c), 1.0f);
    float du = 1.0f * (c > 1.0f) - 1.0f * (c < -1.0f);
    float x_i = x[idx];
    cache[threadIdx.x] = (1.0f / u) * (1.0f - du * c / u) * dy[idx] * x_i * (x_i < 0.0f);*/
    float x_i = x[idx];
    cache[threadIdx.x] = dy[idx] * x_i * (x_i < 0.0f);
  } else {
    cache[threadIdx.x] = 0.0f;
  }
  __syncthreads();
  threadblock1024_reduce_sum_f32(cache);
  if (idx < dim) {
    if (threadIdx.x == 0) {
      atomicAdd(&grad[0], cache[0]);
    }
  }
}

extern "C" void arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32(
    size_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *grad,
    cudaStream_t stream)
{
  symm_unit_clip_param_bwd_f32_atomic_fast_kernel<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, clip, x, dy, grad);
}

__global__ void symm_unit_clip_input_bwd_f32_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *dx)
{
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float a = c / max(fabs(c), 1.0f);*/
    float a = clip[0];
    float x_i = x[idx];
    dx[idx] += dy[idx] * ((x_i > 0.0f) + a * (x_i < 0.0f));
  }
}

extern "C" void arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32(
    size_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *dx,
    cudaStream_t stream)
{
  symm_unit_clip_input_bwd_f32_kernel<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, clip, x, dy, dx);
}

	.file	"tmpxft_0035962d_00000000-6_clip.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2015:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2015:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf
	.type	_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf, @function
_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf:
.LFB2037:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2037:
	.size	_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf, .-_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf
	.globl	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.type	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, @function
_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf:
.LFB2038:
	.cfi_startproc
	endbr64
	jmp	_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf
	.cfi_endproc
.LFE2038:
	.size	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, .-_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.globl	arraydiff_cuda_kernel_symm_unit_clip_fwd_f32
	.type	arraydiff_cuda_kernel_symm_unit_clip_fwd_f32, @function
arraydiff_cuda_kernel_symm_unit_clip_fwd_f32:
.LFB2010:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	1023(%rdi), %rax
	movq	%rcx, %r13
	movq	%r8, %r9
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	shrq	$10, %rax
	movq	%rdx, %r12
	movl	$4194305, %edx
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rsi, %rbp
	xorl	%r8d, %r8d
	salq	$10, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movl	$1, %ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	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	.L8
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movq	%r13, %rcx
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	_Z56__device_stub__Z29symm_unit_clip_fwd_f32_kerneljPKfS0_PfjPKfS0_Pf
.L8:
	.cfi_restore_state
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2010:
	.size	arraydiff_cuda_kernel_symm_unit_clip_fwd_f32, .-arraydiff_cuda_kernel_symm_unit_clip_fwd_f32
	.globl	_Z78__device_stub__Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.type	_Z78__device_stub__Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_PfjPKfS0_S0_Pf, @function
_Z78__device_stub__Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_PfjPKfS0_S0_Pf:
.LFB2039:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L11:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2039:
	.size	_Z78__device_stub__Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_PfjPKfS0_S0_Pf, .-_Z78__device_stub__Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.globl	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.type	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, @function
_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf:
.LFB2040:
	.cfi_startproc
	endbr64
	jmp	_Z78__device_stub__Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.cfi_endproc
.LFE2040:
	.size	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, .-_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.globl	_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.type	_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf, @function
_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf:
.LFB2041:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L16:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2041:
	.size	_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf, .-_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.globl	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.type	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, @function
_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf:
.LFB2042:
	.cfi_startproc
	endbr64
	jmp	_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.cfi_endproc
.LFE2042:
	.size	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, .-_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.globl	arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32
	.type	arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32, @function
arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32:
.LFB2011:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	1023(%rdi), %rax
	movq	%r8, %r14
	xorl	%r8d, %r8d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	shrq	$10, %rax
	movq	%rcx, %r13
	movl	$1, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdx, %r12
	movl	$4194305, %edx
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rsi, %rbp
	salq	$10, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	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	.L21
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movq	%r14, %r8
	movq	%r13, %rcx
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z77__device_stub__Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
.L21:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2011:
	.size	arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32, .-arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32
	.globl	_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.type	_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf, @function
_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf:
.LFB2043:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L24:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2043:
	.size	_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf, .-_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.globl	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.type	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, @function
_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf:
.LFB2044:
	.cfi_startproc
	endbr64
	jmp	_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
	.cfi_endproc
.LFE2044:
	.size	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, .-_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.globl	arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32
	.type	arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32, @function
arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32:
.LFB2012:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	1023(%rdi), %rax
	movq	%r8, %r14
	xorl	%r8d, %r8d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	shrq	$10, %rax
	movq	%rcx, %r13
	movl	$1, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%rdx, %r12
	movl	$4194305, %edx
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rsi, %rbp
	salq	$10, %rdx
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	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	.L29
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movq	%r14, %r8
	movq	%r13, %rcx
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z65__device_stub__Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_PfjPKfS0_S0_Pf
.L29:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2012:
	.size	arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32, .-arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf"
.LC1:
	.string	"_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf"
.LC2:
	.string	"_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf"
.LC3:
	.string	"_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2046:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf(%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	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf(%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	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf(%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	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf(%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
.LFE2046:
	.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:

/*
Copyright 2017 the arraydiff authors

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/


#include <hip/hip_runtime.h>
#include "common.cuh"
#include <hip/hip_runtime_api.h>
#include <stdint.h>
#include <stdlib.h>

__global__ void symm_unit_clip_fwd_f32_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    float *y)
{
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float a = c / max(fabs(c), 1.0f);*/
    float a = clip[0];
    float x_i = x[idx];
    y[idx] = x_i * ((x_i > 0.0f) + a * (x_i < 0.0f));
  }
}

extern "C" void arraydiff_cuda_kernel_symm_unit_clip_fwd_f32(
    size_t dim,
    const float *clip,
    const float *x,
    float *y,
    hipStream_t stream)
{
  symm_unit_clip_fwd_f32_kernel<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, clip, x, y);
}

__global__ void symm_unit_clip_param_bwd_f32_atomic_naive_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *grad)
{
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float u = max(fabs(c), 1.0f);
    float du = 1.0f * (c > 1.0f) - 1.0f * (c < -1.0f);
    float x_i = x[idx];
    atomicAdd(&grad[0], (1.0f / u) * (1.0f - du * c / u) * dy[idx] * x_i * (x_i < 0.0f));*/
    float x_i = x[idx];
    atomicAdd(&grad[0], dy[idx] * x_i * (x_i < 0.0f));
  }
}

__global__ void symm_unit_clip_param_bwd_f32_atomic_fast_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *grad)
{
  __shared__ float cache[1024];
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float u = max(fabs(c), 1.0f);
    float du = 1.0f * (c > 1.0f) - 1.0f * (c < -1.0f);
    float x_i = x[idx];
    cache[threadIdx.x] = (1.0f / u) * (1.0f - du * c / u) * dy[idx] * x_i * (x_i < 0.0f);*/
    float x_i = x[idx];
    cache[threadIdx.x] = dy[idx] * x_i * (x_i < 0.0f);
  } else {
    cache[threadIdx.x] = 0.0f;
  }
  __syncthreads();
  threadblock1024_reduce_sum_f32(cache);
  if (idx < dim) {
    if (threadIdx.x == 0) {
      atomicAdd(&grad[0], cache[0]);
    }
  }
}

extern "C" void arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32(
    size_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *grad,
    hipStream_t stream)
{
  symm_unit_clip_param_bwd_f32_atomic_fast_kernel<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, clip, x, dy, grad);
}

__global__ void symm_unit_clip_input_bwd_f32_kernel(
    uint32_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *dx)
{
  uint32_t idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < dim) {
    /*float c = clip[0];
    float a = c / max(fabs(c), 1.0f);*/
    float a = clip[0];
    float x_i = x[idx];
    dx[idx] += dy[idx] * ((x_i > 0.0f) + a * (x_i < 0.0f));
  }
}

extern "C" void arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32(
    size_t dim,
    const float *clip,
    const float *x,
    const float *dy,
    float *dx,
    hipStream_t stream)
{
  symm_unit_clip_input_bwd_f32_kernel<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, clip, x, dy, dx);
}

	.text
	.file	"clip.hip"
	.globl	_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf # -- Begin function _Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.type	_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf,@function
_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf: # @_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	4(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	32(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	24(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, .Lfunc_end0-_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_symm_unit_clip_fwd_f32 # -- Begin function arraydiff_cuda_kernel_symm_unit_clip_fwd_f32
	.type	arraydiff_cuda_kernel_symm_unit_clip_fwd_f32,@function
arraydiff_cuda_kernel_symm_unit_clip_fwd_f32: # @arraydiff_cuda_kernel_symm_unit_clip_fwd_f32
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.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 %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %r9
	movq	%rcx, %rbx
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %r12
	leaq	1023(%rdi), %rax
	shrq	$10, %rax
	movl	%eax, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB1_2
# %bb.1:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB1_2:
	.cfi_def_cfa_offset 48
	movl	%r12d, %edi
	movq	%r15, %rsi
	movq	%r14, %rdx
	movq	%rbx, %rcx
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf # TAILCALL
.Lfunc_end1:
	.size	arraydiff_cuda_kernel_symm_unit_clip_fwd_f32, .Lfunc_end1-arraydiff_cuda_kernel_symm_unit_clip_fwd_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf # -- Begin function _Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.type	_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf,@function
_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf: # @_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, .Lfunc_end2-_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.cfi_endproc
                                        # -- End function
	.globl	_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf # -- Begin function _Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.type	_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf,@function
_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf: # @_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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_end3:
	.size	_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, .Lfunc_end3-_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32 # -- Begin function arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32
	.type	arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32,@function
arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32: # @arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %rbx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %r13
	leaq	1023(%rdi), %rax
	shrq	$10, %rax
	movl	%eax, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB4_2
# %bb.1:
	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
.LBB4_2:
	.cfi_def_cfa_offset 48
	movl	%r13d, %edi
	movq	%r12, %rsi
	movq	%r15, %rdx
	movq	%r14, %rcx
	movq	%rbx, %r8
	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
	jmp	_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf # TAILCALL
.Lfunc_end4:
	.size	arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32, .Lfunc_end4-arraydiff_cuda_kernel_symm_unit_clip_param_bwd_nondeterministic_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf # -- Begin function _Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.type	_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf,@function
_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf: # @_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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	_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, .Lfunc_end5-_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32 # -- Begin function arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32
	.type	arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32,@function
arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32: # @arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%r8, %rbx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %r13
	leaq	1023(%rdi), %rax
	shrq	$10, %rax
	movl	%eax, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB6_2
# %bb.1:
	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
.LBB6_2:
	.cfi_def_cfa_offset 48
	movl	%r13d, %edi
	movq	%r12, %rsi
	movq	%r15, %rdx
	movq	%r14, %rcx
	movq	%rbx, %r8
	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
	jmp	_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf # TAILCALL
.Lfunc_end6:
	.size	arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32, .Lfunc_end6-arraydiff_cuda_kernel_symm_unit_clip_input_bwd_f32
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB8_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf,@object # @_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.section	.rodata,"a",@progbits
	.globl	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.p2align	3, 0x0
_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf:
	.quad	_Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.size	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, 8

	.type	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf,@object # @_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.globl	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.p2align	3, 0x0
_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf:
	.quad	_Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.size	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, 8

	.type	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf,@object # @_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.globl	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.p2align	3, 0x0
_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf:
	.quad	_Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.size	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, 8

	.type	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf,@object # @_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.globl	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.p2align	3, 0x0
_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf:
	.quad	_Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.size	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf"
	.size	.L__unnamed_1, 43

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf"
	.size	.L__unnamed_2, 65

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf"
	.size	.L__unnamed_3, 64

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf"
	.size	.L__unnamed_4, 52

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

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

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z44__device_stub__symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.addrsig_sym _Z63__device_stub__symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.addrsig_sym _Z62__device_stub__symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.addrsig_sym _Z50__device_stub__symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.addrsig_sym _Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.addrsig_sym _Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.addrsig_sym _Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf ; -- Begin function _Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.globl	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
	.p2align	8
	.type	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf,@function
_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf: ; @_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x8
	s_load_b64 s[0:1], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_load_b32 s2, s[4:5], 0x0
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_add_co_u32 v2, vcc_lo, s6, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v1, vcc_lo
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_f32_e32 vcc_lo, 0, v2
	v_cndmask_b32_e64 v3, 0, 1.0, vcc_lo
	v_cmp_gt_f32_e32 vcc_lo, 0, v2
	v_cndmask_b32_e64 v4, 0, 1.0, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v3, s2, v4
	v_mul_f32_e32 v2, v2, v3
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf, .Lfunc_end0-_Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 208
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf ; -- Begin function _Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.globl	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
	.p2align	8
	.type	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf,@function
_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf: ; @_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x34
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB1_6
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x10
	v_mov_b32_e32 v2, 0
	s_load_b64 s[0:1], s[0:1], 0x20
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	global_load_b32 v2, v[2:3], off
	global_load_b32 v0, v[0:1], off
	s_waitcnt vmcnt(1)
	v_cmp_gt_f32_e32 vcc_lo, 0, v2
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v0, v2, v0
	v_bfrev_b32_e32 v2, 1
	v_cndmask_b32_e64 v1, 0, 1.0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f32_e32 v0, v0, v1
.LBB1_2:                                ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s3, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s4, v0, s3
	s_lshl_b32 s3, 1, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s2, s2, s3
	s_cmp_lg_u32 s2, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v2, s4, v2
	s_cbranch_scc1 .LBB1_2
; %bb.3:                                ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s2, 0
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s3, exec_lo, s3
	s_cbranch_execz .LBB1_6
; %bb.4:
	s_load_b32 s3, s[0:1], 0x0
	v_mov_b32_e32 v3, 0
	s_waitcnt lgkmcnt(0)
	v_mov_b32_e32 v1, s3
.LBB1_5:                                ; %atomicrmw.start
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v0, v1, v2
	global_atomic_cmpswap_b32 v0, v3, v[0:1], s[0:1] glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, v0, v1
	v_mov_b32_e32 v1, v0
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB1_5
.LBB1_6:                                ; %Flow23
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.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 4
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf, .Lfunc_end1-_Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 324
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf ; -- Begin function _Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.globl	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
	.p2align	8
	.type	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf,@function
_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf: ; @_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
; %bb.0:
	s_clause 0x3
	s_load_b32 s6, s[0:1], 0x34
	s_load_b32 s7, s[0:1], 0x0
	s_load_b64 s[4:5], s[0:1], 0x20
	s_load_b128 s[0:3], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s6, s6, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s6, v[0:1]
	v_mov_b32_e32 v2, 0
	v_cmp_gt_u32_e32 vcc_lo, s7, v1
	s_and_saveexec_b32 s6, vcc_lo
	s_cbranch_execz .LBB2_2
; %bb.1:
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v3, s0, s0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, s0, s1, v2, s0
	v_add_co_u32 v1, s0, s2, v1
	v_add_co_ci_u32_e64 v2, s0, s3, v2, s0
	global_load_b32 v3, v[3:4], off
	global_load_b32 v1, v[1:2], off
	s_waitcnt vmcnt(1)
	v_cmp_gt_f32_e64 s0, 0, v3
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v1, v3, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v2, 0, 1.0, s0
	v_mul_f32_e32 v2, v1, v2
.LBB2_2:
	s_or_b32 exec_lo, exec_lo, s6
	v_lshlrev_b32_e32 v1, 2, v0
	s_movk_i32 s1, 0x200
	ds_store_b32 v1, v2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB2_3:                                ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s1, v0
	s_cbranch_execz .LBB2_5
; %bb.4:                                ;   in Loop: Header=BB2_3 Depth=1
	v_lshl_add_u32 v2, s1, 2, v1
	ds_load_b32 v2, v2
	ds_load_b32 v3, v1
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v2, v2, v3
	ds_store_b32 v1, v2
.LBB2_5:                                ;   in Loop: Header=BB2_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_lshr_b32 s0, s1, 1
	s_cmp_gt_u32 s1, 1
	s_mov_b32 s1, s0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB2_3
; %bb.6:                                ; %_Z30threadblock1024_reduce_sum_f32Pf.exit
	v_cmp_eq_u32_e64 s0, 0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, vcc_lo
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB2_10
; %bb.7:
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mbcnt_lo_u32_b32 v0, s0, 0
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	s_and_b32 s1, exec_lo, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 exec_lo, s1
	s_cbranch_execz .LBB2_10
; %bb.8:
	s_bcnt1_i32_b32 s0, s0
	s_load_b32 s1, s[4:5], 0x0
	v_cvt_f32_ubyte0_e32 v1, s0
	v_mov_b32_e32 v2, 0
	s_mov_b32 s0, 0
	ds_load_b32 v0, v2
	s_waitcnt lgkmcnt(0)
	v_mul_f32_e32 v3, v0, v1
	v_mov_b32_e32 v1, s1
.LBB2_9:                                ; %atomicrmw.start
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v0, v1, v3
	global_atomic_cmpswap_b32 v0, v2, v[0:1], s[4:5] glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, v0, v1
	v_mov_b32_e32 v1, v0
	s_or_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s0
	s_cbranch_execnz .LBB2_9
.LBB2_10:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
		.amdhsa_group_segment_fixed_size 4096
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.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 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf, .Lfunc_end2-_Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 460
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 4096 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf ; -- Begin function _Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.globl	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
	.p2align	8
	.type	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf,@function
_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf: ; @_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x34
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB3_2
; %bb.1:
	s_load_b256 s[0:7], s[0:1], 0x8
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v1, vcc_lo
	global_load_b32 v4, v[2:3], off
	v_add_co_u32 v2, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	global_load_b32 v2, v[2:3], off
	global_load_b32 v3, v[0:1], off
	s_load_b32 s0, s[0:1], 0x0
	s_waitcnt vmcnt(2)
	v_cmp_lt_f32_e32 vcc_lo, 0, v4
	v_cndmask_b32_e64 v5, 0, 1.0, vcc_lo
	v_cmp_gt_f32_e32 vcc_lo, 0, v4
	v_cndmask_b32_e64 v4, 0, 1.0, vcc_lo
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, s0, v4
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v3, v2, v5
	global_store_b32 v[0:1], v3, off
.LBB3_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.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 6
		.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_end3:
	.size	_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf, .Lfunc_end3-_Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 228
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z29symm_unit_clip_fwd_f32_kerneljPKfS0_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z48symm_unit_clip_param_bwd_f32_atomic_naive_kerneljPKfS0_S0_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 4096
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z47symm_unit_clip_param_bwd_f32_atomic_fast_kerneljPKfS0_S0_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z35symm_unit_clip_input_bwd_f32_kerneljPKfS0_S0_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
