Advertisement
Not a member of Pastebin yet?
Sign Up,
it unlocks many cool features!
- __kernel void MinMax( __global float * input,
- __global float * output,
- int sizex,
- int sizey,
- int sizez,
- int size,
- int local_size
- )
- {
- int idx = get_global_id(0);
- int idy = get_global_id(1);
- int idz = get_global_id(2);
- int idx_local = get_local_id(0);
- int idy_local = get_local_id(1);
- int idz_local = get_local_id(2);
- int local_x = get_local_size(0);
- int local_y = get_local_size(1);
- int local_z = get_local_size(2);
- int gid = idx + sizex*idy + sizex*sizey*idz;
- int gid_local = idx_local + local_x*idy_local +local_x*local_y*idz_local;
- int group_count = get_num_groups(0) * get_num_groups(1) * get_num_groups(2);
- //int group_id = groupid_x + group_y*groupid_y + groupid_z*group_z*groupid_z;
- if((gid<size) && (gid_local<local_size))
- {
- __local float Minvals[ 64 ];
- __local float Maxvals[ 64 ];
- //local memory
- Minvals[gid_local] = input[gid];
- Maxvals[gid_local] = input[gid];
- int j = gid + gid_local + group_count;
- float ass = Minvals[gid_local] + Maxvals[gid_local];
- int offset = (local_x*local_y*local_z)/2;
- //Local reduction
- for(;gid_local<offset;offset/=2)
- {
- Minvals[gid_local] = fmin(Minvals[gid_local],Minvals[gid_local+offset]);
- Maxvals[gid_local] = fmax(Maxvals[gid_local],Maxvals[gid_local+offset]);
- }
- barrier(CLK_GLOBAL_MEM_FENCE);
- //Global rueckschreiben
- if(gid_local == 0)
- {
- //global Rueckschreiben
- output[gid*2] = Minvals[gid_local];
- output[gid*2 +1] = Maxvals[gid_local];
- }
- barrier(CLK_GLOBAL_MEM_FENCE);
- //Global Reduction
- if(gid==0)
- {
- float mint=256;
- float maxt=0;
- for(int i =0; i < group_count; i=i+2)
- {
- mint = fmin(mint,output[i]);
- maxt = fmax(maxt,output[i+1]);
- }
- output[0] = mint;
- output[1] = maxt;
- }
- }
- }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement