#include <cuda.h>  
#include <float.h>
const float MAX_EXP = 80. ; 

__global__ void kCumulativeSum(
                               float *probabilities, 
                               float *sumProbabilities, 
                               int softMaxWidth, 
                               int signalLength
                              )
{
   extern __shared__ float sData[] ; 
   float *dataPrev = (float *) &sData[0] ; 
   float *dataNext = (float *) &sData[softMaxWidth] ; 

   // first copy data for current softMax
   int signalIndex = signalLength * blockIdx.x ;
   int dataIndex = blockIdx.y * softMaxWidth + threadIdx.x ;

   int maxLength = min(softMaxWidth, signalLength-dataIndex+threadIdx.x) ; 
   for (int index = threadIdx.x ; index < maxLength && dataIndex < signalLength ; 
        index+= blockDim.x, dataIndex += blockDim.x)
   {
      dataPrev[index] = probabilities[signalIndex+dataIndex] ; 
   }
   __syncthreads() ; 

   // Now compute cumulative sum
   dataIndex = blockIdx.y * softMaxWidth  + threadIdx.x ;
   for (int round = 1 ; round <  maxLength ; round = round<<1)
   {
      for (int index = threadIdx.x ; index < maxLength ; index+= blockDim.x)
      {
         float val = dataPrev[index] ; 
         if (index >= round)
            val += dataPrev[index-round] ;
         dataNext[index] = val ; 
      }
      float *temp = dataPrev ; 
      dataPrev = dataNext ; 
      dataNext = temp ; 
      __syncthreads() ; 
   }

   // Write out the data.
   for (int index = threadIdx.x ; index < maxLength && dataIndex < signalLength ; 
        index+= blockDim.x, dataIndex += blockDim.x)
   {
      sumProbabilities[signalIndex+dataIndex] = dataPrev[index] ;
   }
}

__global__ void kMultinomialSample(
                                   float *unifRandNums, 
                                   float *probabilities, 
                                   float *samples, 
                                   int softMaxWidth, 
                                   int startShift,
                                   int signalLength
                                  )
{
   extern __shared__ float sData[] ; 
   float *dataPrev = (float *) &sData[0] ; 
   float *dataNext = (float *) &sData[softMaxWidth] ; 


   // first copy data for current softMax
   int signalIndex = signalLength * blockIdx.x ;
   int dataIndex = blockIdx.y * softMaxWidth + threadIdx.x + startShift ;

   float randNum = unifRandNums[blockIdx.x*gridDim.y + blockIdx.y] ; 

   int maxLength = min(softMaxWidth, signalLength-dataIndex+threadIdx.x) ; 
   for (int index = threadIdx.x ; index < maxLength && dataIndex < signalLength ; 
        index+= blockDim.x, dataIndex += blockDim.x)
   {
      dataPrev[index] = probabilities[signalIndex+dataIndex] ; 
   }
   __syncthreads() ; 

   // Now compute cumulative sum
   dataIndex = blockIdx.y * softMaxWidth  + threadIdx.x + startShift ;
   for (int round = 1 ; round <  maxLength ; round = round<<1)
   {
      for (int index = threadIdx.x ; index < maxLength ; index+= blockDim.x)
      {
         float val = dataPrev[index] ; 
         if (index >= round)
            val += dataPrev[index-round] ;
         dataNext[index] = val ; 
      }
      float *temp = dataPrev ; 
      dataPrev = dataNext ; 
      dataNext = temp ; 
      __syncthreads() ; 
   }

   // Find the appropriate index where cumulative[i-1]<= r < cumulative[i]
   for (int index = threadIdx.x ; index < maxLength && dataIndex < signalLength ; 
        index+= blockDim.x, dataIndex += blockDim.x)
   {
      if (index == 0)
      {
         if (randNum < dataPrev[index])
            samples[signalIndex+dataIndex] = 1 ;
         else
            samples[signalIndex+dataIndex] = 0 ;

      }
      else
      {
         if (randNum >= dataPrev[index-1] && randNum < dataPrev[index])
            samples[signalIndex+dataIndex] = 1 ;
         else
            samples[signalIndex+dataIndex] = 0 ;
      }
   }
}

