-
Notifications
You must be signed in to change notification settings - Fork 0
/
convert.cl
41 lines (34 loc) · 1.11 KB
/
convert.cl
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
__kernel void convert_2bit_4chan(__global const unsigned char *input,
__global float2 *data,__local unsigned int *scratch,
__const float4 lut, __const int spc)
{
int idx = get_global_id(0);
int local_idx = get_local_id(0);
// Load the time sample into local memory
scratch[local_idx] = input[idx];
// Interpet the LUT as an array
float *lp = (float *)&lut;
// Loop over each channel
for (int c = 0; c < 4; c++)
{
data[c*spc + idx].x = lp[(scratch[local_idx] >> (2*c)) & 0x03];
data[c*spc + idx].y = 0;
}
}
__kernel void convert_2bit_8chan(__global const unsigned short *input,
__global float2 *data, __local unsigned int *scratch,
__const float4 lut, __const int spc)
{
int idx = get_global_id(0);
int local_idx = get_local_id(0);
// Load the time sample into local memory
scratch[local_idx] = input[idx];
// Interpet the LUT as an array
float *lp = (float *)&lut;
// Loop over each channel
for (int c = 0; c < 8; c++)
{
data[c*spc + idx].x = lp[(scratch[local_idx] >> (2*c)) & 0x03];
data[c*spc + idx].y = 0;
}
}