
// This file is part of the LITIV framework; visit the original repository at
// https://github.com/plstcharles/litiv for more information.
//
// Copyright 2017 Pierre-Luc St-Charles; pierre-luc.st-charles<at>polymtl.ca
//
// 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.
//
// //////////////////////////////////////////////////////////////////////////
//
//               SLIC Superpixel Oversegmentation Algorithm
//       CUDA implementation of Achanta et al.'s method (TPAMI 2012)
//
// Note: requires CUDA compute architecture >= 3.0
// Author: Francois-Xavier Derue
// Contact: francois.xavier.derue@gmail.com
// Source: https://github.com/fderue/SLIC_CUDA
//
// Copyright (c) 2016 fderue
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
//

#include "SlicCuda.cuh"

__device__ inline float2 operator-(const float2 & a, const float2 & b) {
    return make_float2(a.x - b.x, a.y - b.y);
}

__device__ inline float3 operator-(const float3 & a, const float3 & b) {
    return make_float3(a.x - b.x, a.y - b.y, a.z - b.z);
}

__device__ inline int2 operator+(const int2 & a, const int2 & b) {
    return make_int2(a.x + b.x, a.y + b.y);
}

__device__ inline float computeDistance(float2 c_p_xy, float3 c_p_Lab, float areaSpx, float wc2) {
    float ds2 = pow(c_p_xy.x, 2) + pow(c_p_xy.y, 2);
    float dc2 = pow(c_p_Lab.x, 2) + pow(c_p_Lab.y, 2) + pow(c_p_Lab.z, 2);
    float dist = sqrt(dc2 + ds2/areaSpx*wc2);
    return dist;
}

/////////////////////////////////////////////////////////////////////////

__global__ void device::kRgb2CIELab(const cudaTextureObject_t texFrameBGRA, cudaSurfaceObject_t surfFrameLab, int width, int height) {

    int px = blockIdx.x*blockDim.x + threadIdx.x;
    int py = blockIdx.y*blockDim.y + threadIdx.y;

    if (px<width && py<height) {
        const uchar4 nPixel = tex2D<uchar4>(texFrameBGRA, px, py);//inputImg[offset];

        const float _b = nPixel.x/255.0f;
        const float _g = nPixel.y/255.0f;
        const float _r = nPixel.z/255.0f;

        float x = _r * 0.412453f + _g * 0.357580f + _b * 0.180423f;
        float y = _r * 0.212671f + _g * 0.715160f + _b * 0.072169f;
        float z = _r * 0.019334f + _g * 0.119193f + _b * 0.950227f;

        x /= 0.950456f;
        float y3 = exp(log(y)/3.0f);
        z /= 1.088754f;

        float l, a, b;

        x = x > 0.008856f ? exp(log(x)/3.0f) : (7.787f * x + 0.13793f);
        y = y > 0.008856f ? y3 : 7.787f * y + 0.13793f;
        z = z > 0.008856f ? z /= exp(log(z)/3.0f) : (7.787f * z + 0.13793f);

        l = y > 0.008856f ? (116.0f * y3 - 16.0f) : 903.3f * y;
        a = (x - y) * 500.0f;
        b = (y - z) * 200.0f;

        float4 fPixel;
        fPixel.x = l;
        fPixel.y = a;
        fPixel.z = b;
        fPixel.w = 0;

        surf2Dwrite(fPixel, surfFrameLab, px * 16, py);
    }
}