__global__ void kSoftMaxStackApproxWithPositionBiases(
                                               float *activations, 
                                               float *probabilities, 
                                               float *stdevs, 
                                               float *featureBiases,
                                               float *positionBiases,
                                               int softMaxWidth, 
                                               int signalLength
                                               )
{
   extern __shared__ float sData[] ; 
   float *arrData = (float *) &sData[0] ; 
   float *threadStores = (float *) &sData[softMaxWidth] ; 

   int signalIndex = blockIdx.x*signalLength ; 
   int dataIndex = blockIdx.y*softMaxWidth + threadIdx.x ; 
  
   float maxVal = -FLT_MAX ;  
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      arrData[index] = activations[signalIndex+dataIndex] + positionBiases[index] ; 
      if (maxVal < arrData[index])
         maxVal = arrData[index] ; 
   }
   threadStores[threadIdx.x] = maxVal ; 
   __syncthreads() ; 

   // do a reduction to find the max of all maxes. 
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
      if (threadIdx.x < s)
      {
         threadStores[threadIdx.x] = 
              threadStores[threadIdx.x+s]*(threadStores[threadIdx.x + s] >=threadStores[threadIdx.x]) + 
              threadStores[threadIdx.x] * (threadStores[threadIdx.x] > threadStores[threadIdx.x+s])  ;
      }
      __syncthreads();
   }

   // now we have max. Lets subract it from all elements, and compute intermediate logSumExp over all elements 
   // a thread is responsible for
   float bias = featureBiases[blockIdx.x] ; 
   maxVal = threadStores[0] ; 

   __syncthreads();

   float sumExp = 0. ;

   dataIndex = blockIdx.y*softMaxWidth + threadIdx.x ; 
  
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      arrData[index] = arrData[index] - maxVal ; 
      sumExp += __expf(arrData[index]) ; 
   }

   threadStores[threadIdx.x] = sumExp ;
   __syncthreads() ; 

   // compute normalization constant over sumExp by summing together all the intermediate values.
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
       if (threadIdx.x < s)
          threadStores[threadIdx.x] += threadStores[threadIdx.x + s];
       __syncthreads();
   }
   __syncthreads() ; 

   sumExp = threadStores[0]  ; 

   float maxValNew = fmaxf(maxVal, bias) ;
  
   bias = bias - maxValNew ;  
   float reluSumExp = sumExp * __expf(maxVal-maxValNew) + __expf(bias) ;

   // notice, reusing maxValNew and maxVal variables for new causes.

   // maxValNew is the negative delta from the previous one.
   maxValNew = maxVal-maxValNew ; 
   // maxVal is -log(exp(biases)/(exp(biases)+sum(exp(x_i))))
   maxVal = -bias + __logf(reluSumExp) ;


   // write out probabilities and standard deviations
   dataIndex = blockIdx.y*softMaxWidth + threadIdx.x ; 
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      float r_x_i = __expf(arrData[index])/ sumExp ; 
      float relu_r_x_i = __expf(arrData[index]+maxValNew)/ reluSumExp ;  
      probabilities[signalIndex+dataIndex]  = r_x_i * maxVal ; 
      float variance = r_x_i * ((1-r_x_i)*maxVal + relu_r_x_i) ; 
      if (variance < 0) // stupid overflows, underflows etc, causing a perfectly reasonable calculation to look negative..
         variance = 0.0 ; 
      stdevs[signalIndex+dataIndex]  = sqrtf(variance) ; 
   }
   __syncthreads() ; 
  
}

