Advertisement
Guest User

minmax

a guest
Feb 10th, 2013
70
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C++ 2.63 KB | None
  1. #define WORKGROUPSIZE 64
  2. kernel void MinMax( global float * input,
  3.             global float * output,
  4.                         int sizex,
  5.             int sizey,
  6.             int sizez,
  7.             int 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.         int local_size = local_x*local_y*local_z;
  22.    
  23.         int gid = idx + sizex*idy + sizex*sizey*idz;
  24.         int gid_local = idx_local + local_x*idy_local +local_x*local_y*idz_local;
  25.         int group_count = get_num_groups(0) * get_num_groups(1) * get_num_groups(2);
  26.    
  27.         int groups_x = get_num_groups(0);
  28.         int groups_y = get_num_groups(1);
  29.         int groups_z = get_num_groups(2);
  30.    
  31.         int groupid_x = get_group_id(0);
  32.         int groupid_y = get_group_id(1);
  33.         int groupid_z = get_group_id(2);
  34.    
  35.         int group_id = groupid_x + groups_x*groupid_y + groups_y*groups_x*groupid_z;
  36.         bool isValid = true;
  37.     __local int count_val = 0;
  38.     count_val++;
  39.    
  40.     if(idx > sizex) isValid = false;
  41.     if(idy > sizey) isValid = false;
  42.     if(idz > sizez) isValid = false;
  43.  
  44.  
  45.         if(isValid == true)
  46.         {  
  47.             __local float Minvals[ WORKGROUPSIZE ];
  48.             __local float Maxvals[ WORKGROUPSIZE ];
  49.  
  50.             Minvals[gid_local] = input[gid];
  51.             Maxvals[gid_local] = input[gid];
  52.        
  53.         barrier(CLK_LOCAL_MEM_FENCE);
  54.             int offset = local_size/2;
  55.             //Lokal reduzieren
  56.             for(;gid_local<offset;offset/=2)
  57.             {
  58.        
  59.             Minvals[gid_local] = fmin(Minvals[gid_local],Minvals[gid_local+offset]);
  60.             Maxvals[gid_local] = fmax(Maxvals[gid_local],Maxvals[gid_local+offset]);
  61.        
  62.         }
  63.             barrier(CLK_LOCAL_MEM_FENCE);
  64.        
  65.  
  66.             //Global rueckschreiben
  67.             if(gid_local == 0)
  68.             {
  69.                 //global Rueckschreiben
  70.                 output[group_id*2] =  Minvals[gid_local];
  71.                 output[group_id*2 +1] = Maxvals[gid_local];
  72.            
  73.             }
  74.  
  75.             barrier(CLK_GLOBAL_MEM_FENCE);
  76.        
  77.             //Global Reduzieren
  78.             if(gid==0)
  79.             {
  80.                     float mint=256;
  81.                     float maxt=0;
  82.                 for(int i =0; i < group_count; i=i+2)
  83.                     {
  84.                         mint     = fmin(mint,output[i]);
  85.                         maxt     = fmax(maxt,output[i+1]);
  86.                
  87.                     }  
  88.            
  89.                     output[0] = mint;
  90.                     output[1] = maxt;
  91.             }
  92.        
  93.         }
  94. }
Advertisement
RAW Paste Data Copied
Advertisement