1#ifndef __MATH_MATH_KERNEL_H__
2#define __MATH_MATH_KERNEL_H__
48 for (
int i = idx;
i < n;
i +=
str) {
104 for (
int i = idx;
i < m;
i +=
str) {
105#if __CUDA_ARCH__ >= 600
114template<
typename T >
137 const int mask_size) {
142 for (
int i = idx;
i < mask_size;
i +=
str) {
a[
mask[
i]] = c; }
148template<
typename T >
157 for (
int i = idx;
i < n;
i +=
str) {
165template<
typename T >
173 for (
int i = idx;
i < n;
i +=
str) {
181template<
typename T >
190 for (
int i = idx;
i < n;
i +=
str) {
198template<
typename T >
206 for (
int i = idx;
i < n;
i +=
str) {
214template<
typename T >
223 for (
int i = idx;
i < n;
i +=
str) {
231template<
typename T >
239 for (
int i = idx;
i < n;
i +=
str) {
247template<
typename T >
255 for (
int i = idx;
i < n;
i +=
str) {
263template<
typename T >
272 for (
int i = idx;
i < n;
i +=
str) {
280template<
typename T >
290 for (
int i = idx;
i < n;
i +=
str) {
291 a[
i] = b[
i] + c[
i] + d[
i];
298template<
typename T >
307 for (
int i = idx;
i < n;
i +=
str) {
315template<
typename T >
326 for (
int i = idx;
i < n;
i+=
str) {
329 tmp += p[
j][
i]*alpha[
j];
338template<
typename T >
347 for (
int i = idx;
i < n;
i +=
str) {
355template<
typename T >
364 for (
int i = idx;
i < n;
i +=
str) {
372template<
typename T >
383 for (
int i = idx;
i < n;
i +=
str) {
391template<
typename T >
399 for (
int i = idx;
i < n;
i +=
str) {
407template<
typename T >
415 for (
int i = idx;
i < n;
i +=
str) {
423template<
typename T >
432 for (
int i = idx;
i < n;
i +=
str) {
440template<
typename T >
448 for (
int i = idx;
i < n;
i +=
str) {
456template<
typename T >
465 for (
int i = idx;
i < n;
i +=
str) {
473template<
typename T >
482 for (
int i = idx;
i < n;
i +=
str) {
483 a[
i] =
a[
i] - b[
i] * c[
i];
490template<
typename T >
498 for (
int i = idx;
i < n;
i +=
str) {
506template<
typename T >
515 for (
int i = idx;
i < n;
i +=
str) {
523template<
typename T >
532 for (
int i = idx;
i < n;
i +=
str) {
533 a[
i] =
a[
i] + b[
i] * c[
i];
541template<
typename T >
551 for (
int i = idx;
i < n;
i +=
str) {
552 a[
i] =
a[
i] + b[
i] * c[
i] * d[
i];
560template<
typename T >
573 for (
int i = idx;
i < n;
i +=
str) {
582template<
typename T >
597 for (
int i = idx;
i < n;
i +=
str) {
600 u3[
i] = v1[
i]*w2[
i] - v2[
i]*w1[
i];
622template<
typename T >
628 for (
int i = idx;
i<n ;
i +=
str)
655template<
typename T >
666 for (
int i=idx ;
i<n ;
i+=step)
690template<
typename T >
705 for (
int i = idx;
i < n;
i+=
str) {
725template<
typename T >
740 for (
int i = idx;
i < n;
i+=
str) {
766template<
typename T >
780 for (
int i = idx;
i < n;
i+=
str) {
801template<
typename T >
815 for (
int i = idx;
i < n;
i+=
str) {
836template<
typename T >
849 for (
int i = idx;
i<n ;
i +=
str)
871template<
typename T >
878 for (
int i = idx;
i < n;
i +=
str) {
__global__ void const T *__restrict__ x
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
__global__ void addcol4_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T *__restrict__ d, const int n)
__global__ void pwmin_vec3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void reduce_kernel(T *bufred, const int n)
__global__ void cdiv2_kernel(T *__restrict__ a, T *__restrict__ b, const T c, const int n)
__global__ void invcol2_kernel(T *__restrict__ a, const T *__restrict__ b, const int n)
__global__ void add2_kernel(T *__restrict__ a, const T *__restrict__ b, const int n)
__inline__ __device__ T reduce_warp(T val)
__global__ void masked_atomic_reduction_kernel(T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int m)
__global__ void pwmax_vec3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void glsc3_many_kernel(const T *a, const T **b, const T *c, T *buf_h, const int j, const int n)
__global__ void cfill_mask_kernel(T *__restrict__ a, const T c, const int size, int *__restrict__ mask, const int mask_size)
__global__ void cdiv_kernel(T *__restrict__ a, const T c, const int n)
__global__ void glsc3_reduce_kernel(T *bufred, const int n, const int j)
__global__ void pwmax_sca2_kernel(T *__restrict__ a, const T c, const int n)
__global__ void glsubnorm2_kernel(const T *a, const T *b, T *buf_h, const int n)
__global__ void pwmin_vec2_kernel(T *__restrict__ a, const T *__restrict__ b, const int n)
__global__ void add3s2_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T c1, const T c2, const int n)
__global__ void add2s1_kernel(T *__restrict__ a, const T *__restrict__ b, const T c1, const int n)
__global__ void masked_gather_copy_kernel(T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
__global__ void add2s2_many_kernel(T *__restrict__ x, const T **p, const T *alpha, const int p_cur, const int n)
__global__ void pwmax_vec2_kernel(T *__restrict__ a, const T *__restrict__ b, const int n)
__global__ void cmult_kernel(T *__restrict__ a, const T c, const int n)
__global__ void addcol3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void pwmin_sca3_kernel(T *__restrict__ a, const T *__restrict b, const T c, const int n)
__global__ void pwmax_sca3_kernel(T *__restrict__ a, const T *__restrict b, const T c, const int n)
__global__ void col2_kernel(T *__restrict__ a, const T *__restrict__ b, const int n)
__global__ void col3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void sub2_kernel(T *__restrict__ a, const T *__restrict__ b, const int n)
__global__ void glsc2_kernel(const T *a, const T *b, T *buf_h, const int n)
__global__ void cmult2_kernel(T *__restrict__ a, T *__restrict__ b, const T c, const int n)
__global__ void pwmin_sca2_kernel(T *__restrict__ a, const T c, const int n)
__global__ void sub3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void glsum_kernel(const T *a, T *buf_h, const int n)
__global__ void masked_copy_kernel(T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
__global__ void glsc3_kernel(const T *a, const T *b, const T *c, T *buf_h, const int n)
__global__ void add2s2_kernel(T *__restrict__ a, const T *__restrict__ b, const T c1, const int n)
__global__ void vdot3_kernel(T *__restrict__ dot, const T *__restrict__ u1, const T *__restrict__ u2, const T *__restrict__ u3, const T *__restrict__ v1, const T *__restrict__ v2, const T *__restrict__ v3, const int n)
__global__ void invcol1_kernel(T *__restrict__ a, const int n)
__global__ void add3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void add4_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const T *__restrict__ d, const int n)
__global__ void cfill_kernel(T *__restrict__ a, const T c, const int n)
__global__ void masked_scatter_copy_kernel(T *__restrict__ a, T *__restrict__ b, int *__restrict__ mask, const int n, const int n_mask)
__global__ void vcross_kernel(T *__restrict__ u1, T *__restrict__ u2, T *__restrict__ u3, const T *__restrict__ v1, const T *__restrict__ v2, const T *__restrict__ v3, const T *__restrict__ w1, const T *__restrict__ w2, const T *__restrict__ w3, const int n)
__global__ void addsqr2s2_kernel(T *__restrict__ a, const T *__restrict__ b, const T c1, const int n)
__global__ void cadd2_kernel(T *__restrict__ a, T *__restrict__ b, const T c, const int n)
__global__ void absval_kernel(T *__restrict__ a, const int n)
__global__ void invcol3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void subcol3_kernel(T *__restrict__ a, const T *__restrict__ b, const T *__restrict__ c, const int n)
__global__ void cadd_kernel(T *__restrict__ a, const T c, const int n)
Object for handling masks in Neko.