<commit_before>#ifndef TESTS_SIXTRACKLIB_OPENCL_TEST_PARTICLES_KERNEL_OPENCL_CL__
#define TESTS_SIXTRACKLIB_OPENCL_TEST_PARTICLES_KERNEL_OPENCL_CL__

#if !defined( SIXTRL_NO_INCLUDES )
    #include "sixtracklib/opencl/internal/default_compile_options.h"

    #include "sixtracklib/common/definitions.h"
    #include "sixtracklib/common/internal/buffer_main_defines.h"
    #include "sixtracklib/common/buffer/managed_buffer_minimal.h"
    #include "sixtracklib/common/buffer/managed_buffer_remap.h"
    #include "sixtracklib/common/particles.h"
#endif /* !defined( SIXTRL_NO_INCLUDES ) */

#pragma OPENCL_EXTENSION cl_khr_int32_extended_atomics

__kernel void NS(Particles_copy_buffer_opencl)(
    SIXTRL_BUFFER_DATAPTR_DEC unsigned char const* SIXTRL_RESTRICT in_buffer_begin,
    SIXTRL_BUFFER_DATAPTR_DEC unsigned char* SIXTRL_RESTRICT out_buffer_begin,
    SIXTRL_BUFFER_DATAPTR_DEC int* SIXTRL_RESTRICT ptr_success_flag )
{
    typedef NS(buffer_size_t) buf_size_t;
    typedef SIXTRL_BUFFER_OBJ_ARGPTR_DEC NS(Object)*       obj_iter_t;
    typedef SIXTRL_BUFFER_OBJ_ARGPTR_DEC NS(Object) const* obj_const_iter_t;

    typedef SIXTRL_PARTICLE_ARGPTR_DEC NS(Particles) const* ptr_const_particles_t;
    typedef SIXTRL_PARTICLE_ARGPTR_DEC NS(Particles)*       ptr_particles_t;

    int success_flag = ( int )0u;
    buf_size_t const slot_size = ( buf_size_t )8u;

    if( ( !NS(ManagedBuffer_needs_remapping)(  in_buffer_begin,  slot_size ) ) &&
        ( !NS(ManagedBuffer_needs_remapping)( out_buffer_begin,  slot_size ) ) &&
        (  NS(ManagedBuffer_get_num_objects)( out_buffer_begin, slot_size ) ==
           NS(ManagedBuffer_get_num_objects)(  in_buffer_begin, slot_size ) ) )
    {
        obj_const_iter_t in_obj_it  = NS(ManagedBuffer_get_const_objects_index_begin)(
            in_buffer_begin, slot_size );

        obj_const_iter_t in_obj_end = NS(ManagedBuffer_get_const_objects_index_end)(
            in_buffer_begin, slot_size );

        obj_iter_t out_obj_it = NS(ManagedBuffer_get_objects_index_begin)(
            out_buffer_begin, slot_size );

        buf_size_t obj_begin_index = ( buf_size_t )0u;

        size_t const global_id = get_global_id( 0 );

        for( ; in_obj_it != in_obj_end ; ++in_obj_it, ++out_obj_it )
        {
            ptr_const_particles_t in_particles = ( ptr_const_particles_t )(
                uintptr_t )NS(Object_get_begin_addr)( in_obj_it );

            ptr_particles_t out_particles = ( ptr_particles_t )(
                uintptr_t )NS(Object_get_begin_addr)( out_obj_it );

            buf_size_t const nn =
                NS(Particles_get_num_of_particles)( in_particles );

            buf_size_t const obj_end_index = obj_begin_index + nn;

            if( nn == NS(Particles_get_num_of_particles)( out_particles ) )
            {
                if( ( global_id >= obj_begin_index ) &&
                    ( global_id <  obj_end_index ) )
                {
                    buf_size_t const particle_id = global_id - obj_begin_index;

                    if( particle_id < nn )
                    {
                        NS(Particles_copy_single)( out_particles, particle_id,
                            in_particles, particle_id );
                    }
                    else
                    {
                        success_flag |= -1;
                    }

                    break;<commit_msg>testlib/opencl: add additional test in particle buffer copy kernel (directly)<commit_after>#ifndef TESTS_SIXTRACKLIB_OPENCL_TEST_PARTICLES_KERNEL_OPENCL_CL__
#define TESTS_SIXTRACKLIB_OPENCL_TEST_PARTICLES_KERNEL_OPENCL_CL__

#if !defined( SIXTRL_NO_INCLUDES )
    #include "sixtracklib/opencl/internal/default_compile_options.h"

    #include "sixtracklib/common/definitions.h"
    #include "sixtracklib/common/internal/buffer_main_defines.h"
    #include "sixtracklib/common/buffer/managed_buffer_minimal.h"
    #include "sixtracklib/common/buffer/managed_buffer_remap.h"
    #include "sixtracklib/common/particles.h"
#endif /* !defined( SIXTRL_NO_INCLUDES ) */

#pragma OPENCL_EXTENSION cl_khr_int32_extended_atomics

__kernel void NS(Particles_copy_buffer_opencl)(
    SIXTRL_BUFFER_DATAPTR_DEC unsigned char const* SIXTRL_RESTRICT in_buffer_begin,
    SIXTRL_BUFFER_DATAPTR_DEC unsigned char* SIXTRL_RESTRICT out_buffer_begin,
    SIXTRL_BUFFER_DATAPTR_DEC int* SIXTRL_RESTRICT ptr_success_flag )
{
    typedef NS(buffer_size_t) buf_size_t;
    typedef SIXTRL_BUFFER_OBJ_ARGPTR_DEC NS(Object)*       obj_iter_t;
    typedef SIXTRL_BUFFER_OBJ_ARGPTR_DEC NS(Object) const* obj_const_iter_t;

    typedef SIXTRL_PARTICLE_ARGPTR_DEC NS(Particles) const* ptr_const_particles_t;
    typedef SIXTRL_PARTICLE_ARGPTR_DEC NS(Particles)*       ptr_particles_t;

    int success_flag = ( int )0u;
    buf_size_t const slot_size = ( buf_size_t )8u;

    if( ( !NS(ManagedBuffer_needs_remapping)(  in_buffer_begin,  slot_size ) ) &&
        ( !NS(ManagedBuffer_needs_remapping)( out_buffer_begin,  slot_size ) ) &&
        (  NS(ManagedBuffer_get_num_objects)( out_buffer_begin, slot_size ) ==
           NS(ManagedBuffer_get_num_objects)(  in_buffer_begin, slot_size ) ) )
    {
        obj_const_iter_t in_obj_it  = NS(ManagedBuffer_get_const_objects_index_begin)(
            in_buffer_begin, slot_size );

        obj_const_iter_t in_obj_end = NS(ManagedBuffer_get_const_objects_index_end)(
            in_buffer_begin, slot_size );

        obj_iter_t out_obj_it = NS(ManagedBuffer_get_objects_index_begin)(
            out_buffer_begin, slot_size );

        buf_size_t obj_begin_index = ( buf_size_t )0u;

        size_t const global_id = get_global_id( 0 );

        for( ; in_obj_it != in_obj_end ; ++in_obj_it, ++out_obj_it )
        {
            ptr_const_particles_t in_particles = ( ptr_const_particles_t )(
                uintptr_t )NS(Object_get_begin_addr)( in_obj_it );

            ptr_particles_t out_particles = ( ptr_particles_t )(
                uintptr_t )NS(Object_get_begin_addr)( out_obj_it );

            buf_size_t const nn =
                NS(Particles_get_num_of_particles)( in_particles );

            buf_size_t const obj_end_index = obj_begin_index + nn;

            if( nn == NS(Particles_get_num_of_particles)( out_particles ) )
            {
                if( ( global_id >= obj_begin_index ) &&
                    ( global_id <  obj_end_index ) )
                {
                    buf_size_t const particle_id = global_id - obj_begin_index;

                    if( particle_id < nn )
                    {
                        NS(Particles_copy_single)( out_particles, particle_id,
                            in_particles, particle_id );

                        if( NS(Particles_get_state_value)( out_particles, particle_id ) !=
                            NS(Particles_get_state_value)( in_particles,  particle_id ) )
                        {
                            printf( "ERROR %d\r\n", ( int )particle_id );
                        }
                    }
                    else
                    {
                        success_flag |= -1;
                    }

                    break;
{'fraction_non_alphanumeric': 0.07326982347996638, 'fraction_numerical': 0.001961333706920706, 'mean_word_length': 2.6404895461499236, 'pattern_counts': {'":': 0, '<': 7, '<?xml version=': 0, '>': 5, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2893079', 'n_tokens_mistral': 2454, 'n_tokens_neox': 2376, 'n_words': 346}