__global__ void device::kInitClusters(const cudaSurfaceObject_t surfFrameLab, float* clusters, int width, int height, int nSpxPerRow, int nSpxPerCol, float diamSpxD2) {
    int centroidIdx = blockIdx.x*blockDim.x + threadIdx.x;
    int nSpx = nSpxPerCol*nSpxPerRow;

    float diamSpx = diamSpxD2 * 2;
    if (centroidIdx<nSpx){
        int i = centroidIdx / nSpxPerRow;
        int j = centroidIdx%nSpxPerRow;

        int x = fminf(j*diamSpx + diamSpxD2, width-1);
        int y = fminf(i*diamSpx + diamSpxD2, height-1);

        float4 color;
        surf2Dread(&color, surfFrameLab, x * 16, y);
        clusters[centroidIdx] = color.x;
        clusters[centroidIdx + nSpx] = color.y;
        clusters[centroidIdx + 2 * nSpx] = color.z;
        clusters[centroidIdx + 3 * nSpx] = x;
        clusters[centroidIdx + 4 * nSpx] = y;
    }
}

__global__ void device::kAssignment(const cudaSurfaceObject_t surfFrameLab, const float* clusters, const int width, const int height, const int nClustPerRow, const int nbSpx,  const int diamSpx, const float wc2, cudaSurfaceObject_t surfLabels, float* accAtt_g) {
    // gather NNEIGH surrounding clusters
    const int NNEIGH = 3;
    __shared__ float4 sharedLab[NNEIGH][NNEIGH];
    __shared__ float2 sharedXY[NNEIGH][NNEIGH];

    int nn2 = NNEIGH / 2;
    if (threadIdx.x<NNEIGH && threadIdx.y<NNEIGH){
        int stepX = threadIdx.x - nn2; //[-1 0 1]
        int stepY = threadIdx.y - nn2; //[-1 0 1]

        int neighClusterX = blockIdx.x + stepX;
        int neighClusterY = blockIdx.y + stepY;

        int neighClusterLinIdx = neighClusterY*nClustPerRow + neighClusterX;

        if (neighClusterX < 0 || neighClusterX >= nClustPerRow ||
            neighClusterLinIdx < 0 || neighClusterLinIdx >= nbSpx ){
            sharedXY[threadIdx.y][threadIdx.x].x = -1;
        }
        else {
            sharedLab[threadIdx.y][threadIdx.x].x = clusters[neighClusterLinIdx];
            sharedLab[threadIdx.y][threadIdx.x].y = clusters[neighClusterLinIdx + nbSpx];
            sharedLab[threadIdx.y][threadIdx.x].z = clusters[neighClusterLinIdx + 2 * nbSpx];

            sharedXY[threadIdx.y][threadIdx.x].x = clusters[neighClusterLinIdx + 3 * nbSpx];
            sharedXY[threadIdx.y][threadIdx.x].y = clusters[neighClusterLinIdx + 4 * nbSpx];
        }
    }

    __syncthreads();

    // Find nearest neighbour
    float areaSpx = diamSpx*diamSpx;
    float distanceMin = 100000;
    float labelMin = -1;

    int px = blockIdx.x*blockDim.x + threadIdx.x;
    int py = blockIdx.y*blockDim.y*gridDim.z + blockIdx.z*blockDim.y + threadIdx.y;
    
    if (py<height && px<width){
        float4 color;
        surf2Dread(&color, surfFrameLab, px * 16, py);
        float3 px_Lab = make_float3(color.x, color.y, color.z);
        float2 px_xy = make_float2(px, py);
        for (int i = 0; i<NNEIGH; i++){
            int i_nn2 = i - nn2;
            for (int j = 0; j<NNEIGH; j++){
                if (sharedXY[i][j].x != -1){
                    float2 cluster_xy = make_float2(sharedXY[i][j].x, sharedXY[i][j].y);
                    float3 cluster_Lab = make_float3(sharedLab[i][j].x, sharedLab[i][j].y, sharedLab[i][j].z);

                    float2 px_c_xy = px_xy - cluster_xy;
                    float3 px_c_Lab = px_Lab - cluster_Lab;

                    float distTmp = fminf(computeDistance(px_c_xy, px_c_Lab, areaSpx, wc2), distanceMin);

                    if (distTmp != distanceMin){
                        distanceMin = distTmp;
                        labelMin = (blockIdx.y + i_nn2)*gridDim.x + blockIdx.x + (j - nn2);
                    }
                }
            }
        }
        
        surf2Dwrite(labelMin, surfLabels, px * 4, py);

        int iLabelMin = int(labelMin);
        atomicAdd(&accAtt_g[iLabelMin], px_Lab.x);
        atomicAdd(&accAtt_g[iLabelMin + nbSpx], px_Lab.y);
        atomicAdd(&accAtt_g[iLabelMin + 2 * nbSpx], px_Lab.z);
        atomicAdd(&accAtt_g[iLabelMin + 3 * nbSpx], px);
        atomicAdd(&accAtt_g[iLabelMin + 4 * nbSpx], py);
        atomicAdd(&accAtt_g[iLabelMin + 5 * nbSpx], 1); 
    }
}

