1 /*
2 * Copyright (C) by Argonne National Laboratory
3 * See COPYRIGHT in top-level directory
4 *
5 * DO NOT EDIT: AUTOMATICALLY GENERATED FILE !!
6 */
7
8 #include <string.h>
9 #include <stdint.h>
10 #include <wchar.h>
11 #include <assert.h>
12 #include <cuda.h>
13 #include <cuda_runtime.h>
14 #include "yaksuri_cudai_base.h"
15 #include "yaksuri_cudai_pup.h"
16
yaksuri_cudai_kernel_pack_blkhindx__Bool(const void * inbuf,void * outbuf,uintptr_t count,const yaksuri_cudai_md_s * __restrict__ md)17 __global__ void yaksuri_cudai_kernel_pack_blkhindx__Bool(const void *inbuf, void *outbuf, uintptr_t count, const yaksuri_cudai_md_s *__restrict__ md)
18 {
19 const char *__restrict__ sbuf = (const char *) inbuf;
20 char *__restrict__ dbuf = (char *) outbuf;
21 uintptr_t extent = md->extent;
22 uintptr_t idx = blockIdx.x * blockDim.x + threadIdx.x;
23 uintptr_t res = idx;
24 uintptr_t inner_elements = md->num_elements;
25
26 if (idx >= (count * inner_elements))
27 return;
28
29 uintptr_t x0 = res / inner_elements;
30 res %= inner_elements;
31 inner_elements /= md->u.blkhindx.count;
32
33 uintptr_t x1 = res / inner_elements;
34 res %= inner_elements;
35 inner_elements /= md->u.blkhindx.blocklength;
36 uintptr_t x2 = res;
37
38 intptr_t *array_of_displs1 = md->u.blkhindx.array_of_displs;
39 *((_Bool *) (void *) (dbuf + idx * sizeof(_Bool))) = *((const _Bool *) (const void *) (sbuf + x0 * extent + array_of_displs1[x1] + x2 * sizeof(_Bool)));
40 }
41
yaksuri_cudai_pack_blkhindx__Bool(const void * inbuf,void * outbuf,uintptr_t count,yaksuri_cudai_md_s * md,int n_threads,int n_blocks_x,int n_blocks_y,int n_blocks_z,int device)42 void yaksuri_cudai_pack_blkhindx__Bool(const void *inbuf, void *outbuf, uintptr_t count, yaksuri_cudai_md_s *md, int n_threads, int n_blocks_x, int n_blocks_y, int n_blocks_z, int device)
43 {
44 void *args[] = { &inbuf, &outbuf, &count, &md };
45 cudaError_t cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_pack_blkhindx__Bool,
46 dim3(n_blocks_x, n_blocks_y, n_blocks_z), dim3(n_threads), args, 0, yaksuri_cudai_global.stream[device]);
47 YAKSURI_CUDAI_CUDA_ERR_CHECK(cerr);
48 }
49
yaksuri_cudai_kernel_unpack_blkhindx__Bool(const void * inbuf,void * outbuf,uintptr_t count,const yaksuri_cudai_md_s * __restrict__ md)50 __global__ void yaksuri_cudai_kernel_unpack_blkhindx__Bool(const void *inbuf, void *outbuf, uintptr_t count, const yaksuri_cudai_md_s *__restrict__ md)
51 {
52 const char *__restrict__ sbuf = (const char *) inbuf;
53 char *__restrict__ dbuf = (char *) outbuf;
54 uintptr_t extent = md->extent;
55 uintptr_t idx = blockIdx.x * blockDim.x + threadIdx.x;
56 uintptr_t res = idx;
57 uintptr_t inner_elements = md->num_elements;
58
59 if (idx >= (count * inner_elements))
60 return;
61
62 uintptr_t x0 = res / inner_elements;
63 res %= inner_elements;
64 inner_elements /= md->u.blkhindx.count;
65
66 uintptr_t x1 = res / inner_elements;
67 res %= inner_elements;
68 inner_elements /= md->u.blkhindx.blocklength;
69 uintptr_t x2 = res;
70
71 intptr_t *array_of_displs1 = md->u.blkhindx.array_of_displs;
72 *((_Bool *) (void *) (dbuf + x0 * extent + array_of_displs1[x1] + x2 * sizeof(_Bool))) = *((const _Bool *) (const void *) (sbuf + idx * sizeof(_Bool)));
73 }
74
yaksuri_cudai_unpack_blkhindx__Bool(const void * inbuf,void * outbuf,uintptr_t count,yaksuri_cudai_md_s * md,int n_threads,int n_blocks_x,int n_blocks_y,int n_blocks_z,int device)75 void yaksuri_cudai_unpack_blkhindx__Bool(const void *inbuf, void *outbuf, uintptr_t count, yaksuri_cudai_md_s *md, int n_threads, int n_blocks_x, int n_blocks_y, int n_blocks_z, int device)
76 {
77 void *args[] = { &inbuf, &outbuf, &count, &md };
78 cudaError_t cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_unpack_blkhindx__Bool,
79 dim3(n_blocks_x, n_blocks_y, n_blocks_z), dim3(n_threads), args, 0, yaksuri_cudai_global.stream[device]);
80 YAKSURI_CUDAI_CUDA_ERR_CHECK(cerr);
81 }
82
83