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_contig__Bool(const void * inbuf,void * outbuf,uintptr_t count,const yaksuri_cudai_md_s * __restrict__ md)17 __global__ void yaksuri_cudai_kernel_pack_contig__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.contig.count;
32
33 uintptr_t x1 = res;
34
35 intptr_t stride1 = md->u.contig.child->extent;
36 *((_Bool *) (void *) (dbuf + idx * sizeof(_Bool))) = *((const _Bool *) (const void *) (sbuf + x0 * extent + x1 * stride1));
37 }
38
yaksuri_cudai_pack_contig__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)39 void yaksuri_cudai_pack_contig__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)
40 {
41 void *args[] = { &inbuf, &outbuf, &count, &md };
42 cudaError_t cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_pack_contig__Bool,
43 dim3(n_blocks_x, n_blocks_y, n_blocks_z), dim3(n_threads), args, 0, yaksuri_cudai_global.stream[device]);
44 YAKSURI_CUDAI_CUDA_ERR_CHECK(cerr);
45 }
46
yaksuri_cudai_kernel_unpack_contig__Bool(const void * inbuf,void * outbuf,uintptr_t count,const yaksuri_cudai_md_s * __restrict__ md)47 __global__ void yaksuri_cudai_kernel_unpack_contig__Bool(const void *inbuf, void *outbuf, uintptr_t count, const yaksuri_cudai_md_s *__restrict__ md)
48 {
49 const char *__restrict__ sbuf = (const char *) inbuf;
50 char *__restrict__ dbuf = (char *) outbuf;
51 uintptr_t extent = md->extent;
52 uintptr_t idx = blockIdx.x * blockDim.x + threadIdx.x;
53 uintptr_t res = idx;
54 uintptr_t inner_elements = md->num_elements;
55
56 if (idx >= (count * inner_elements))
57 return;
58
59 uintptr_t x0 = res / inner_elements;
60 res %= inner_elements;
61 inner_elements /= md->u.contig.count;
62
63 uintptr_t x1 = res;
64
65 intptr_t stride1 = md->u.contig.child->extent;
66 *((_Bool *) (void *) (dbuf + x0 * extent + x1 * stride1)) = *((const _Bool *) (const void *) (sbuf + idx * sizeof(_Bool)));
67 }
68
yaksuri_cudai_unpack_contig__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)69 void yaksuri_cudai_unpack_contig__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)
70 {
71 void *args[] = { &inbuf, &outbuf, &count, &md };
72 cudaError_t cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_unpack_contig__Bool,
73 dim3(n_blocks_x, n_blocks_y, n_blocks_z), dim3(n_threads), args, 0, yaksuri_cudai_global.stream[device]);
74 YAKSURI_CUDAI_CUDA_ERR_CHECK(cerr);
75 }
76
77