__global__ void device::kUpdate(int nbSpx, float* clusters, float* accAtt_g) {
    int cluster_idx = blockIdx.x*blockDim.x + threadIdx.x;

    if (cluster_idx<nbSpx){
        int nbSpx2 = nbSpx * 2;
        int nbSpx3 = nbSpx * 3;
        int nbSpx4 = nbSpx * 4;
        int nbSpx5 = nbSpx * 5;
        float counter = accAtt_g[cluster_idx + nbSpx5];
        if (counter != 0){
            clusters[cluster_idx] = accAtt_g[cluster_idx] / counter;
            clusters[cluster_idx + nbSpx] = accAtt_g[cluster_idx + nbSpx] / counter;
            clusters[cluster_idx + nbSpx2] = accAtt_g[cluster_idx + nbSpx2] / counter;
            clusters[cluster_idx + nbSpx3] = accAtt_g[cluster_idx + nbSpx3] / counter;
            clusters[cluster_idx + nbSpx4] = accAtt_g[cluster_idx + nbSpx4] / counter;

            //reset accumulator
            accAtt_g[cluster_idx] = 0;
            accAtt_g[cluster_idx + nbSpx] = 0;
            accAtt_g[cluster_idx + nbSpx2] = 0;
            accAtt_g[cluster_idx + nbSpx3] = 0;
            accAtt_g[cluster_idx + nbSpx4] = 0;
            accAtt_g[cluster_idx + nbSpx5] = 0;
        }
    }
}

/////////////////////////////////////////////////////////////////////////

void host::kRgb2CIELab(const KernelParams& oKParams, const cudaTextureObject_t inputImg, cudaSurfaceObject_t outputImg, int width, int height) {
    cudaKernelWrap(kRgb2CIELab,oKParams,inputImg,outputImg,width,height);
}

void host::kInitClusters(const KernelParams& oKParams, const cudaSurfaceObject_t frameLab, float* clusters, int width, int height, int nSpxPerRow, int nSpxPerCol, float diamSpxD2) {
    cudaKernelWrap(kInitClusters,oKParams,frameLab,clusters,width,height,nSpxPerRow,nSpxPerCol,diamSpxD2);
}

void host::kAssignment(const KernelParams& oKParams,const cudaSurfaceObject_t frameLab, const float* clusters, const int width, const int height, const int nClustPerRow, const int nbSpx, const int diamSpx, const float wc2, cudaSurfaceObject_t labels, float* accAtt_g) {
    cudaKernelWrap(kAssignment,oKParams,frameLab,clusters,width,  height, nClustPerRow, nbSpx,diamSpx,wc2,labels,accAtt_g);
}

void host::kUpdate(const KernelParams& oKParams, int nbSpx, float* clusters, float* accAtt_g) {
    cudaKernelWrap(kUpdate,oKParams,nbSpx,clusters,accAtt_g);
}
	.file	"tmpxft_002da8f4_00000000-6_SlicCuda.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2043:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2043:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii
	.type	_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii, @function
