Skip to content

Instantly share code, notes, and snippets.

@M0nteCarl0
Created August 3, 2023 19:01
Show Gist options
  • Save M0nteCarl0/7753f858fcf4924133bc07851b417b56 to your computer and use it in GitHub Desktop.
Save M0nteCarl0/7753f858fcf4924133bc07851b417b56 to your computer and use it in GitHub Desktop.
CUDA C linear & cubic interpolation kernels
__global__ void cubic_interp(float* x, float* y, int n, float* x_new, float* y_new, int m) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= m) return;
int i = 1;
while (i < n && x[i] <= x_new[idx]) {
i++;
}
i--;
float h = x[i + 1] - x[i];
float a = (x[i + 1] - x_new[idx]) / h;
float b = (x_new[idx] - x[i]) / h;
float c = ((y[i + 1] - y[i]) / h - (2 * (y[i + 1] - y[i])) / (h * h) * h * h) / h;
float d = ((2 * (y[i + 1] - y[i])) / (h * h) * h * h - (y[i + 1] - y[i]) / h) / h;
y_new[idx] = y[i] + a * (y[i + 1] - y[i]) + b * ((1.0 / 3.0) * c * h * h * h + (1.0 / 2.0) * d * h * h);
}
__global__ void linear_interp(float* x, float* y, int n, float* x_new, float* y_new, int m) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= m) return;
int i = 0;
while (i < n - 1 && x_new[idx] > x[i + 1]) {
i++;
}
float slope = (y[i + 1] - y[i]) / (x[i + 1] - x[i]);
y_new[idx] = y[i] + slope * (x_new[idx] - x[i]);
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment