math_ext.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>
#include <neko/math/bcknd/device/device_mpi_op.h>
#include <neko/math/bcknd/device/device_mpi_reduce.h>
 
// Local includes
#include "math_ext_kernel.h"
 
extern "C" {
 
void hip_copy_mask(void* a, void* b, int* size, int* mask, int* mask_size) {
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
 
    if(*mask_size == 0) return;
    hipLaunchKernelGGL(copy_mask_kernel<real>, nblcks, nthrds, 0,
         (hipStream_t)glb_cmd_queue,
         (real*)a, (real*)b, *size, mask, *mask_size);
    HIP_CHECK(hipGetLastError());
}
 
void hip_cadd_mask(void* a, real* c, int* size, int* mask, int* mask_size) {
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
 
    if(*mask_size == 0) return;
    hipLaunchKernelGGL(cadd_mask_kernel<real>, nblcks, nthrds, 0,
         (hipStream_t)glb_cmd_queue, (real*)a, *c, *size, mask, *mask_size);
    HIP_CHECK(hipGetLastError());
}
 
void hip_invcol1_mask(void* a, int* size, int* mask, int* mask_size) {
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
 
    if(*mask_size == 0) return;
    hipLaunchKernelGGL(invcol1_mask_kernel<real>, nblcks, nthrds, 0,
         (hipStream_t)glb_cmd_queue, (real*)a, *size, mask, *mask_size);
    HIP_CHECK(hipGetLastError());
}
 
void hip_col2_mask(void* a, void* b, int* size, int* mask, int* mask_size) {
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
 
    if(*mask_size == 0) return;
    hipLaunchKernelGGL(col2_mask_kernel<real>, nblcks, nthrds, 0,
         (hipStream_t)glb_cmd_queue,
         (real*)a, (real*)b, *size, mask, *mask_size);
    HIP_CHECK(hipGetLastError());
}
 
void hip_col3_mask(
    void* a, void* b, void* c, int* size, int* mask, int* mask_size) {
 
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
 
    if(*mask_size == 0) return;
    hipLaunchKernelGGL(col3_mask_kernel<real>, nblcks, nthrds, 0,
         (hipStream_t)glb_cmd_queue,
         (real*)a, (real*)b, (real*)c, *size, mask, *mask_size);
    HIP_CHECK(hipGetLastError());
}
 
void hip_sub3_mask(
    void* a, void* b, void* c, int* size, int* mask, int* mask_size) {
 
    const dim3 nthrds(1024, 1, 1);
    const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
 
    if(*mask_size == 0) return;
    hipLaunchKernelGGL(sub3_mask_kernel<real>, nblcks, nthrds, 0,
         (hipStream_t)glb_cmd_queue,
         (real*)a, (real*)b, (real*)c, *size, mask, *mask_size);
    HIP_CHECK(hipGetLastError());
}
}