Neko  0.8.1
A portable framework for high-order spectral element flow simulations
gs.c
Go to the documentation of this file.
1 /*
2  Copyright (c) 2021, The Neko Authors
3  All rights reserved.
4 
5  Redistribution and use in source and binary forms, with or without
6  modification, are permitted provided that the following conditions
7  are met:
8 
9  * Redistributions of source code must retain the above copyright
10  notice, this list of conditions and the following disclaimer.
11 
12  * Redistributions in binary form must reproduce the above
13  copyright notice, this list of conditions and the following
14  disclaimer in the documentation and/or other materials provided
15  with the distribution.
16 
17  * Neither the name of the authors nor the names of its
18  contributors may be used to endorse or promote products derived
19  from this software without specific prior written permission.
20 
21  THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22  "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23  LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24  FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25  COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26  INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27  BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28  LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29  CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30  LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31  ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32  POSSIBILITY OF SUCH DAMAGE.
33 */
34 
35 #ifdef __APPLE__
36 #include <OpenCL/cl.h>
37 #else
38 #include <CL/cl.h>
39 #endif
40 
41 #include <limits.h>
42 #include <stdio.h>
43 #include <device/device_config.h>
44 #include <device/opencl/jit.h>
45 #include <device/opencl/prgm_lib.h>
46 #include <device/opencl/check.h>
47 
48 #include "gs_kernels.cl.h"
49 
50 #define GS_OP_ADD 1
51 #define GS_OP_MUL 2
52 #define GS_OP_MIN 3
53 #define GS_OP_MAX 4
54 
58 void opencl_gather_kernel(void *v, int *m, int *o, void *dg,
59  void *u, int *n, void *gd, int *nb,
60  void *b, void *bo, int *op) {
61  cl_int err;
62 
63  if (gs_program == NULL)
64  opencl_kernel_jit(gs_kernels, (cl_program *) &gs_program);
65 
66  const int nblks = ((*m) + 256 - 1) / 256;
67  const size_t global_item_size = 256 * nblks;
68  const size_t local_item_size = 256;
69 
70 
71  switch (*op) {
72  case GS_OP_ADD:
73  {
74  cl_kernel kernel = clCreateKernel(gs_program,
75  "gather_kernel_add", &err);
76  CL_CHECK(err);
77 
78  CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &v));
79  CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), m));
80  CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), o));
81  CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dg));
82  CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &u));
83  CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
84  CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &gd));
85  CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), nb));
86  CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &b));
87  CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &bo));
88 
89  CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel,
90  1, NULL, &global_item_size,
91  &local_item_size, 0, NULL, NULL));
92  }
93  break;
94  case GS_OP_MUL:
95  {
96  cl_kernel kernel = clCreateKernel(gs_program,
97  "gather_kernel_mul", &err);
98  CL_CHECK(err);
99 
100  CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &v));
101  CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), m));
102  CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), o));
103  CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dg));
104  CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &u));
105  CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
106  CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &gd));
107  CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), nb));
108  CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &b));
109  CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &bo));
110 
111  CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel,
112  1, NULL, &global_item_size,
113  &local_item_size, 0, NULL, NULL));
114  }
115  break;
116  case GS_OP_MIN:
117  {
118  cl_kernel kernel = clCreateKernel(gs_program,
119  "gather_kernel_min", &err);
120  CL_CHECK(err);
121 
122  CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &v));
123  CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), m));
124  CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), o));
125  CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dg));
126  CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &u));
127  CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
128  CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &gd));
129  CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), nb));
130  CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &b));
131  CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &bo));
132 
133  CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel,
134  1, NULL, &global_item_size,
135  &local_item_size, 0, NULL, NULL));
136  }
137  break;
138  case GS_OP_MAX:
139  {
140  cl_kernel kernel = clCreateKernel(gs_program,
141  "gather_kernel_max", &err);
142  CL_CHECK(err);
143 
144  CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &v));
145  CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), m));
146  CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), o));
147  CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dg));
148  CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &u));
149  CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
150  CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &gd));
151  CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), nb));
152  CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &b));
153  CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &bo));
154 
155  CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel,
156  1, NULL, &global_item_size,
157  &local_item_size, 0, NULL, NULL));
158  }
159  break;
160  }
161 }
162 
166 void opencl_scatter_kernel(void *v, int *m, void *dg,
167  void *u, int *n, void *gd,
168  int *nb, void *b, void *bo) {
169  cl_int err;
170 
171  if (gs_program == NULL)
172  opencl_kernel_jit(gs_kernels, (cl_program *) &gs_program);
173 
174  cl_kernel kernel = clCreateKernel(gs_program, "scatter_kernel", &err);
175  CL_CHECK(err);
176 
177  CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &v));
178  CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), m));
179  CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dg));
180  CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &u));
181  CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
182  CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &gd));
183  CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), nb));
184  CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &b));
185  CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &bo));
186 
187  const int nblks = ((*m) + 256 - 1) / 256;
188  const size_t global_item_size = 256 * nblks;
189  const size_t local_item_size = 256;
190 
191  CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
192  NULL, &global_item_size, &local_item_size,
193  0, NULL, NULL));
194 }
u
__global__ void const T *__restrict__ u
Definition: conv1_kernel.h:132
v
__global__ void const T *__restrict__ const T *__restrict__ v
Definition: lambda2_kernel.h:243
device_config.h
glb_cmd_queue
void * glb_cmd_queue
GS_OP_MAX
#define GS_OP_MAX
Definition: gs.c:53
GS_OP_ADD
#define GS_OP_ADD
Definition: gs.c:50
opencl_gather_kernel
void opencl_gather_kernel(void *v, int *m, int *o, void *dg, void *u, int *n, void *gd, int *nb, void *b, void *bo, int *op)
Definition: gs.c:58
opencl_scatter_kernel
void opencl_scatter_kernel(void *v, int *m, void *dg, void *u, int *n, void *gd, int *nb, void *b, void *bo)
Definition: gs.c:166
GS_OP_MUL
#define GS_OP_MUL
Definition: gs.c:51
GS_OP_MIN
#define GS_OP_MIN
Definition: gs.c:52
opencl_kernel_jit
void opencl_kernel_jit(const char *kernel, cl_program *program)
Definition: jit.c:50
CL_CHECK
#define CL_CHECK(err)
Definition: check.h:12
prgm_lib.h
gs_program
void * gs_program