__global__ void kSoftMaxReluWithPositionBiases(
                                                 float *activations, 
                                                 float *probabilities, 
                                                 float *meanValues, 
                                                 float *featureStdevs, 
                                                 float *featureBiases,
                                                 float *positionBiases,
                                                 int softMaxWidth, 
                                                 int shift, 
                                                 int signalLength, 
                                                 float minExpForSum
                                              )
{
   extern __shared__ float sData[] ; 
   float *arrData = (float *) &sData[0] ; 
   float *arrActivation = (float *) &sData[softMaxWidth] ; 
   float *threadStores = (float *) &sData[2*softMaxWidth] ; 

   int signalIndex = blockIdx.x*signalLength ; 
   int dataIndex = blockIdx.y*softMaxWidth + threadIdx.x + shift ; 
  
   float maxVal = -FLT_MAX ;  
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      float x_i = activations[signalIndex+dataIndex] + positionBiases[index] ; 
      arrActivation[index] = x_i ; 
      float val = x_i ; 
      x_i = x_i - .5 ; 
      while (x_i >= minExpForSum)
      {
        if (x_i > MAX_EXP)
           val = val + x_i ;
        else
           val = val +  __logf(1. + __expf(x_i)) ; 
        x_i = x_i - 1.0 ; 
      }
      
      arrData[index] = val ; 
      if (maxVal < arrData[index])
         maxVal = arrData[index] ; 
   }
   threadStores[threadIdx.x] = maxVal ; 
   __syncthreads() ; 

   // do a reduction to find the max of all maxes. 
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
      if (threadIdx.x < s)
      {
         threadStores[threadIdx.x] = 
              threadStores[threadIdx.x+s]*(threadStores[threadIdx.x + s] >=threadStores[threadIdx.x]) + 
              threadStores[threadIdx.x] * (threadStores[threadIdx.x] > threadStores[threadIdx.x+s])  ;
      }
      __syncthreads();
   }

   // now we have max. Lets subract it from all elements, and compute intermediate logSumExp over all elements 
   // a thread is responsible for
   float bias = featureBiases[blockIdx.x] ; 
   maxVal = fmaxf(threadStores[0], bias) ; 

   __syncthreads();

   float sumExp = 0. ;

   dataIndex = blockIdx.y*softMaxWidth + threadIdx.x + shift ; 
  
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      arrData[index] = arrData[index] - maxVal ; 
      arrActivation[index] = arrActivation[index] - bias ;
      sumExp += __expf(arrData[index]) ; 
   }

   threadStores[threadIdx.x] = sumExp ;
   __syncthreads() ; 

   // compute normalization constant over sumExp by summing together all the intermediate values.
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
       if (threadIdx.x < s)
          threadStores[threadIdx.x] += threadStores[threadIdx.x + s];
       __syncthreads();
   }
   __syncthreads() ; 
   sumExp = threadStores[0] + __expf(bias-maxVal) ;

   // write out probabilities. 
   dataIndex = blockIdx.y*softMaxWidth + threadIdx.x + shift ; 
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      float prob = __expf(arrData[index])/ sumExp ; 
      probabilities[signalIndex+dataIndex]  = prob ; 
      // using the relu approximation. Remember to threshold
      meanValues[signalIndex+dataIndex]  = arrActivation[index] ; 
      featureStdevs[signalIndex+dataIndex]  = sqrtf(1. / (1 + __expf(-arrActivation[index])));
   }
   __syncthreads() ; 
  
}

__global__ void kSoftMaxWithOffAndPositionBiases(
                                                 float *activations, 
                                                 float *probabilities, 
                                                 float *featureBiases,
                                                 float *positionBiases,
                                                 int softMaxWidth, 
                                                 int signalLength
                                                )
{
   extern __shared__ float sData[] ; 
   float *arrData = (float *) &sData[0] ; 
   float *threadStores = (float *) &sData[softMaxWidth] ; 

   int signalIndex = blockIdx.x*signalLength ; 
   int dataIndex = blockIdx.y*softMaxWidth + threadIdx.x ; 
  
   float maxVal = -FLT_MAX ;  
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      arrData[index] = activations[signalIndex+dataIndex] + positionBiases[index] ; 
      if (maxVal < arrData[index])
         maxVal = arrData[index] ; 
   }
   threadStores[threadIdx.x] = maxVal ; 
   __syncthreads() ; 

   // do a reduction to find the max of all maxes. 
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
      if (threadIdx.x < s)
      {
         threadStores[threadIdx.x] = 
              threadStores[threadIdx.x+s]*(threadStores[threadIdx.x + s] >=threadStores[threadIdx.x]) + 
              threadStores[threadIdx.x] * (threadStores[threadIdx.x] > threadStores[threadIdx.x+s])  ;
      }
      __syncthreads();
   }

   // now we have max. Lets subract it from all elements, and compute intermediate logSumExp over all elements 
   // a thread is responsible for
   float bias = featureBiases[blockIdx.x] ; 
   maxVal = fmaxf(threadStores[0], bias) ; 

   __syncthreads();

   float sumExp = 0. ;

   dataIndex = blockIdx.y*softMaxWidth + threadIdx.x ; 
  
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      arrData[index] = arrData[index] - maxVal ; 
      sumExp += __expf(arrData[index]) ; 
   }

   threadStores[threadIdx.x] = sumExp ;
   __syncthreads() ; 

   // compute normalization constant over sumExp by summing together all the intermediate values.
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
       if (threadIdx.x < s)
          threadStores[threadIdx.x] += threadStores[threadIdx.x + s];
       __syncthreads();
   }
   __syncthreads() ; 
   sumExp = threadStores[0] + __expf(bias-maxVal) ;

   // write out probabilities. 
   dataIndex = blockIdx.y*softMaxWidth + threadIdx.x ; 
   for (int index = threadIdx.x ; index < softMaxWidth && dataIndex < signalLength ; 
            index += blockDim.x, dataIndex += blockDim.x)
   {
      probabilities[signalIndex+dataIndex]  = __expf(arrData[index])/ sumExp ; 
   }
   __syncthreads() ; 
  
}



