Advertisement
Guest User

Untitled

a guest
Jan 23rd, 2013
108
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. __kernel void MinMax(  __global float * input,
  2.                         __global float * output,
  3.                         int sizex,
  4.                         int sizey,
  5.                         int sizez,
  6.                         int size,
  7.                         int local_size
  8.                     )
  9. {
  10.     int idx = get_global_id(0);
  11.     int idy = get_global_id(1);
  12.     int idz = get_global_id(2);
  13.  
  14.     int idx_local = get_local_id(0);
  15.     int idy_local = get_local_id(1);
  16.     int idz_local = get_local_id(2);
  17.  
  18.     int local_x = get_local_size(0);
  19.     int local_y = get_local_size(1);
  20.     int local_z = get_local_size(2);
  21.  
  22.     int gid = idx + sizex*idy + sizex*sizey*idz;
  23.     int gid_local = idx_local + local_x*idy_local +local_x*local_y*idz_local;
  24.     int group_count = get_num_groups(0) * get_num_groups(1) * get_num_groups(2);
  25.  
  26.     //int group_id = groupid_x + group_y*groupid_y + groupid_z*group_z*groupid_z;
  27.  
  28.  
  29.  
  30.     if((gid<size) && (gid_local<local_size))
  31.     {
  32.         __local float Minvals[ 64 ];
  33.         __local float Maxvals[ 64 ];
  34.  
  35.     //local memory
  36.         Minvals[gid_local] = input[gid];
  37.         Maxvals[gid_local] = input[gid];
  38.  
  39.  
  40.         int j = gid + gid_local + group_count;
  41.         float ass = Minvals[gid_local] + Maxvals[gid_local];
  42.  
  43.         int offset = (local_x*local_y*local_z)/2;
  44.         //Local reduction
  45.         for(;gid_local<offset;offset/=2)
  46.         {
  47.             Minvals[gid_local] = fmin(Minvals[gid_local],Minvals[gid_local+offset]);
  48.             Maxvals[gid_local] = fmax(Maxvals[gid_local],Maxvals[gid_local+offset]);
  49.         }
  50.         barrier(CLK_GLOBAL_MEM_FENCE);
  51.         //Global rueckschreiben
  52.         if(gid_local == 0)
  53.         {
  54.             //global Rueckschreiben
  55.             output[gid*2] =  Minvals[gid_local];
  56.             output[gid*2 +1] = Maxvals[gid_local];
  57.         }
  58.         barrier(CLK_GLOBAL_MEM_FENCE);
  59.         //Global Reduction
  60.         if(gid==0)
  61.         {
  62.             float mint=256;
  63.             float maxt=0;
  64.             for(int i =0; i < group_count; i=i+2)
  65.             {
  66.                 mint    = fmin(mint,output[i]);
  67.                 maxt    = fmax(maxt,output[i+1]);
  68.             }
  69.             output[0] = mint;
  70.             output[1] = maxt;
  71.         }
  72.     }
  73. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement