210_check_list_performance.doxy 20 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443
  1. /* StarPU --- Runtime system for heterogeneous multicore architectures.
  2. *
  3. * Copyright (C) 2011-2013,2015,2017 Inria
  4. * Copyright (C) 2010-2017 CNRS
  5. * Copyright (C) 2009-2011,2013-2017 Université de Bordeaux
  6. *
  7. * StarPU is free software; you can redistribute it and/or modify
  8. * it under the terms of the GNU Lesser General Public License as published by
  9. * the Free Software Foundation; either version 2.1 of the License, or (at
  10. * your option) any later version.
  11. *
  12. * StarPU is distributed in the hope that it will be useful, but
  13. * WITHOUT ANY WARRANTY; without even the implied warranty of
  14. * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
  15. *
  16. * See the GNU Lesser General Public License in COPYING.LGPL for more details.
  17. */
  18. /*! \page CheckListWhenPerformanceAreNotThere Check List When Performance Are Not There
  19. TODO: improve!
  20. To achieve good
  21. performance, we give below a list of features which should be checked.
  22. For a start, you can use \ref OfflinePerformanceTools to get a Gantt chart which
  23. will show roughly where time is spent, and focus correspondingly.
  24. \section ConfigurationImprovePerformance Configuration That May Improve Performance
  25. The \ref enable-fast "--enable-fast" configuration option disables all
  26. assertions. This makes StarPU more performant for really small tasks by
  27. disabling all sanity checks. Only use this for measurements and production, not for development, since this will drop all basic checks.
  28. \section DataRelatedFeaturesToImprovePerformance Data Related Features That May Improve Performance
  29. link to \ref DataManagement
  30. link to \ref DataPrefetch
  31. \section TaskRelatedFeaturesToImprovePerformance Task Related Features That May Improve Performance
  32. link to \ref TaskGranularity
  33. link to \ref TaskSubmission
  34. link to \ref TaskPriorities
  35. \section SchedulingRelatedFeaturesToImprovePerformance Scheduling Related Features That May Improve Performance
  36. link to \ref TaskSchedulingPolicy
  37. link to \ref TaskDistributionVsDataTransfer
  38. link to \ref Energy-basedScheduling
  39. link to \ref StaticScheduling
  40. \section CUDA-specificOptimizations CUDA-specific Optimizations
  41. Due to CUDA limitations, StarPU will have a hard time overlapping its own
  42. communications and the codelet computations if the application does not use a
  43. dedicated CUDA stream for its computations instead of the default stream,
  44. which synchronizes all operations of the GPU. StarPU provides one by the use
  45. of starpu_cuda_get_local_stream() which can be used by all CUDA codelet
  46. operations to avoid this issue. For instance:
  47. \code{.c}
  48. func <<<grid,block,0,starpu_cuda_get_local_stream()>>> (foo, bar);
  49. cudaStreamSynchronize(starpu_cuda_get_local_stream());
  50. \endcode
  51. Unfortunately, some CUDA libraries do not have stream variants of
  52. kernels. That will lower the potential for overlapping.
  53. Calling starpu_cublas_init() makes StarPU already do appropriate calls for the
  54. CUBLAS library. Some libraries like Magma may however change the current stream of CUBLAS v1,
  55. one then has to call <c>cublasSetKernelStream(starpu_cuda_get_local_stream())</c> at
  56. the beginning of the codelet to make sure that CUBLAS is really using the proper
  57. stream. When using CUBLAS v2, starpu_cublas_get_local_handle() can be called to queue CUBLAS
  58. kernels with the proper configuration.
  59. Similarly, calling starpu_cusparse_init() makes StarPU create CUSPARSE handles
  60. on each CUDA device, starpu_cusparse_get_local_handle() can then be used to
  61. queue CUSPARSE kernels with the proper configuration.
  62. If the kernel can be made to only use this local stream or other self-allocated
  63. streams, i.e. the whole kernel submission can be made asynchronous, then
  64. one should enable asynchronous execution of the kernel. That means setting
  65. the flag ::STARPU_CUDA_ASYNC in the corresponding field starpu_codelet::cuda_flags, and dropping the
  66. <c>cudaStreamSynchronize()</c> call at the end of the <c>cuda_func</c> function, so that it
  67. returns immediately after having queued the kernel to the local stream. That way, StarPU will be
  68. able to submit and complete data transfers while kernels are executing, instead of only at each
  69. kernel submission. The kernel just has to make sure that StarPU can use the
  70. local stream to synchronize with the kernel startup and completion.
  71. If the kernel uses its own non-default stream, one can synchronize that stream
  72. with the StarPU-provided stream this way:
  73. \code{.c}
  74. cudaEvent_t event;
  75. call_kernel_with_its_own_stream()
  76. cudaEventCreateWithFlags(&event, cudaEventDisableTiming);
  77. cudaEventRecord(event, get_kernel_stream());
  78. cudaStreamWaitEvent(starpu_cuda_get_local_stream(), event, 0);
  79. cudaEventDestroy(event);
  80. \endcode
  81. That code makes the StarPU-provided stream wait for a new event, which will be
  82. triggered by the completion of the kernel.
  83. Using the flag ::STARPU_CUDA_ASYNC also permits to enable concurrent kernel
  84. execution, on cards which support it (Kepler and later, notably). This is
  85. enabled by setting the environment variable \ref STARPU_NWORKER_PER_CUDA to the
  86. number of kernels to execute concurrently. This is useful when kernels are
  87. small and do not feed the whole GPU with threads to run.
  88. \section OpenCL-specificOptimizations OpenCL-specific Optimizations
  89. If the kernel can be made to only use the StarPU-provided command queue or other self-allocated
  90. queues, i.e. the whole kernel submission can be made asynchronous, then
  91. one should enable asynchronous execution of the kernel. This means setting
  92. the flag ::STARPU_OPENCL_ASYNC in the corresponding field starpu_codelet::opencl_flags and dropping the
  93. <c>clFinish()</c> and starpu_opencl_collect_stats() calls at the end of the kernel, so
  94. that it returns immediately after having queued the kernel to the provided queue.
  95. That way, StarPU will be able to submit and complete data transfers while kernels are executing, instead of
  96. only at each kernel submission. The kernel just has to make sure
  97. that StarPU can use the command queue it has provided to synchronize with the
  98. kernel startup and completion.
  99. \section DetectionStuckConditions Detecting Stuck Conditions
  100. It may happen that for some reason, StarPU does not make progress for a long
  101. period of time. Reason are sometimes due to contention inside StarPU, but
  102. sometimes this is due to external reasons, such as stuck MPI driver, or CUDA
  103. driver, etc.
  104. <c>export STARPU_WATCHDOG_TIMEOUT=10000</c> (\ref STARPU_WATCHDOG_TIMEOUT)
  105. allows to make StarPU print an error message whenever StarPU does not terminate
  106. any task for 10ms, but lets the application continue normally. In addition to that,
  107. <c>export STARPU_WATCHDOG_CRASH=1</c> (\ref STARPU_WATCHDOG_CRASH)
  108. raises <c>SIGABRT</c> in that condition, thus allowing to catch the situation in gdb.
  109. It can also be useful to type <c>handle SIGABRT nopass</c> in <c>gdb</c> to be able to let
  110. the process continue, after inspecting the state of the process.
  111. \section HowToLimitMemoryPerNode How to Limit Memory Used By StarPU And Cache Buffer Allocations
  112. By default, StarPU makes sure to use at most 90% of the memory of GPU devices,
  113. moving data in and out of the device as appropriate and with prefetch and
  114. writeback optimizations. Concerning the main memory, by default it will not
  115. limit its consumption, since by default it has nowhere to push the data to when
  116. memory gets tight. This also means that by default StarPU will not cache buffer
  117. allocations in main memory, since it does not know how much of the system memory
  118. it can afford.
  119. In the case of GPUs, the \ref STARPU_LIMIT_CUDA_MEM, \ref STARPU_LIMIT_CUDA_devid_MEM,
  120. \ref STARPU_LIMIT_OPENCL_MEM, and \ref STARPU_LIMIT_OPENCL_devid_MEM environment variables
  121. can be used to control how
  122. much (in MiB) of the GPU device memory should be used at most by StarPU (their
  123. default values are 90% of the available memory).
  124. In the case of the main memory, the \ref STARPU_LIMIT_CPU_MEM environment
  125. variable can be used to specify how much (in MiB) of the main memory should be
  126. used at most by StarPU for buffer allocations. This way, StarPU will be able to
  127. cache buffer allocations (which can be a real benefit if a lot of bufferes are
  128. involved, or if allocation fragmentation can become a problem), and when using
  129. \ref OutOfCore, StarPU will know when it should evict data out to the disk.
  130. It should be noted that by default only buffer allocations automatically
  131. done by StarPU are accounted here, i.e. allocations performed through
  132. starpu_malloc_on_node() which are used by the data interfaces
  133. (matrix, vector, etc.). This does not include allocations performed by
  134. the application through e.g. malloc(). It does not include allocations
  135. performed through starpu_malloc() either, only allocations
  136. performed explicitly with the \ref STARPU_MALLOC_COUNT flag, i.e. by calling
  137. \code{.c}
  138. starpu_malloc_flags(STARPU_MALLOC_COUNT)
  139. \endcode
  140. are taken into account. If the
  141. application wants to make StarPU aware of its own allocations, so that StarPU
  142. knows precisely how much data is allocated, and thus when to evict allocation
  143. caches or data out to the disk, starpu_memory_allocate() can be used to
  144. specify an amount of memory to be accounted for. starpu_memory_deallocate()
  145. can be used to account freed memory back. Those can for instance be used by data
  146. interfaces with dynamic data buffers: instead of using starpu_malloc_on_node(),
  147. they would dynamically allocate data with malloc/realloc, and notify starpu of
  148. the delta thanks to starpu_memory_allocate() and starpu_memory_deallocate() calls.
  149. starpu_memory_get_total() and starpu_memory_get_available()
  150. can be used to get an estimation of how much memory is available.
  151. starpu_memory_wait_available() can also be used to block until an
  152. amount of memory becomes available, but it may be preferrable to call
  153. \code{.c}
  154. starpu_memory_allocate(STARPU_MEMORY_WAIT)
  155. \endcode
  156. to reserve that amount immediately.
  157. \section HowToReduceTheMemoryFootprintOfInternalDataStructures How To Reduce The Memory Footprint Of Internal Data Structures
  158. It is possible to reduce the memory footprint of the task and data internal
  159. structures of StarPU by describing the shape of your machine and/or your
  160. application at the configure step.
  161. To reduce the memory footprint of the data internal structures of StarPU, one
  162. can set the
  163. \ref enable-maxcpus "--enable-maxcpus",
  164. \ref enable-maxnumanodes "--enable-maxnumanodes",
  165. \ref enable-maxcudadev "--enable-maxcudadev",
  166. \ref enable-maxopencldev "--enable-maxopencldev" and
  167. \ref enable-maxnodes "--enable-maxnodes"
  168. configure parameters to give StarPU
  169. the architecture of the machine it will run on, thus tuning the size of the
  170. structures to the machine.
  171. To reduce the memory footprint of the task internal structures of StarPU, one
  172. can set the \ref enable-maxbuffers "--enable-maxbuffers" configure parameter to
  173. give StarPU the maximum number of buffers that a task can use during an
  174. execution. For example, in the Cholesky factorization (dense linear algebra
  175. application), the GEMM task uses up to 3 buffers, so it is possible to set the
  176. maximum number of task buffers to 3 to run a Cholesky factorization on StarPU.
  177. The size of the various structures of StarPU can be printed by
  178. <c>tests/microbenchs/display_structures_size</c>.
  179. It is also often useless to submit *all* the tasks at the same time. One can
  180. make the starpu_task_submit() function block when a reasonable given number of
  181. tasks have been submitted, by setting the \ref STARPU_LIMIT_MIN_SUBMITTED_TASKS and
  182. \ref STARPU_LIMIT_MAX_SUBMITTED_TASKS environment variables, for instance:
  183. <c>
  184. export STARPU_LIMIT_MAX_SUBMITTED_TASKS=10000
  185. export STARPU_LIMIT_MIN_SUBMITTED_TASKS=9000
  186. </c>
  187. To make StarPU block submission when 10000 tasks are submitted, and unblock
  188. submission when only 9000 tasks are still submitted, i.e. 1000 tasks have
  189. completed among the 10000 that were submitted when submission was blocked. Of
  190. course this may reduce parallelism if the threshold is set too low. The precise
  191. balance depends on the application task graph.
  192. An idea of how much memory is used for tasks and data handles can be obtained by
  193. setting the \ref STARPU_MAX_MEMORY_USE environment variable to <c>1</c>.
  194. \section HowtoReuseMemory How To Reuse Memory
  195. When your application needs to allocate more data than the available amount of
  196. memory usable by StarPU (given by starpu_memory_get_available()), the
  197. allocation cache system can reuse data buffers used by previously executed
  198. tasks. For that system to work with MPI tasks, you need to submit tasks progressively instead
  199. of as soon as possible, because in the case of MPI receives, the allocation cache check for reusing data
  200. buffers will be done at submission time, not at execution time.
  201. You have two options to control the task submission flow. The first one is by
  202. controlling the number of submitted tasks during the whole execution. This can
  203. be done whether by setting the environment variables
  204. \ref STARPU_LIMIT_MAX_SUBMITTED_TASKS and \ref STARPU_LIMIT_MIN_SUBMITTED_TASKS to
  205. tell StarPU when to stop submitting tasks and when to wake up and submit tasks
  206. again, or by explicitely calling starpu_task_wait_for_n_submitted() in
  207. your application code for finest grain control (for example, between two
  208. iterations of a submission loop).
  209. The second option is to control the memory size of the allocation cache. This
  210. can be done in the application by using jointly
  211. starpu_memory_get_available() and starpu_memory_wait_available() to submit
  212. tasks only when there is enough memory space to allocate the data needed by the
  213. task, i.e when enough data are available for reuse in the allocation cache.
  214. \section PerformanceModelCalibration Performance Model Calibration
  215. Most schedulers are based on an estimation of codelet duration on each kind
  216. of processing unit. For this to be possible, the application programmer needs
  217. to configure a performance model for the codelets of the application (see
  218. \ref PerformanceModelExample for instance). History-based performance models
  219. use on-line calibration. StarPU will automatically calibrate codelets
  220. which have never been calibrated yet, and save the result in
  221. <c>$STARPU_HOME/.starpu/sampling/codelets</c>.
  222. The models are indexed by machine name.
  223. By default, StarPU stores separate performance models according to the hostname
  224. of the system. To avoid having to calibrate performance models for each node
  225. of a homogeneous cluster for instance, the model can be shared by using
  226. <c>export STARPU_HOSTNAME=some_global_name</c> (\ref STARPU_HOSTNAME), where
  227. <c>some_global_name</c> is the name of the cluster for instance, which thus
  228. overrides the hostname of the system.
  229. By default, StarPU stores separate performance models for each GPU. To avoid
  230. having to calibrate performance models for each GPU of a homogeneous set of GPU
  231. devices for instance, the model can be shared by setting
  232. <c>export STARPU_PERF_MODEL_HOMOGENEOUS_CUDA=1</c> ,
  233. <c>export STARPU_PERF_MODEL_HOMOGENEOUS_OPENCL=1</c> ,
  234. <c>export STARPU_PERF_MODEL_HOMOGENEOUS_MIC=1</c> ,
  235. <c>export STARPU_PERF_MODEL_HOMOGENEOUS_MPI_MS=1</c> , or
  236. <c>export STARPU_PERF_MODEL_HOMOGENEOUS_SCC=1</c> (depending on your GPU device type).
  237. To force continuing calibration,
  238. use <c>export STARPU_CALIBRATE=1</c> (\ref STARPU_CALIBRATE). This may be necessary if your application
  239. has not-so-stable performance. StarPU will force calibration (and thus ignore
  240. the current result) until 10 (<c>_STARPU_CALIBRATION_MINIMUM</c>) measurements have been
  241. made on each architecture, to avoid badly scheduling tasks just because the
  242. first measurements were not so good. Details on the current performance model status
  243. can be obtained from the tool <c>starpu_perfmodel_display</c>: the <c>-l</c>
  244. option lists the available performance models, and the <c>-s</c> option permits
  245. to choose the performance model to be displayed. The result looks like:
  246. \verbatim
  247. $ starpu_perfmodel_display -s starpu_slu_lu_model_11
  248. performance model for cpu_impl_0
  249. # hash size flops mean dev n
  250. 914f3bef 1048576 0.000000e+00 2.503577e+04 1.982465e+02 8
  251. 3e921964 65536 0.000000e+00 5.527003e+02 1.848114e+01 7
  252. e5a07e31 4096 0.000000e+00 1.717457e+01 5.190038e+00 14
  253. ...
  254. \endverbatim
  255. Which shows that for the LU 11 kernel with a 1MiB matrix, the average
  256. execution time on CPUs was about 25ms, with a 0.2ms standard deviation, over
  257. 8 samples. It is a good idea to check this before doing actual performance
  258. measurements.
  259. A graph can be drawn by using the tool <c>starpu_perfmodel_plot</c>:
  260. \verbatim
  261. $ starpu_perfmodel_plot -s starpu_slu_lu_model_11
  262. 4096 16384 65536 262144 1048576 4194304
  263. $ gnuplot starpu_starpu_slu_lu_model_11.gp
  264. $ gv starpu_starpu_slu_lu_model_11.eps
  265. \endverbatim
  266. \image html starpu_starpu_slu_lu_model_11.png
  267. \image latex starpu_starpu_slu_lu_model_11.eps "" width=\textwidth
  268. If a kernel source code was modified (e.g. performance improvement), the
  269. calibration information is stale and should be dropped, to re-calibrate from
  270. start. This can be done by using <c>export STARPU_CALIBRATE=2</c> (\ref STARPU_CALIBRATE).
  271. Note: history-based performance models get calibrated
  272. only if a performance-model-based scheduler is chosen.
  273. The history-based performance models can also be explicitly filled by the
  274. application without execution, if e.g. the application already has a series of
  275. measurements. This can be done by using starpu_perfmodel_update_history(),
  276. for instance:
  277. \code{.c}
  278. static struct starpu_perfmodel perf_model =
  279. {
  280. .type = STARPU_HISTORY_BASED,
  281. .symbol = "my_perfmodel",
  282. };
  283. struct starpu_codelet cl =
  284. {
  285. .cuda_funcs = { cuda_func1, cuda_func2 },
  286. .nbuffers = 1,
  287. .modes = {STARPU_W},
  288. .model = &perf_model
  289. };
  290. void feed(void)
  291. {
  292. struct my_measure *measure;
  293. struct starpu_task task;
  294. starpu_task_init(&task);
  295. task.cl = &cl;
  296. for (measure = &measures[0]; measure < measures[last]; measure++)
  297. {
  298. starpu_data_handle_t handle;
  299. starpu_vector_data_register(&handle, -1, 0, measure->size, sizeof(float));
  300. task.handles[0] = handle;
  301. starpu_perfmodel_update_history(&perf_model, &task,
  302. STARPU_CUDA_DEFAULT + measure->cudadev, 0,
  303. measure->implementation, measure->time);
  304. starpu_task_clean(&task);
  305. starpu_data_unregister(handle);
  306. }
  307. }
  308. \endcode
  309. Measurement has to be provided in milliseconds for the completion time models,
  310. and in Joules for the energy consumption models.
  311. \section Profiling Profiling
  312. A quick view of how many tasks each worker has executed can be obtained by setting
  313. <c>export STARPU_WORKER_STATS=1</c> (\ref STARPU_WORKER_STATS). This is a convenient way to check that
  314. execution did happen on accelerators, without penalizing performance with
  315. the profiling overhead.
  316. A quick view of how much data transfers have been issued can be obtained by setting
  317. <c>export STARPU_BUS_STATS=1</c> (\ref STARPU_BUS_STATS).
  318. More detailed profiling information can be enabled by using <c>export STARPU_PROFILING=1</c> (\ref STARPU_PROFILING)
  319. or by
  320. calling starpu_profiling_status_set() from the source code.
  321. Statistics on the execution can then be obtained by using <c>export
  322. STARPU_BUS_STATS=1</c> and <c>export STARPU_WORKER_STATS=1</c> .
  323. More details on performance feedback are provided in the next chapter.
  324. \section OverheadProfiling Overhead Profiling
  325. \ref OfflinePerformanceTools can already provide an idea of to what extent and
  326. which part of StarPU bring overhead on the execution time. To get a more precise
  327. analysis of the parts of StarPU which bring most overhead, <c>gprof</c> can be used.
  328. First, recompile and reinstall StarPU with <c>gprof</c> support:
  329. \code
  330. ./configure --enable-perf-debug --disable-shared --disable-build-tests --disable-build-examples
  331. \endcode
  332. Make sure not to leave a dynamic version of StarPU in the target path: remove
  333. any remaining <c>libstarpu-*.so</c>
  334. Then relink your application with the static StarPU library, make sure that
  335. running <c>ldd</c> on your application does not mention any libstarpu
  336. (i.e. it's really statically-linked).
  337. \code
  338. gcc test.c -o test $(pkg-config --cflags starpu-1.3) $(pkg-config --libs starpu-1.3)
  339. \endcode
  340. Now you can run your application, and a <c>gmon.out</c> file should appear in the
  341. current directory, you can process it by running <c>gprof</c> on your application:
  342. \code
  343. gprof ./test
  344. \endcode
  345. That will dump an analysis of the time spent in StarPU functions.
  346. */