__global__ void kSoftMaxWithOff(float *activations, float *probabilities, float *biases,
                       int softMaxWidth, int signalLength, int numPtsPerThread)
{
   extern __shared__ float sData[] ; 
   float *arrData = (float *) &sData[0] ; 
   float *threadStores = (float *) &sData[softMaxWidth] ; 

   int signalIndex = blockIdx.x*signalLength ; 
   int blockIndex = blockIdx.y*softMaxWidth ; 
   int softMaxIndex = blockIndex + threadIdx.x*numPtsPerThread ;

   // copy to local memory
   int numToCopy = numPtsPerThread ; 
   if (softMaxIndex + numToCopy > signalLength)
      numToCopy = signalLength - softMaxIndex ; 
   if (softMaxIndex + numToCopy > blockIndex + softMaxWidth)
      numToCopy = blockIndex + softMaxWidth - softMaxIndex ; 
  
   float maxVal = -FLT_MAX ;  
   for (int index = 0 ; index < numToCopy ; index++)
   {
      arrData[softMaxIndex+index-blockIndex] = activations[signalIndex+softMaxIndex+index] ; 
      if (maxVal < arrData[softMaxIndex+index])
         maxVal = arrData[softMaxIndex+index-blockIndex] ; 
   }
   threadStores[threadIdx.x] = maxVal ; 
   __syncthreads() ; 

   // do a reduction to find the max of all maxes. 
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
      if (threadIdx.x < s)
      {
         threadStores[threadIdx.x] = 
              threadStores[threadIdx.x+s]*(threadStores[threadIdx.x + s] >=threadStores[threadIdx.x]) + 
              threadStores[threadIdx.x] * (threadStores[threadIdx.x] > threadStores[threadIdx.x+s])  ;
      }
      __syncthreads();
   }

   // now we have max. Lets subract it from all elements, and compute intermediate logSumExp over all elements 
   // a thread is responsible for
   float bias = biases[blockIdx.x] ; 
   maxVal = threadStores[0] * (threadStores[0] > bias) + bias * (bias >= threadStores[0]) ; 
   __syncthreads();

   float sumExp = 0. ;  
   for (int index = 0 ; index < numToCopy ; index++)
   {
      arrData[softMaxIndex+index-blockIndex] = arrData[softMaxIndex+index-blockIndex] - maxVal ; 
      sumExp += __expf(arrData[softMaxIndex+index-blockIndex]) ;
   }

   threadStores[threadIdx.x] = sumExp ;
   __syncthreads() ; 

   // compute normalization constant over sumExp by summing together all the intermediate values.
   for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
   {
       if (threadIdx.x < s)
          threadStores[threadIdx.x] += threadStores[threadIdx.x + s];
       __syncthreads();
   }
   __syncthreads() ; 
   sumExp = threadStores[0] + __expf(bias-maxVal) ;

   for (int index = 0 ; index < numToCopy ; index++)
   {
      probabilities[signalIndex+softMaxIndex+index]  = __expf(arrData[softMaxIndex+index-blockIndex])/ sumExp ; 
   }
   __syncthreads() ; 
  
}


	.file	"tmpxft_002d188b_00000000-6_SoftMax_kernels.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z38__device_stub__Z14kCumulativeSumPfS_iiPfS_ii
	.type	_Z38__device_stub__Z14kCumulativeSumPfS_iiPfS_ii, @function
