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