#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
                      CLK_ADDRESS_CLAMP | //Clamp to zeros
                      CLK_FILTER_NEAREST; //Don't interpolate

__inline void PreSystem(int tsysSystemType, double* tsysNumberProperties)
{{
{1}
}}

__inline void PostInteraction(int tsysSystemType, double* tsysNumberProperties, double* fsysNumberProperties)
{{
{2}
}}

__inline void PostSystem(int tsysSystemType, double* tsysNumberProperties)
{{
{3}
}}

__inline void Interaction(int tsysSystemType, int fsysSystemType, double* tsysNumberProperties, double* fsysNumberProperties)
{{
{4}
}}

__kernel void PreInteractionForward(__global int* systemTypeBuffer, __read_only image3d_t numberPropertiesBuffer1, __write_only image3d_t numberPropertiesBuffer2)
{{
    int tsys = get_global_id(0);
    int size = get_global_size(0);

    double tsysNumberProperties[{0}];

    for(int i = 0; i < {0}; i++)
    {{
        tsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer1, smp, (int4)(tsys, i, tsys, 0))).x;
    }}

    PreSystem(systemTypeBuffer[tsys], tsysNumberProperties);

    for(int i = 0; i < {0}; i++)
    {{
        write_imagei(numberPropertiesBuffer2, (int4)(tsys, i, tsys, 0), as_int4((double2)(tsysNumberProperties[i], 0.0)));
    }}
}}

__kernel void PostInteractionForward(__global int* systemTypeBuffer, __read_only image3d_t numberPropertiesBuffer1, __write_only image3d_t numberPropertiesBuffer2)
{{
    int tsys = get_global_id(0);
    int size = get_global_size(0);

    double tsysNumberProperties[{0}];
    double fsysNumberProperties[{0}];

    for(int i = 0; i < {0}; i++)
    {{
        tsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer1, smp, (int4)(tsys, i, tsys, 0))).x;
    }}

    for(int fsys = 0; fsys < size; fsys++)
    {{
        if(tsys != fsys)
        {{
            for (int i = 0; i < {0}; i++)
            {{
                fsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer1, smp, (int4)(tsys, i, fsys, 0))).x;
            }}

            PostInteraction(systemTypeBuffer[tsys], tsysNumberProperties, fsysNumberProperties);

            for(int i = 0; i < {0}; i++)
            {{
                write_imagei(numberPropertiesBuffer2, (int4)(tsys, i, fsys, 0), as_int4((double2)(fsysNumberProperties[i], 0.0)));
            }}
        }}
    }}

    PostSystem(systemTypeBuffer[tsys], tsysNumberProperties);

    for(int i = 0; i < {0}; i++)
    {{
        write_imagei(numberPropertiesBuffer2, (int4)(tsys, i, tsys, 0), as_int4((double2)(tsysNumberProperties[i], 0.0)));
    }}
}}

__kernel void InteractionForward(__global int* systemTypeBuffer, __read_only image3d_t numberPropertiesBuffer1, __write_only image3d_t numberPropertiesBuffer2)
{{
    int tsys = get_global_id(0);
    int fsys = get_global_id(1);

    if(tsys != fsys)
    {{
        double tsysNumberProperties[{0}];
        double fsysNumberProperties[{0}];

        for(int i = 0; i < {0}; i++)
        {{
            tsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer1, smp, (int4)(tsys, i, tsys, 0))).x;
            fsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer1, smp, (int4)(fsys, i, fsys, 0))).x;
        }}

        Interaction(systemTypeBuffer[tsys], systemTypeBuffer[fsys], tsysNumberProperties, fsysNumberProperties);

        for(int i = 0; i < {0}; i++)
        {{
            write_imagei(numberPropertiesBuffer2, (int4)(tsys, i, fsys, 0), as_int4((double2)(tsysNumberProperties[i], 0.0)));
            write_imagei(numberPropertiesBuffer2, (int4)(fsys, i, fsys, 0), as_int4((double2)(fsysNumberProperties[i], 0.0)));
        }}
    }}
}}

__kernel void PreInteractionReverse(__global int* systemTypeBuffer, __write_only image3d_t numberPropertiesBuffer1, __read_only image3d_t numberPropertiesBuffer2)
{{
    int tsys = get_global_id(0);
    int size = get_global_size(0);

    double tsysNumberProperties[{0}];

    for(int i = 0; i < {0}; i++)
    {{
        tsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer2, smp, (int4)(tsys, i, tsys, 0))).x;
    }}

    PreSystem(systemTypeBuffer[tsys], tsysNumberProperties);

    for(int i = 0; i < {0}; i++)
    {{
        write_imagei(numberPropertiesBuffer1, (int4)(tsys, i, tsys, 0), as_int4((double2)(tsysNumberProperties[i], 0.0)));
    }}
}}

__kernel void PostInteractionReverse(__global int* systemTypeBuffer, __write_only image3d_t numberPropertiesBuffer1, __read_only image3d_t numberPropertiesBuffer2)
{{
    int tsys = get_global_id(0);
    int size = get_global_size(0);

    double tsysNumberProperties[{0}];
    double fsysNumberProperties[{0}];

    for(int i = 0; i < {0}; i++)
    {{
        tsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer2, smp, (int4)(tsys, i, tsys, 0))).x;
    }}

    for(int fsys = 0; fsys < size; fsys++)
    {{
        if(tsys != fsys)
        {{
            for (int i = 0; i < {0}; i++)
            {{
                fsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer2, smp, (int4)(tsys, i, fsys, 0))).x;
            }}

            PostInteraction(systemTypeBuffer[tsys], tsysNumberProperties, fsysNumberProperties);

            for(int i = 0; i < {0}; i++)
            {{
                write_imagei(numberPropertiesBuffer1, (int4)(tsys, i, fsys, 0), as_int4((double2)(fsysNumberProperties[i], 0.0)));
            }}
        }}
    }}

    PostSystem(systemTypeBuffer[tsys], tsysNumberProperties);

    for(int i = 0; i < {0}; i++)
    {{
        write_imagei(numberPropertiesBuffer1, (int4)(tsys, i, tsys, 0), as_int4((double2)(tsysNumberProperties[i], 0.0)));
    }}
}}

__kernel void InteractionReverse(__global int* systemTypeBuffer, __write_only image3d_t numberPropertiesBuffer1, __read_only image3d_t numberPropertiesBuffer2)
{{
    int tsys = get_global_id(0);
    int fsys = get_global_id(1);

    if(tsys != fsys)
    {{
        double tsysNumberProperties[{0}];
        double fsysNumberProperties[{0}];

        for(int i = 0; i < {0}; i++)
        {{
            tsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer2, smp, (int4)(tsys, i, tsys, 0))).x;
            fsysNumberProperties[i] = as_double2(read_imagei(numberPropertiesBuffer2, smp, (int4)(fsys, i, fsys, 0))).x;
        }}

        Interaction(systemTypeBuffer[tsys], systemTypeBuffer[fsys], tsysNumberProperties, fsysNumberProperties);

        for(int i = 0; i < {0}; i++)
        {{
            write_imagei(numberPropertiesBuffer1, (int4)(tsys, i, fsys, 0), as_int4((double2)(tsysNumberProperties[i], 0.0)));
            write_imagei(numberPropertiesBuffer1, (int4)(fsys, i, fsys, 0), as_int4((double2)(fsysNumberProperties[i], 0.0)));
        }}
    }}
}}