1 Reply Latest reply on Nov 17, 2013 2:31 AM by Gilad.Yarnitzky

    Kernel not found (Kernel with Subroutines and structs)

    kaedo

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