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