RAMP_mapping.hip

Go to the documentation of this file.

 
// System includes
#include <stdio.h>
#include <stdlib.h>
 
// Device includes
#include <hip/hip_runtime.h>
 
// Neko includes
#include <neko/device/device_config.h>
#include <neko/device/hip/check.h>
 
// Local includes
#include "RAMP_mapping_kernel.h"
 
extern "C" {
 
void hip_convex_down_RAMP_mapping_apply(real* f_min, real* f_max, real* q,
    void* X_out_d, void* X_in_d, int* n) {
 
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
 
    hipLaunchKernelGGL(convex_down_RAMP_mapping_apply_kernel<real>,
        nblcks, nthrds, 0, (hipStream_t)glb_cmd_queue,
        *f_min, *f_max, *q, (real*)X_out_d, (real*)X_in_d, *n);
    HIP_CHECK(hipGetLastError());
}
 
void hip_convex_down_RAMP_mapping_apply_backward(real* f_min, real* f_max,
    real* q, void* sens_out_d, void* sens_in_d, void* X_in_d, int* n) {
 
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
 
    hipLaunchKernelGGL(convex_down_RAMP_mapping_apply_backward_kernel<real>,
        nblcks, nthrds, 0, (hipStream_t)glb_cmd_queue,
        *f_min, *f_max, *q, (real*)sens_out_d, (real*)sens_in_d,
        (real*)X_in_d, *n);
    HIP_CHECK(hipGetLastError());
}
 
void hip_convex_up_RAMP_mapping_apply(real* f_min, real* f_max, real* q,
    void* X_out_d, void* X_in_d, int* n) {
 
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
 
    hipLaunchKernelGGL(convex_up_RAMP_mapping_apply_kernel<real>,
        nblcks, nthrds, 0, (hipStream_t)glb_cmd_queue,
        *f_min, *f_max, *q, (real*)X_out_d, (real*)X_in_d, *n);
    HIP_CHECK(hipGetLastError());
}
 
void hip_convex_up_RAMP_mapping_apply_backward(real* f_min, real* f_max,
    real* q, void* sens_out_d, void* sens_in_d, void* X_in_d, int* n) {
 
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
 
    hipLaunchKernelGGL(convex_up_RAMP_mapping_apply_backward_kernel<real>,
        nblcks, nthrds, 0, (hipStream_t)glb_cmd_queue,
        *f_min, *f_max, *q, (real*)sens_out_d, (real*)sens_in_d,
        (real*)X_in_d, *n);
    HIP_CHECK(hipGetLastError());
}
}