_Z38__device_stub__Z14kCumulativeSumPfS_iiPfS_ii:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%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	_Z14kCumulativeSumPfS_ii(%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
.LFE2032:
	.size	_Z38__device_stub__Z14kCumulativeSumPfS_iiPfS_ii, .-_Z38__device_stub__Z14kCumulativeSumPfS_iiPfS_ii
	.globl	_Z14kCumulativeSumPfS_ii
	.type	_Z14kCumulativeSumPfS_ii, @function
_Z14kCumulativeSumPfS_ii:
.LFB2033:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z14kCumulativeSumPfS_iiPfS_ii
	.cfi_endproc
.LFE2033:
	.size	_Z14kCumulativeSumPfS_ii, .-_Z14kCumulativeSumPfS_ii
	.globl	_Z45__device_stub__Z18kMultinomialSamplePfS_S_iiiPfS_S_iii
	.type	_Z45__device_stub__Z18kMultinomialSamplePfS_S_iiiPfS_S_iii, @function
_Z45__device_stub__Z18kMultinomialSamplePfS_S_iiiPfS_S_iii:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z18kMultinomialSamplePfS_S_iii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L8:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z45__device_stub__Z18kMultinomialSamplePfS_S_iiiPfS_S_iii, .-_Z45__device_stub__Z18kMultinomialSamplePfS_S_iiiPfS_S_iii
	.globl	_Z18kMultinomialSamplePfS_S_iii
	.type	_Z18kMultinomialSamplePfS_S_iii, @function
_Z18kMultinomialSamplePfS_S_iii:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z18kMultinomialSamplePfS_S_iiiPfS_S_iii
	.cfi_endproc
.LFE2035:
	.size	_Z18kMultinomialSamplePfS_S_iii, .-_Z18kMultinomialSamplePfS_S_iii
	.globl	_Z67__device_stub__Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_iiPfS_S_S_S_ii
	.type	_Z67__device_stub__Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_iiPfS_S_S_S_ii, @function
_Z67__device_stub__Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_iiPfS_S_S_S_ii:
.LFB2036:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%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)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%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)
	leaq	4(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	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 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L13:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z67__device_stub__Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_iiPfS_S_S_S_ii, .-_Z67__device_stub__Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_iiPfS_S_S_S_ii
	.globl	_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii
	.type	_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii, @function
_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii:
.LFB2037:
	.cfi_startproc
	endbr64
	jmp	_Z67__device_stub__Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_iiPfS_S_S_S_ii
	.cfi_endproc
.LFE2037:
	.size	_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii, .-_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii
	.globl	_Z64__device_stub__Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiifPfS_S_S_S_S_iiif
	.type	_Z64__device_stub__Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiifPfS_S_S_S_S_iiif, @function
_Z64__device_stub__Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiifPfS_S_S_S_S_iiif:
.LFB2038:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	248(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	80(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 256
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L18:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z64__device_stub__Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiifPfS_S_S_S_S_iiif, .-_Z64__device_stub__Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiifPfS_S_S_S_S_iiif
	.globl	_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif
	.type	_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif, @function
_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif:
.LFB2039:
	.cfi_startproc
	endbr64
	jmp	_Z64__device_stub__Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiifPfS_S_S_S_S_iiif
	.cfi_endproc
.LFE2039:
	.size	_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif, .-_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif
	.globl	_Z60__device_stub__Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_iiPfS_S_S_ii
	.type	_Z60__device_stub__Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_iiPfS_S_S_ii, @function
_Z60__device_stub__Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_iiPfS_S_S_ii:
.LFB2040:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movl	%r9d, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L23:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_Z60__device_stub__Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_iiPfS_S_S_ii, .-_Z60__device_stub__Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_iiPfS_S_S_ii
	.globl	_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii
	.type	_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii, @function
_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii:
.LFB2041:
	.cfi_startproc
	endbr64
	jmp	_Z60__device_stub__Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_iiPfS_S_S_ii
	.cfi_endproc
.LFE2041:
	.size	_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii, .-_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii
	.globl	_Z42__device_stub__Z15kSoftMaxWithOffPfS_S_iiiPfS_S_iii
	.type	_Z42__device_stub__Z15kSoftMaxWithOffPfS_S_iiiPfS_S_iii, @function
_Z42__device_stub__Z15kSoftMaxWithOffPfS_S_iiiPfS_S_iii:
.LFB2042:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z15kSoftMaxWithOffPfS_S_iii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L28:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2042:
	.size	_Z42__device_stub__Z15kSoftMaxWithOffPfS_S_iiiPfS_S_iii, .-_Z42__device_stub__Z15kSoftMaxWithOffPfS_S_iiiPfS_S_iii
	.globl	_Z15kSoftMaxWithOffPfS_S_iii
	.type	_Z15kSoftMaxWithOffPfS_S_iii, @function
_Z15kSoftMaxWithOffPfS_S_iii:
.LFB2043:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z15kSoftMaxWithOffPfS_S_iiiPfS_S_iii
	.cfi_endproc
.LFE2043:
	.size	_Z15kSoftMaxWithOffPfS_S_iii, .-_Z15kSoftMaxWithOffPfS_S_iii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z15kSoftMaxWithOffPfS_S_iii"
.LC1:
	.string	"_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii"
.LC2:
	.string	"_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif"
.LC3:
	.string	"_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii"
.LC4:
	.string	"_Z18kMultinomialSamplePfS_S_iii"
.LC5:
	.string	"_Z14kCumulativeSumPfS_ii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2045:
	.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	_Z15kSoftMaxWithOffPfS_S_iii(%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	_Z32kSoftMaxWithOffAndPositionBiasesPfS_S_S_ii(%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	_Z30kSoftMaxReluWithPositionBiasesPfS_S_S_S_S_iiif(%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	_Z37kSoftMaxStackApproxWithPositionBiasesPfS_S_S_S_ii(%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	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z18kMultinomialSamplePfS_S_iii(%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	.LC5(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14kCumulativeSumPfS_ii(%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
.LFE2045:
	.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:

#pragma once
#include <hip/hip_runtime.h>


__global__ void kMultinomialSample(
                                   float *unifRandNums, 
                                   float *probabilities, 
                                   float *samples, 
                                   int softMaxWidth, 
                                   int startShift,
                                   int signalLength
                                  ) ; 

__global__ void kCumulativeSum(
                               float *probabilities, 
                               float *sumProbabilities, 
                               int softMaxWidth, 
                               int signalLength
                              ) ; 

__global__ void kSoftMaxWithOff(
                                float *activations, 
                                float *probabilities, 
                                float *biases,
                                int softMaxWidth,
                                int signalLength, 
                                int numPtsPerThread
                               ) ; 
__global__ void kSoftMaxWithOffAndPositionBiases(
                                                 float *activations, 
                                                 float *probabilities, 
                                                 float *featureBiases,
                                                 float *positionBiases,
                                                 int softMaxWidth, 
                                                 int signalLength
                                                ) ; 
__global__ void kSoftMaxStackApproxWithPositionBiases(
                                               float *activations, 
                                               float *probabilities, 
                                               float *stdevs, 
                                               float *featureBiases,
                                               float *positionBiases,
                                               int softMaxWidth, 
                                               int signalLength
                                               ) ; 

__global__ void kSoftMaxReluWithPositionBiases(
                                                 float *activations, 
                                                 float *probabilities, 
                                                 float *meanValues, 
                                                 float *featureStdevs, 
                                                 float *featureBiases,
                                                 float *positionBiases,
                                                 int softMaxWidth, 
                                                 int shift, 
                                                 int signalLength, 
                                                 float minExpForSum
                                              ) ; 


	.text
	.file	"SoftMax_kernels.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

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

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

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

	.end_amdgpu_metadata
