compute-runtime/shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel

/*
 * Copyright (C) 2018-2025 Intel Corporation
 *
 * SPDX-License-Identifier: MIT
 *
 */

R"===(

__kernel void CopyBufferRectBytes2d(
    __global const char* src,
    __global char* dst,
    uint2 SrcOrigin,
    uint2 DstOrigin,
    uint SrcPitch,
    uint DstPitch )

{
    int x = get_global_id(0);
    int y = get_global_id(1);

    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );

    *( dst + LDstOffset )  = *( src + LSrcOffset ); 

}

__kernel void CopyBufferRectBytesMiddle2d(
    const __global uint* src,
    __global uint* dst,
    uint2 SrcOrigin,
    uint2 DstOrigin,
    uint SrcPitch,
    uint DstPitch )

{
    int x = get_global_id(0);
    int y = get_global_id(1);

    uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch );
    uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch );

    src += LSrcOffset >> 2;
    dst += LDstOffset >> 2;
    
    uint4 loaded = vload4(x,src);
    vstore4(loaded,x,dst);
}

__kernel void CopyBufferRectBytes3d(
    __global const char* src, 
    __global char* dst, 
    uint4 SrcOrigin, 
    uint4 DstOrigin, 
    uint2 SrcPitch, 
    uint2 DstPitch ) 
 
{ 
    int x = get_global_id(0); 
    int y = get_global_id(1); 
    int z = get_global_id(2); 
 
    uint LSrcOffset = x + SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ) + ( ( z + SrcOrigin.z ) * SrcPitch.y ); 
    uint LDstOffset = x + DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ) + ( ( z + DstOrigin.z ) * DstPitch.y ); 
 
    *( dst + LDstOffset )  = *( src + LSrcOffset );  
 
}

__kernel void CopyBufferRectBytesMiddle3d(
    const __global uint* src,
    __global uint* dst,
    uint4 SrcOrigin,
    uint4 DstOrigin,
    uint2 SrcPitch,
    uint2 DstPitch )

{
    int x = get_global_id(0); 
    int y = get_global_id(1); 
    int z = get_global_id(2); 
 
    uint LSrcOffset = SrcOrigin.x + ( ( y + SrcOrigin.y ) * SrcPitch.x ) + ( ( z + SrcOrigin.z ) * SrcPitch.y ); 
    uint LDstOffset = DstOrigin.x + ( ( y + DstOrigin.y ) * DstPitch.x ) + ( ( z + DstOrigin.z ) * DstPitch.y ); 

    src += LSrcOffset >> 2;
    dst += LDstOffset >> 2;
    
    uint4 loaded = vload4(x,src);
    vstore4(loaded,x,dst);
}
)==="