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_hvector_wchar_t(const void * inbuf,void * outbuf,uintptr_t count,const yaksuri_cudai_md_s * __restrict__ md)17 __global__ void yaksuri_cudai_kernel_pack_hvector_wchar_t(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.hvector.count;
32
33 uintptr_t x1 = res / inner_elements;
34 res %= inner_elements;
35 inner_elements /= md->u.hvector.blocklength;
36 uintptr_t x2 = res;
37
38 intptr_t stride1 = md->u.hvector.stride;
39 *((wchar_t *) (void *) (dbuf + idx * sizeof(wchar_t))) = *((const wchar_t *) (const void *) (sbuf + x0 * extent + x1 * stride1 + x2 * sizeof(wchar_t)));
40 }
41
yaksuri_cudai_pack_hvector_wchar_t(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_hvector_wchar_t(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_hvector_wchar_t,
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_hvector_wchar_t(const void * inbuf,void * outbuf,uintptr_t count,const yaksuri_cudai_md_s * __restrict__ md)50 __global__ void yaksuri_cudai_kernel_unpack_hvector_wchar_t(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.hvector.count;
65
66 uintptr_t x1 = res / inner_elements;
67 res %= inner_elements;
68 inner_elements /= md->u.hvector.blocklength;
69 uintptr_t x2 = res;
70
71 intptr_t stride1 = md->u.hvector.stride;
72 *((wchar_t *) (void *) (dbuf + x0 * extent + x1 * stride1 + x2 * sizeof(wchar_t))) = *((const wchar_t *) (const void *) (sbuf + idx * sizeof(wchar_t)));
73 }
74
yaksuri_cudai_unpack_hvector_wchar_t(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_hvector_wchar_t(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_hvector_wchar_t,
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