summaryrefslogtreecommitdiff
path: root/chroma/cuda/interpolate.h
blob: e7335babedefd1ff0ad7d6159ef3a88a73ac3073 (plain)
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
42
43
44
45
46
47
48
49
50
51
#ifndef __INTERPOLATE_H__
#define __INTERPOLATE_H__

__device__ float
interp(float x, int n, float *xp, float *fp)
{
    int lower = 0;
    int upper = n-1;

    if (x <= xp[lower])
	return fp[lower];

    if (x >= xp[upper])
	return fp[upper];

    while (lower < upper-1)
    {
	int half = (lower+upper)/2;

	if (x < xp[half])
	    upper = half;
	else
	    lower = half;
    }

    float df = fp[upper] - fp[lower];
    float dx = xp[upper] - xp[lower];

    return fp[lower] + df*(x-xp[lower])/dx;
}

__device__ float
interp_uniform(float x, int n, float x0, float dx, float *fp)
{
    if (x <= x0)
	return x0;

    float xmax = x0 + dx*(n-1);

    if (x >= xmax)
	return xmax;

    int lower = (x - x0)/dx;
    int upper = lower + 1;

    float df = fp[upper] - fp[lower];

    return fp[lower] + df*(x-(x0+dx*lower))/dx;
}

#endif