Neko  0.9.99
A portable framework for high-order spectral element flow simulations
opr_opgrad.hip
Go to the documentation of this file.
1 /*
2  Copyright (c) 2021-2022, 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 #include <string.h>
36 #include <stdlib.h>
37 #include <stdio.h>
38 #include <hip/hip_runtime.h>
39 #include <device/device_config.h>
40 #include <device/hip/check.h>
41 #include "opgrad_kernel.h"
42 
43 extern "C" {
44  #include <common/neko_log.h>
45 }
46 
47 template < const int >
48 int tune_opgrad(void *ux, void *uy, void *uz, void *u,
49  void *dx, void *dy, void *dz,
50  void *drdx, void *dsdx, void *dtdx,
51  void *drdy, void *dsdy, void *dtdy,
52  void *drdz, void *dsdz, void *dtdz,
53  void *w3, int *nel, int *lx);
54 
55 extern "C" {
56 
60  void hip_opgrad(void *ux, void *uy, void *uz, void *u,
61  void *dx, void *dy, void *dz,
62  void *drdx, void *dsdx, void *dtdx,
63  void *drdy, void *dsdy, void *dtdy,
64  void *drdz, void *dsdz, void *dtdz,
65  void *w3, int *nel, int *lx) {
66 
67  static int autotune[19] = { 0 };
68 
69  const dim3 nthrds_1d(1024, 1, 1);
70  const dim3 nthrds_kstep((*lx), (*lx), 1);
71  const dim3 nblcks((*nel), 1, 1);
72 
73 #define CASE_1D(LX) \
74  hipLaunchKernelGGL( HIP_KERNEL_NAME(opgrad_kernel_1d<real, LX, 1024> ), \
75  nblcks, nthrds_1d, 0, (hipStream_t) glb_cmd_queue, \
76  (real *) ux, (real *) uy, (real *) uz, (real *) u, \
77  (real *) dx, (real *) dy, (real *) dz, \
78  (real *) drdx, (real *) dsdx, (real *) dtdx, \
79  (real *) drdy, (real *) dsdy, (real *) dtdy, \
80  (real *) drdz, (real *) dsdz, (real *) dtdz, \
81  (real *) w3); \
82  HIP_CHECK(hipGetLastError());
83 
84 
85 #define CASE_KSTEP(LX) \
86  hipLaunchKernelGGL( HIP_KERNEL_NAME(opgrad_kernel_kstep<real, LX> ), \
87  nblcks, nthrds_kstep, 0, (hipStream_t) glb_cmd_queue, \
88  (real *) ux, (real *) uy, (real *) uz, (real *) u, \
89  (real *) dx, (real *) dy, (real *) dz, \
90  (real *) drdx, (real *) dsdx, (real *) dtdx, \
91  (real *) drdy, (real *) dsdy, (real *) dtdy, \
92  (real *) drdz, (real *) dsdz, (real *) dtdz, \
93  (real *) w3); \
94  HIP_CHECK(hipGetLastError());
95 
96 #define CASE(LX) \
97  case LX: \
98  if(autotune[LX] == 0 ) { \
99  autotune[LX]=tune_opgrad<LX>(ux, uy, uz, u, \
100  dx, dy, dz, \
101  drdx, dsdx, dtdx, \
102  drdy, dsdy, dtdy, \
103  drdz, dsdz, dtdz, \
104  w3, nel, lx); \
105  } else if (autotune[LX] == 1 ) { \
106  CASE_1D(LX); \
107  } else if (autotune[LX] == 2 ) { \
108  CASE_KSTEP(LX); \
109  } \
110  break
111 
112  switch(*lx) {
113  CASE(2);
114  CASE(3);
115  CASE(4);
116  CASE(5);
117  CASE(6);
118  CASE(7);
119  CASE(8);
120  CASE(9);
121  CASE(10);
122  CASE(11);
123  CASE(12);
124  CASE(13);
125  CASE(14);
126  CASE(15);
127  CASE(16);
128  CASE(17);
129  CASE(18);
130  default:
131  {
132  fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
133  exit(1);
134  }
135  }
136  }
137 }
138 
139 template < const int LX >
140 int tune_opgrad(void *ux, void *uy, void *uz, void *u,
141  void *dx, void *dy, void *dz,
142  void *drdx, void *dsdx, void *dtdx,
143  void *drdy, void *dsdy, void *dtdy,
144  void *drdz, void *dsdz, void *dtdz,
145  void *w3, int *nel, int *lx) {
146  hipEvent_t start,stop;
147  float time1,time2;
148  int retval;
149 
150  const dim3 nthrds_1d(1024, 1, 1);
151  const dim3 nthrds_kstep((*lx), (*lx), 1);
152  const dim3 nblcks((*nel), 1, 1);
153 
154  char *env_value = NULL;
155  char neko_log_buf[80];
156 
157  env_value=getenv("NEKO_AUTOTUNE");
158 
159  sprintf(neko_log_buf, "Autotune opgrad (lx: %d)", *lx);
160  log_section(neko_log_buf);
161 
162  if(env_value) {
163  if( !strcmp(env_value,"1D") ) {
164  CASE_1D(LX);
165  sprintf(neko_log_buf,"Set by env : 1 (1D)");
166  log_message(neko_log_buf);
167  log_end_section();
168  return 1;
169  } else if( !strcmp(env_value,"KSTEP") ) {
170  CASE_KSTEP(LX);
171  sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
172  log_message(neko_log_buf);
173  log_end_section();
174  return 2;
175  } else {
176  sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
177  log_error(neko_log_buf);
178  }
179  }
180 
181  HIP_CHECK(hipEventCreate(&start));
182  HIP_CHECK(hipEventCreate(&stop));
183 
184  HIP_CHECK(hipEventRecord(start,0));
185 
186  for(int i = 0; i < 100; i++) {
187  CASE_1D(LX);
188  }
189 
190  HIP_CHECK(hipEventRecord(stop,0));
191  HIP_CHECK(hipEventSynchronize(stop));
192  HIP_CHECK(hipEventElapsedTime(&time1, start, stop));;
193 
194  HIP_CHECK(hipEventRecord(start,0));
195 
196  for(int i = 0; i < 100; i++) {
197  CASE_KSTEP(LX);
198  }
199 
200  HIP_CHECK(hipEventRecord(stop,0));
201  HIP_CHECK(hipEventSynchronize(stop));
202  HIP_CHECK(hipEventElapsedTime(&time2, start, stop));
203 
204  if(time1 < time2) {
205  retval = 1;
206  } else {
207  retval = 2;
208  }
209 
210  sprintf(neko_log_buf, "Chose : %d (%s)", retval,
211  (retval > 1 ? "KSTEP" : "1D"));
212  log_message(neko_log_buf);
213  log_end_section();
214  return retval;
215 }
drdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdy
Definition: ax_helm_full_kernel.h:54
drdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdz
Definition: ax_helm_full_kernel.h:55
dsdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdz
Definition: ax_helm_full_kernel.h:58
dsdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdy
Definition: ax_helm_full_kernel.h:57
i
const int i
Definition: ax_helm_full_kernel.h:95
dtdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdy
Definition: ax_helm_full_kernel.h:60
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition: ax_helm_full_kernel.h:46
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition: ax_helm_full_kernel.h:49
drdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdx
Definition: ax_helm_full_kernel.h:53
dtdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdz
Definition: ax_helm_full_kernel.h:61
dsdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdx
Definition: ax_helm_full_kernel.h:56
dtdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdx
Definition: ax_helm_full_kernel.h:59
dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition: ax_helm_full_kernel.h:51
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition: ax_helm_full_kernel.h:50
w3
__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
Definition: cdtp_kernel.h:113
uy
__global__ void T *__restrict__ uy
Definition: opgrad_kernel.h:126
uz
__global__ void T *__restrict__ T *__restrict__ uz
Definition: opgrad_kernel.h:127
device_config.h
HIP_CHECK
#define HIP_CHECK(err)
Definition: check.h:8
opgrad_kernel.h
neko_log.h
log_error
void log_error(char *msg)
log_message
void log_message(char *msg)
log_end_section
void log_end_section()
log_section
void log_section(char *msg)
CASE
#define CASE(LX)
tune_opgrad
int tune_opgrad(void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx)
Definition: opr_opgrad.hip:140
CASE_KSTEP
#define CASE_KSTEP(LX)
CASE_1D
#define CASE_1D(LX)
hip_opgrad
void hip_opgrad(void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx)
Definition: opr_opgrad.hip:60