steepest_descent_kernel.h

Go to the documentation of this file.

 
#ifndef __NEKO_TOP_CUDA_STEEPEST_DESCENT_KERNELS__
#define __NEKO_TOP_CUDA_STEEPEST_DESCENT_KERNELS__
 
template <typename T>
__global__ void sd_direction_kernel(
    T* __restrict__ s, const T* __restrict__ x, const T xmin, const T xmax,
    const int n) {
 
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int str = blockDim.x * gridDim.x;
 
    for (int i = idx; i < n; i += str) {
        const T xi = x[i];
        T si = -s[i];
        if ((xi >= xmax && si > T(0.0)) || (xi <= xmin && si < T(0.0))) {
            si = T(0.0);
        }
        s[i] = si;
    }
}
 
template <typename T>
__global__ void sd_update_move_limits_kernel(
    T* __restrict__ delta, const T* __restrict__ x, const T* __restrict__ x1,
    const T* __restrict__ x2, const T dec, const T inc, const T dmin,
    const T dmax, const int n) {
 
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int str = blockDim.x * gridDim.x;
 
    for (int i = idx; i < n; i += str) {
        const T d = (x[i] - x1[i]) * (x1[i] - x2[i]);
        T di = delta[i];
        if (d < T(0.0)) {
            di = dec * di;
        } else if (d > T(0.0)) {
            di = inc * di;
        }
        delta[i] = min(dmax, max(dmin, di));
    }
}
 
template <typename T>
__global__ void sd_update_design_kernel(
    T* __restrict__ x_new, const T* __restrict__ x, const T* __restrict__ s,
    const T* __restrict__ delta, const T psi, const T smax, const T xmin,
    const T xmax, const int n) {
 
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int str = blockDim.x * gridDim.x;
 
    for (int i = idx; i < n; i += str) {
        const T xi = x[i];
        const T di = delta[i];
        const T trial = xi + psi * (s[i] / smax);
        x_new[i] = min(xmax, max(xmin, min(xi + di, max(xi - di, trial))));
    }
}
 
#endif // __NEKO_TOP_CUDA_STEEPEST_DESCENT_KERNELS__