cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

kaedo
Journeyman III

Kernel not found (Kernel with Subroutines and structs)

Hi folks,

I'm a little bit lost with my code. I programmed some other kernels before and they all works. That is my first advanced one.

I looked in the specifications and googled for using subroutines and structes in opencl and could not found any error.

But CodeXL Analyze Mode says that it could not find a kernel in the source file. But compilation finish without any errors.

I don't get any results when I run the kernel.

The Code is a porting of the Ahor Corasick algorithm from Multifast (MultiFast | Free Development software downloads at SourceForge.net)

OS: Arch Linux Kernel 3.11.6-1-x86_64

AMD Driver: catalyst-total 13.4-18


typedef struct


{


  uint  position;


  unsigned char null;


}offset;




typedef struct


{


  unsigned char alpha;


  offset oset;


}outgoing;




typedef struct


{


  uint id;


}matched_patterns;




typedef struct


{


  uint id;


  unsigned char final;


  uint depth;


  uint  failure_offset;


  unsigned char failure_NULL;


  uint matched_patterns_start;


  uint matched_patterns_degree;


  uint outgoing_start;


  uint outgoing_degree;


}offset_node;




typedef struct


{


  uint node_count;


  uint edge_count;


  offset current;


  unsigned short position;


  unsigned short base_position;


  uint total_patterns;


}offset_automata;




typedef struct


{


  uint patterns [10];


    uint positions[10];


    unsigned char   match_num;


}match_offset;




typedef char AC_ALPHABET_t;








offset node_find_next_offset   ( offset_node thiz, AC_ALPHABET_t alpha, global outgoing edges[])


{


    int i;


     offset of;




    for (i=0; i < thiz.outgoing_degree; i++)


    {


        if(edges[thiz.outgoing_start + i].alpha == alpha)


        {


        of.position = edges[thiz.outgoing_start + i].oset.position;


        of.null = 0;


        return of;


        }


    }


  of.position = 0;


  of.null = 1;




    return of;


}




offset node_findbs_next_offset ( offset_node thiz, AC_ALPHABET_t alpha, global outgoing edges[])


{


    int min, max, mid;


     offset of;


    AC_ALPHABET_t amid;




    min = 0;


    max = thiz.outgoing_degree - 1;




    while (min <= max)


    {


        mid = (min+max) >> 1;


        amid = edges[thiz.outgoing_start + mid].alpha;


        if (alpha > amid)


            min = mid + 1;


        else if (alpha < amid)


            max = mid - 1;


        else


        {


        of.position = edges[thiz.outgoing_start + mid].oset.position;


        of.null = 0;


        return of;


        }


    }


  of.position = 0;


  of.null = 1;




    return of;


}






__kernel void signature_matching( __global uchar *packets, /* all packets */


  __global short2 *packet_index, /* start and length of packet in buff*/


  __global matched_patterns *all_patterns, /* holds all patterns which occur in the tree */


  __global match_offset *g_match, /* match part 1 - used to save the matches */


  __global offset_node *node_array, /* all nodes in AC tree*/


  __global outgoing *edges_array, /* all edges in AC tree*/


  __global offset_automata *of_automata) /* the automata, hold some usefull infos*/


{


    uint row = get_global_id(0);




  local match_offset l_match;



    uint start  = packet_index[row].s0;


    uint length = packet_index[row].s1;






    unsigned char payload[1500];




    offset_automata thiz = *of_automata;




    unsigned short in_dex;




    for(in_dex = 0; in_dex < length; in_dex++)


    {


    payload[in_dex] = payload[start + in_dex];


    }




  offset current;


  offset next;




    unsigned short position = 0;


    unsigned char  i;




    current = thiz.current;




    l_match.match_num = 0;




    offset_node cur_of_node;






    while (position < length)


    {


    cur_of_node = node_array[current.position];


    next = node_findbs_next_offset(cur_of_node, payload[position], edges_array);


        if (next.null == 1)


        {


            if (cur_of_node.id != 1 )


  {


  current.position = cur_of_node.failure_offset;


  }else


                position++;


        }


        else


        {


            current = next;


            position++;


            cur_of_node = node_array[current.position];


        }




        if ((cur_of_node.final == 1) && (next.null != 1))


        {


 


        if(l_match.match_num  < 10)


        {


 


        for(i = 0; i < cur_of_node.matched_patterns_degree; i++)


        {




 


        l_match.positions[l_match.match_num + i] = position + thiz.base_position;




 


        l_match.patterns [l_match.match_num + i] = all_patterns[cur_of_node.matched_patterns_start + i].id;




        }


 


        l_match.match_num += cur_of_node.matched_patterns_degree;


        }else


        {


        break;


        }


        }


    }






    thiz.current = current;


    thiz.position = position;




    if (!l_match.match_num)


    {


  thiz.base_position+= position;


    }



};









0 Likes
1 Reply

Currently there is a known issue that will be addressed in one of the upcoming versions:

If there are comments in the parameters of the kernels then it might cause CodeXL not to recognize the kernel and not list it in the CL file's overview tab.

So if you remove the comments it will identify the kernel correctly.

0 Likes