Don't like ads? PRO users don't see any ads ;-)
Public Pastes
Guest

minmax

By: a guest on Feb 10th, 2013  |  syntax: C++  |  size: 2.63 KB  |  hits: 26  |  expires: Never
download  |  raw  |  embed  |  report abuse  |  print
Text below is selected. Please press Ctrl+C to copy to your clipboard. (⌘+C on Mac)
  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. }
create a new version of this paste RAW Paste Data