_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii:
.LFB2065:
	.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	_ZN6device11kRgb2CIELabEyyii(%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
.LFE2065:
	.size	_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii, .-_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii
	.globl	_ZN6device11kRgb2CIELabEyyii
	.type	_ZN6device11kRgb2CIELabEyyii, @function
_ZN6device11kRgb2CIELabEyyii:
.LFB2066:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii
	.cfi_endproc
.LFE2066:
	.size	_ZN6device11kRgb2CIELabEyyii, .-_ZN6device11kRgb2CIELabEyyii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/fderue/SPIKES/master/module/SLIC_CUDA/cuda/SlicCuda.cu"
.LC1:
	.string	"void host::kRgb2CIELab(const KernelParams&, cudaTextureObject_t, cudaSurfaceObject_t, int, int)"
.LC2:
	.string	"cuda kernel 'kRgb2CIELab' execution failed [code=%d, msg=%s]\n\t... in function '%s' from %s(%d)\n"
	.text
	.globl	_ZN4host11kRgb2CIELabERK12KernelParamsyyii
	.type	_ZN4host11kRgb2CIELabERK12KernelParamsyyii, @function
_ZN4host11kRgb2CIELabERK12KernelParamsyyii:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rdi, %rax
	movl	%r8d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rsi, %rbx
	pushq	%r9
	.cfi_def_cfa_offset 48
	movq	12(%rdi), %rdx
	movl	20(%rdi), %ecx
	movl	8(%rax), %esi
	movq	(%rdi), %rdi
	movq	32(%rax), %r9
	movq	24(%rax), %r8
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movl	%r13d, %ecx
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z42__device_stub__ZN6device11kRgb2CIELabEyyiiyyii
.L9:
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L8
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC0(%rip), %r9
	leaq	.LC1(%rip), %r8
	pushq	%rdx
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movq	%rax, %rcx
	movl	%ebx, %edx
	leaq	.LC2(%rip), %rsi
	pushq	$241
	.cfi_def_cfa_offset 64
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%ebx, %edi
	call	exit@PLT
.L8:
	.cfi_restore_state
	popq	%rax
	.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
.LFE2037:
	.size	_ZN4host11kRgb2CIELabERK12KernelParamsyyii, .-_ZN4host11kRgb2CIELabERK12KernelParamsyyii
	.globl	_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif
	.type	_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif, @function
_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif:
.LFB2067:
	.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
	movl	%edx, 28(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 24(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 20(%rsp)
	movl	%r9d, 16(%rsp)
	movss	%xmm0, 12(%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	28(%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, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_ZN6device13kInitClustersEyPfiiiif(%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
.L12:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2067:
	.size	_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif, .-_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif
	.globl	_ZN6device13kInitClustersEyPfiiiif
	.type	_ZN6device13kInitClustersEyPfiiiif, @function
_ZN6device13kInitClustersEyPfiiiif:
.LFB2068:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif
	.cfi_endproc
.LFE2068:
	.size	_ZN6device13kInitClustersEyPfiiiif, .-_ZN6device13kInitClustersEyPfiiiif
	.section	.rodata.str1.1
.LC3:
	.string	"void host::kInitClusters(const KernelParams&, cudaSurfaceObject_t, float*, int, int, int, int, float)"
.LC4:
	.string	"cuda kernel 'kInitClusters' execution failed [code=%d, msg=%s]\n\t... in function '%s' from %s(%d)\n"
	.text
	.globl	_ZN4host13kInitClustersERK12KernelParamsyPfiiiif
	.type	_ZN4host13kInitClustersERK12KernelParamsyPfiiiif, @function
_ZN4host13kInitClustersERK12KernelParamsyPfiiiif:
.LFB2038:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %rax
	movl	%r9d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%r8d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rsi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 64
	movq	12(%rdi), %rdx
	movl	20(%rdi), %ecx
	movl	8(%rax), %esi
	movq	(%rdi), %rdi
	movss	%xmm0, 12(%rsp)
	movq	32(%rax), %r9
	movq	24(%rax), %r8
	call	__cudaPushCallConfiguration@PLT
	movss	12(%rsp), %xmm0
	testl	%eax, %eax
	jne	.L18
	movl	64(%rsp), %r9d
	movl	%r14d, %r8d
	movl	%r13d, %ecx
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z48__device_stub__ZN6device13kInitClustersEyPfiiiifyPfiiiif
.L18:
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L17
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC0(%rip), %r9
	movl	%ebx, %edx
	pushq	%rsi
	.cfi_remember_state
	.cfi_def_cfa_offset 72
	movq	%rax, %rcx
	leaq	.LC3(%rip), %r8
	xorl	%eax, %eax
	pushq	$245
	.cfi_def_cfa_offset 80
	leaq	.LC4(%rip), %rsi
	call	__printf_chk@PLT
	movl	%ebx, %edi
	call	exit@PLT
.L17:
	.cfi_restore_state
	addq	$16, %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
.LFE2038:
	.size	_ZN4host13kInitClustersERK12KernelParamsyPfiiiif, .-_ZN4host13kInitClustersERK12KernelParamsyPfiiiif
	.globl	_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf
	.type	_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf, @function
_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf:
.LFB2069:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	240(%rsp), %rax
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movl	%edx, 28(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 24(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 20(%rsp)
	movl	%r9d, 16(%rsp)
	movq	%rax, (%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%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	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	pushq	64(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_ZN6device11kAssignmentEyPKfiiiiifyPf(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 240
	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 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L21:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2069:
	.size	_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf, .-_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf
	.globl	_ZN6device11kAssignmentEyPKfiiiiifyPf
	.type	_ZN6device11kAssignmentEyPKfiiiiifyPf, @function
_ZN6device11kAssignmentEyPKfiiiiifyPf:
.LFB2070:
	.cfi_startproc
	endbr64
	jmp	_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf
	.cfi_endproc
.LFE2070:
	.size	_ZN6device11kAssignmentEyPKfiiiiifyPf, .-_ZN6device11kAssignmentEyPKfiiiiifyPf
	.section	.rodata.str1.1
.LC5:
	.string	"void host::kAssignment(const KernelParams&, cudaSurfaceObject_t, const float*, int, int, int, int, int, float, cudaSurfaceObject_t, float*)"
.LC6:
	.string	"cuda kernel 'kAssignment' execution failed [code=%d, msg=%s]\n\t... in function '%s' from %s(%d)\n"
	.text
	.globl	_ZN4host11kAssignmentERK12KernelParamsyPKfiiiiifyPf
	.type	_ZN4host11kAssignmentERK12KernelParamsyPKfiiiiifyPf, @function
_ZN4host11kAssignmentERK12KernelParamsyPKfiiiiifyPf:
.LFB2039:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %rax
	movl	%r9d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%r8d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rsi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 64
	movq	12(%rdi), %rdx
	movl	20(%rdi), %ecx
	movl	8(%rax), %esi
	movq	(%rdi), %rdi
	movss	%xmm0, 12(%rsp)
	movq	32(%rax), %r9
	movq	24(%rax), %r8
	call	__cudaPushCallConfiguration@PLT
	movss	12(%rsp), %xmm0
	testl	%eax, %eax
	jne	.L27
	pushq	%rdi
	.cfi_def_cfa_offset 72
	movl	%r14d, %r8d
	movl	%r13d, %ecx
	movl	%r12d, %edx
	pushq	96(%rsp)
	.cfi_def_cfa_offset 80
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	pushq	96(%rsp)
	.cfi_def_cfa_offset 88
	movl	96(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 96
	movl	96(%rsp), %r9d
	call	_Z51__device_stub__ZN6device11kAssignmentEyPKfiiiiifyPfyPKfiiiiifyPf
	addq	$32, %rsp
	.cfi_def_cfa_offset 64
.L27:
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L26
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC0(%rip), %r9
	movl	%ebx, %edx
	pushq	%rsi
	.cfi_remember_state
	.cfi_def_cfa_offset 72
	movq	%rax, %rcx
	leaq	.LC5(%rip), %r8
	xorl	%eax, %eax
	pushq	$249
	.cfi_def_cfa_offset 80
	leaq	.LC6(%rip), %rsi
	call	__printf_chk@PLT
	movl	%ebx, %edi
	call	exit@PLT
.L26:
	.cfi_restore_state
	addq	$16, %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
.LFE2039:
	.size	_ZN4host11kAssignmentERK12KernelParamsyPKfiiiiifyPf, .-_ZN4host11kAssignmentERK12KernelParamsyPKfiiiiifyPf
	.globl	_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_
	.type	_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_, @function
_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_:
.LFB2071:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 28(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_ZN6device7kUpdateEiPfS0_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L30:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
	call	__stack_chk_fail@PLT
.L32:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2071:
	.size	_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_, .-_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_
	.globl	_ZN6device7kUpdateEiPfS0_
	.type	_ZN6device7kUpdateEiPfS0_, @function
_ZN6device7kUpdateEiPfS0_:
.LFB2072:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_
	.cfi_endproc
.LFE2072:
	.size	_ZN6device7kUpdateEiPfS0_, .-_ZN6device7kUpdateEiPfS0_
	.section	.rodata.str1.1
.LC7:
	.string	"void host::kUpdate(const KernelParams&, int, float*, float*)"
.LC8:
	.string	"cuda kernel 'kUpdate' execution failed [code=%d, msg=%s]\n\t... in function '%s' from %s(%d)\n"
	.text
	.globl	_ZN4host7kUpdateERK12KernelParamsiPfS3_
	.type	_ZN4host7kUpdateERK12KernelParamsiPfS3_, @function
_ZN4host7kUpdateERK12KernelParamsiPfS3_:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rcx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	(%rdi), %rax
	movl	%esi, %ebx
	movq	12(%rdi), %rdx
	movl	20(%rdi), %ecx
	movl	8(%rdi), %esi
	movq	32(%rdi), %r9
	movq	24(%rdi), %r8
	movq	%rax, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	%ebx, %edi
	call	_Z39__device_stub__ZN6device7kUpdateEiPfS0_iPfS_
.L36:
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L35
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC0(%rip), %r9
	movl	%ebx, %edx
	movq	%rax, %rcx
	pushq	%rax
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	leaq	.LC7(%rip), %r8
	pushq	$253
	.cfi_def_cfa_offset 48
	leaq	.LC8(%rip), %rsi
	call	__printf_chk@PLT
	movl	%ebx, %edi
	call	exit@PLT
.L35:
	.cfi_restore_state
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	_ZN4host7kUpdateERK12KernelParamsiPfS3_, .-_ZN4host7kUpdateERK12KernelParamsiPfS3_
	.section	.rodata.str1.1
.LC9:
	.string	"_ZN6device7kUpdateEiPfS0_"
.LC10:
	.string	"_ZN6device11kAssignmentEyPKfiiiiifyPf"
.LC11:
	.string	"_ZN6device13kInitClustersEyPfiiiif"
.LC12:
	.string	"_ZN6device11kRgb2CIELabEyyii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2074:
	.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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_ZN6device7kUpdateEiPfS0_(%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	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_ZN6device11kAssignmentEyPKfiiiiifyPf(%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	.LC11(%rip), %rdx
	orl	$-1, %r8d
	leaq	_ZN6device13kInitClustersEyPfiiiif(%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	.LC12(%rip), %rdx
	orl	$-1, %r8d
	leaq	_ZN6device11kRgb2CIELabEyyii(%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
.LFE2074:
	.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:


// This file is part of the LITIV framework; visit the original repository at
// https://github.com/plstcharles/litiv for more information.
//
// Copyright 2017 Pierre-Luc St-Charles; pierre-luc.st-charles<at>polymtl.ca
//
// 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.
//
// //////////////////////////////////////////////////////////////////////////
//
//               SLIC Superpixel Oversegmentation Algorithm
//       CUDA implementation of Achanta et al.'s method (TPAMI 2012)
//
// Note: requires CUDA compute architecture >= 3.0
// Author: Francois-Xavier Derue
// Contact: francois.xavier.derue@gmail.com
// Source: https://github.com/fderue/SLIC_CUDA
//
// Copyright (c) 2016 fderue
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
//

#pragma once

#ifdef __HIPCC__
#include <cstdio>
#define cudaKernelWrap(func,kparams,...) do { \
    device::func<<<kparams.vGridSize,kparams.vBlockSize,kparams.nSharedMemSize,kparams.nStream>>>(__VA_ARGS__); \
    const hipError_t __errn = hipGetLastError(); \
    if(__errn!=hipSuccess) { \
        printf("cuda kernel '" #func "' execution failed [code=%d, msg=%s]\n\t... in function '%s' from %s(%d)\n",(int)__errn,hipGetErrorString(__errn),__PRETTY_FUNCTION__,__FILE__,__LINE__); \
        std::exit((int)__errn); \
    } \
} while(0)
#endif //def(__CUDACC__)
#include <hip/hip_runtime.h>

#define cudaErrorCheck(test) do { \
        const hipError_t __errn = test; \
        if(__errn!=hipSuccess) { \
            std::cout<<"error in cuda function"<<std::endl; \
            std::exit((int)__errn); \
        } \
    } while(0)

struct KernelParams {
    KernelParams() :
            vGridSize(0),vBlockSize(0),nSharedMemSize(0),nStream(0) {}
    KernelParams(const dim3& _vGridSize, const dim3& _vBlockSize, size_t _nSharedMemSize=0, hipStream_t _nStream=0) :
            vGridSize(_vGridSize),vBlockSize(_vBlockSize),nSharedMemSize(_nSharedMemSize),nStream(_nStream) {}
    dim3 vGridSize;
    dim3 vBlockSize;
    size_t nSharedMemSize;
    hipStream_t nStream;
};

namespace device {

    __global__ void kRgb2CIELab(const hipTextureObject_t inputImg, hipSurfaceObject_t outputImg, int width, int height);
    __global__ void kInitClusters(const hipSurfaceObject_t frameLab, float* clusters, int width, int height, int nSpxPerRow, int nSpxPerCol, float nbSpxD2);
    __global__ void kAssignment(const hipSurfaceObject_t surfFrameLab, const float* clusters, const int width, const int height, const int nClustPerRow, const int nbSpx, const int diamSpx, const float wc2, hipSurfaceObject_t surfLabels, float* accAtt_g);
    __global__ void kUpdate(int nbSpx, float* clusters, float* accAtt_g);

} // namespace device

namespace host {

    void kRgb2CIELab(const KernelParams& oKParams, const hipTextureObject_t inputImg, hipSurfaceObject_t outputImg, int width, int height);
    void kInitClusters(const KernelParams& oKParams, const hipSurfaceObject_t frameLab, float* clusters, int width, int height, int nSpxPerRow, int nSpxPerCol,float diamSpxD2);
    void kAssignment(const KernelParams& oKParams, const hipSurfaceObject_t surfFrameLab, const float* clusters, const int width, const int height, const int nClustPerRow, const int nbSpx, const int diamSpx, const float wc2, hipSurfaceObject_t surfLabels, float* accAtt_g);
    void kUpdate(const KernelParams& oKParams, int nbSpx, float* clusters, float* accAtt_g);

} // namespace host
	.text
	.file	"SlicCuda.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
