-
Notifications
You must be signed in to change notification settings - Fork 548
Expand file tree
/
Copy pathdebug_cuda.hpp
More file actions
74 lines (63 loc) · 2.91 KB
/
debug_cuda.hpp
File metadata and controls
74 lines (63 loc) · 2.91 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
/*******************************************************
* Copyright (c) 2014, ArrayFire
* All rights reserved.
*
* This file is distributed under 3-clause BSD license.
* The complete license agreement can be obtained at:
* http://arrayfire.com/licenses/BSD-3-Clause
********************************************************/
#pragma once
#include <common/Logger.hpp>
#include <err_cuda.hpp>
#include <platform.hpp>
#include <string>
namespace arrayfire {
namespace cuda {
namespace kernel_logger {
inline auto getLogger() {
static auto logger = common::loggerFactory("kernel");
return logger;
}
} // namespace kernel_logger
} // namespace cuda
} // namespace arrayfire
template<>
struct fmt::formatter<dim3> : fmt::formatter<std::string> {
// parse is inherited from formatter<string_view>.
template<typename FormatContext>
auto format(dim3 c, FormatContext& ctx) {
std::string name = fmt::format("{} {} {}", c.x, c.y, c.z);
return formatter<std::string>::format(name, ctx);
}
};
#define CUDA_LAUNCH_SMEM(fn, blks, thrds, smem_size, ...) \
do { \
{ \
using namespace arrayfire::cuda::kernel_logger; \
AF_TRACE( \
"Launching {}: Blocks: [{}] Threads: [{}] " \
"Shared Memory: {}", \
#fn, blks, thrds, smem_size); \
} \
fn<<<blks, thrds, smem_size, arrayfire::cuda::getActiveStream()>>>( \
__VA_ARGS__); \
} while (false)
#define CUDA_LAUNCH(fn, blks, thrds, ...) \
CUDA_LAUNCH_SMEM(fn, blks, thrds, 0, __VA_ARGS__)
// FIXME: Add a special flag for debug
#ifndef NDEBUG
#define POST_LAUNCH_CHECK() \
do { \
CUDA_CHECK(cudaStreamSynchronize(arrayfire::cuda::getActiveStream())); \
} while (0)
#else
#define POST_LAUNCH_CHECK() \
do { \
if (arrayfire::cuda::synchronize_calls()) { \
CUDA_CHECK( \
cudaStreamSynchronize(arrayfire::cuda::getActiveStream())); \
} else { \
CUDA_CHECK(cudaPeekAtLastError()); \
} \
} while (0)
#endif