AnsweredAssumed Answered

Kernel not found (Kernel with Subroutines and structs)

Question asked by kaedo on Nov 15, 2013
Latest reply on Nov 17, 2013 by Gilad.Yarnitzky

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;
    }

};







Outcomes