/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/host_utility/kernel_launch.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/host_utility/kernel_launch.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/host_utility/kernel_launch.hpp Source File
kernel_launch.hpp
Go to the documentation of this file.
1 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
2 // SPDX-License-Identifier: MIT
3 
4 #pragma once
5 #ifndef __HIPCC_RTC__
6 #include <hip/hip_runtime.h>
7 
8 #include "ck/ck.hpp"
9 #include "ck/utility/env.hpp"
10 #include "ck/stream_config.hpp"
11 #include "ck/host_utility/hip_check_error.hpp"
12 
13 namespace ck {
14 
15 template <typename... Args, typename F>
16 float launch_and_time_kernel(const StreamConfig& stream_config,
17  F kernel,
18  dim3 grid_dim,
19  dim3 block_dim,
20  std::size_t lds_byte,
21  Args... args)
22 {
23 #if CK_TIME_KERNEL
24  if(stream_config.time_kernel_)
25  {
26  if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))
27  {
28  printf("%s: grid_dim {%u, %u, %u}, block_dim {%u, %u, %u} \n",
29  __func__,
30  grid_dim.x,
31  grid_dim.y,
32  grid_dim.z,
33  block_dim.x,
34  block_dim.y,
35  block_dim.z);
36 
37  printf("Warm up %d times\n", stream_config.cold_niters_);
38  }
39  // warm up
40  for(int i = 0; i < stream_config.cold_niters_; ++i)
41  {
42  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
43  hip_check_error(hipGetLastError());
44  }
45 
46  const int nrepeat = stream_config.nrepeat_;
47  if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))
48  {
49  printf("Start running %d times...\n", nrepeat);
50  }
51  hipEvent_t start, stop;
52 
53  hip_check_error(hipEventCreate(&start));
54  hip_check_error(hipEventCreate(&stop));
55 
56  hip_check_error(hipDeviceSynchronize());
57  hip_check_error(hipEventRecord(start, stream_config.stream_id_));
58 
59  for(int i = 0; i < nrepeat; ++i)
60  {
61  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
62  hip_check_error(hipGetLastError());
63  }
64 
65  hip_check_error(hipEventRecord(stop, stream_config.stream_id_));
66  hip_check_error(hipEventSynchronize(stop));
67 
68  float total_time = 0;
69 
70  hip_check_error(hipEventElapsedTime(&total_time, start, stop));
71 
72  hip_check_error(hipEventDestroy(start));
73  hip_check_error(hipEventDestroy(stop));
74 
75  return total_time / nrepeat;
76  }
77  else
78  {
79  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
80  hip_check_error(hipGetLastError());
81 
82  return 0;
83  }
84 #else
85  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
86  hip_check_error(hipGetLastError());
87 
88  return 0;
89 #endif
90 }
91 
92 template <typename... Args, typename F, typename PreProcessFunc>
93 float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
94  PreProcessFunc preprocess,
95  F kernel,
96  dim3 grid_dim,
97  dim3 block_dim,
98  std::size_t lds_byte,
99  Args... args)
100 {
101 #if CK_TIME_KERNEL
102  if(stream_config.time_kernel_)
103  {
104  if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))
105  {
106  printf("%s: grid_dim {%u, %u, %u}, block_dim {%u, %u, %u} \n",
107  __func__,
108  grid_dim.x,
109  grid_dim.y,
110  grid_dim.z,
111  block_dim.x,
112  block_dim.y,
113  block_dim.z);
114 
115  printf("Warm up %d times\n", stream_config.cold_niters_);
116  }
117  // warm up
118  preprocess();
119  for(int i = 0; i < stream_config.cold_niters_; ++i)
120  {
121  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
122  hip_check_error(hipGetLastError());
123  }
124 
125  const int nrepeat = stream_config.nrepeat_;
126  if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))
127  {
128  printf("Start running %d times...\n", nrepeat);
129  }
130  hipEvent_t start, stop;
131 
132  hip_check_error(hipEventCreate(&start));
133  hip_check_error(hipEventCreate(&stop));
134 
135  hip_check_error(hipDeviceSynchronize());
136  hip_check_error(hipEventRecord(start, stream_config.stream_id_));
137 
138  for(int i = 0; i < nrepeat; ++i)
139  {
140  preprocess();
141  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
142  hip_check_error(hipGetLastError());
143  }
144 
145  hip_check_error(hipEventRecord(stop, stream_config.stream_id_));
146  hip_check_error(hipEventSynchronize(stop));
147 
148  float total_time = 0;
149 
150  hip_check_error(hipEventElapsedTime(&total_time, start, stop));
151 
152  hip_check_error(hipEventDestroy(start));
153  hip_check_error(hipEventDestroy(stop));
154 
155  return total_time / nrepeat;
156  }
157  else
158  {
159  preprocess();
160  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
161  hip_check_error(hipGetLastError());
162 
163  return 0;
164  }
165 #else
166  kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
167  hip_check_error(hipGetLastError());
168 
169  return 0;
170 #endif
171 }
172 
173 } // namespace ck
174 
175 #endif
stream_config.hpp
hip_check_error.hpp
ck
Definition: ck.hpp:270
ck::launch_and_time_kernel_with_preprocess
float launch_and_time_kernel_with_preprocess(const StreamConfig &stream_config, PreProcessFunc preprocess, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
Definition: kernel_launch.hpp:93
ck::hip_check_error
void hip_check_error(hipError_t x)
Definition: hip_check_error.hpp:12
ck::EnvIsEnabled
bool EnvIsEnabled(EnvVar)
Definition: env.hpp:139
ck::launch_and_time_kernel
float launch_and_time_kernel(const StreamConfig &stream_config, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
Definition: kernel_launch.hpp:16
StreamConfig
Definition: stream_config.hpp:10
StreamConfig::cold_niters_
int cold_niters_
Definition: stream_config.hpp:14
StreamConfig::time_kernel_
bool time_kernel_
Definition: stream_config.hpp:12
StreamConfig::nrepeat_
int nrepeat_
Definition: stream_config.hpp:15
StreamConfig::stream_id_
hipStream_t stream_id_
Definition: stream_config.hpp:11
CK_ENV
#define CK_ENV(name)
Definition: env.hpp:128