math_ext.cu

Go to the documentation of this file.

 
// System includes
#include <stdio.h>
#include <stdlib.h>
 
// Device includes
#include <cuda_runtime.h>
 
// Neko includes
#include <neko/device/device_config.h>
#include <neko/device/cuda/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 cuda_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;
    cadd_mask_kernel<real><<<nblcks, nthrds, 0, (cudaStream_t)glb_cmd_queue>>>(
        (real*)a, *c, *size, mask, *mask_size);
    CUDA_CHECK(cudaGetLastError());
}
 
void cuda_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;
    invcol1_mask_kernel<real>
        <<<nblcks, nthrds, 0, (cudaStream_t)glb_cmd_queue>>>(
            (real*)a, *size, mask, *mask_size);
    CUDA_CHECK(cudaGetLastError());
}
 
void cuda_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;
    col2_mask_kernel<real><<<nblcks, nthrds, 0, (cudaStream_t)glb_cmd_queue>>>(
        (real*)a, (real*)b, *size, mask, *mask_size);
    CUDA_CHECK(cudaGetLastError());
}
 
void cuda_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;
    col3_mask_kernel<real><<<nblcks, nthrds, 0, (cudaStream_t)glb_cmd_queue>>>(
        (real*)a, (real*)b, (real*)c, *size, mask, *mask_size);
    CUDA_CHECK(cudaGetLastError());
}
 
void cuda_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;
    sub3_mask_kernel<real><<<nblcks, nthrds, 0, (cudaStream_t)glb_cmd_queue>>>(
        (real*)a, (real*)b, (real*)c, *size, mask, *mask_size);
    CUDA_CHECK(cudaGetLastError());
}
}