AnsweredAssumed Answered

Weird behavior of barrier()

Question asked by t-man on Sep 5, 2012
Latest reply on Sep 6, 2012 by notzed
"__kernel void square(                                                   \n" \
"   __global int* nbr,                                          \n" \
"   __global double* BC,                                         \n" \
"   const int ne,                                       \n" \
"   const int nv, \n" \
"   __global int* firstnbr, \n" \
"   __global double* sigma,\n" \
"   __global int* queue, \n" \
"   __global int* queue2, \n" \
"   __global int* nr_queue, \n" \
"   __global int* nr_queue2, \n" \
"   __global int* stack, \n" \
"   __global int* nr_stack, \n" \
"   __global int* succ, \n" \
"   __global int* level, \n" \
"   __global int* nr_level, \n" \
"   __global int *succ_index, \n" \
"   __global int* found, \n" \
"   __global double* delta, \n" \
"  __global int* d, \n" \
"   __global int* sem)\n" \
"{                                                                  \n" \
"   int i = get_local_id(0);                                      \n" \
"int glob = get_global_id(0); \n" \
"   int j = get_num_groups(0);\n" \
" __local int k; k = get_group_id(0); \n" \
"  __local int size; size = get_local_size(0); \n"\
"   sem[0] = 0; \n" \
" \n" \
"   *found = 1; \n" \
   " __local int nr_roots, count,rest,root,nr_neigh; \n " \
   " int neigh_per_thread,h,node,dw,f,gh,temp,nr=0; \n " \
   " queue[0] = 0; \n" \
   " //barrier(CLK_LOCAL_MEM_FENCE); \n" \
   " while ( *found != 0){ \n" \
    "    *found = 0;\n" \
    "  \n" \
     "   //barrier(CLK_LOCAL_MEM_FENCE); \n " \
     "BC[0] = 1; BC[1] = *nr_level%2;\n" \
   "     if(*nr_level%2 == 0) \n" \
    "    {      \n" \
    "            nr_roots = (*nr_queue+1)/j; count = 0; rest = (*nr_queue+1)%j; \n" \
     "           if(k<rest  && i==0) \n" \
      "                  nr_roots=nr_roots + 1; \n" \
       "         *nr_queue2 = 0; if(k==0&&i==0) BC[5] = nr_roots;\n" \
        "        barrier(CLK_LOCAL_MEM_FENCE);   if(k==0&&i==0) BC[6] = nr_roots; \n" \

 

}

}

So what happends is that I need each workgroup to take an equal number of elements from a queue, that has nr_queue elements. 1st step it has 1 element so only group 0 will get 1 element. Thus "nr_roots" becomes 1 and I put it in BC[5]. Then I do a barrier to make sure that all the workitems on the workgroup know that nr_roots is 1, but when I out nr_roots in BC[6] the result is 0. Any1 has any idea why this might be? Thanks!

Outcomes