Skip to content

Commit

Permalink
fix bitonic kernel
Browse files Browse the repository at this point in the history
  • Loading branch information
AvvALlV committed Oct 21, 2023
1 parent 2ef082f commit 473d0d9
Showing 1 changed file with 3 additions and 7 deletions.
10 changes: 3 additions & 7 deletions src/cl/bitonic.cl
Original file line number Diff line number Diff line change
Expand Up @@ -2,19 +2,15 @@
#include <libgpu/opencl/cl/clion_defines.cl>
#endif

void swap(float* a, float* b) {
float tmp = *a;
*a = *b;
*b = tmp;
}

__kernel void bitonic(__global float *as, int K, int local_k) {
int gid = get_global_id(0);
int r_id = gid / local_k;
int rl_id = gid - r_id * local_k;
int idx = 2 * local_k * r_id + rl_id;
bool is_left = (gid / K) % 2 == 0 ? as[idx + local_k] < as[idx] : as[idx + local_k] > as[idx];
if (is_left) {
swap(&as[idx], &as[idx + local_k]);
float tmp = as[idx];
as[idx] = as[idx + local_k];
as[idx + local_k] = tmp;
}
}

0 comments on commit 473d0d9

Please sign in to comment.