GCC Code Coverage Report

chrono_batch_render/
Coverage:
low: ≥ 0%
medium: ≥ 75.0%
high: ≥ 90.0%
Lines:
0 / 214
0.0%
Branches:
0 / 32
0.0%
Functions:
0 / 53
0.0%
device/hip/device.cpp
Line Branch Exec Source
1 // SPDX-License-Identifier: BSD-3-Clause
2 #include <algorithm>
3 #include <cstddef>
4 #include <string>
5 #include <vector>
6
7 #include "chrono_batch_render/device/compute/compute.h"
8 #include "chrono_batch_render/device/hip/device.h"
9 #include "chrono_batch_render/device/queue.h"
10
11 extern "C" {
12 extern const unsigned char chrono_batch_render_hip_prelude[];
13 extern const std::size_t chrono_batch_render_hip_prelude_size;
14 }
15
16 namespace chrono::batch_render::device::hip {
17
18 ✗ DeviceScope::DeviceScope(const Api& hip, hipDevice_t device) : m_hip(hip) {
19 ✗ hip.Check(hip.hipGetDevice(&m_previous), "hipGetDevice");
20 ✗ hip.Check(hip.hipSetDevice(device), "hipSetDevice");
21 ✗ }
22
23 ✗ DeviceScope::~DeviceScope() {
24 ✗ m_hip.Warn(m_hip.hipSetDevice(m_previous), "hipSetDevice");
25 ✗ }
26
27 ✗ HipDevice::HipDevice(const AdapterInfo& adapter, const DeviceDesc& desc, Traversal traversal)
28
0/3
✗ Branch 2 not taken.
✗ Branch 6 not taken.
✗ Branch 9 not taken.
 ✗ : Device(adapter, desc, traversal), m_hip(GetApi()), m_device(int(adapter.ordinal)) {
29
0/1
✗ Branch 1 not taken.
 ✗ m_queue = std::make_unique<compute::ComputeQueue<Traits>>(*this);
30 ✗ }
31
32 ✗ Ref<Buffer> HipDevice::NewBuffer(const BufferDesc& desc) {
33 ✗ return compute::ComputeBuffer<Traits>::Create(*this, desc);
34 }
35
36 ✗ Ref<Texture> HipDevice::NewTexture(const TextureDesc& desc) {
37
0/1
✗ Branch 2 not taken.
 ✗ return Ref<Texture>(new compute::ComputeTexture<Traits>(*this, desc));
38 }
39
40 ✗ void HipDevice::ExternalCopy(BufferRange dst, uint64_t source, Stream producer) {
41 ✗ compute::ExternalCopy<Traits>(*this, dst, source, producer);
42 ✗ }
43
44 ✗ std::unique_ptr<Kernel> HipDevice::NewKernel() {
45 ✗ return std::make_unique<compute::ComputeKernel<Traits>>(*this);
46 }
47
48 ✗ std::vector<MemoryHeap> HipDevice::Heaps() const {
49 ✗ const DeviceScope scope = Traits::Enter(*this);
50 ✗ size_t free = 0;
51 ✗ size_t total = 0;
52
0/3
✗ Branch 1 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
 ✗ if (Functions().hipMemGetInfo(&free, &total) != 0) {
53 ✗ return {};
54 }
55
0/1
✗ Branch 1 not taken.
 ✗ return {{.usage = total - free, .budget = total, .size = total, .deviceLocal = true}};
56 ✗ }
57
58 ✗ DeviceScope Traits::Enter(const HipDevice& device) {
59 ✗ return DeviceScope(device.Functions(), device.Ordinal());
60 }
61
62 ✗ uint64_t Traits::Allocate(const Api& hip, uint64_t bytes) {
63 ✗ void* pointer = nullptr;
64 ✗ hip.Check(hip.hipMalloc(&pointer, bytes), "hipMalloc");
65 ✗ return reinterpret_cast<uint64_t>(pointer);
66 }
67
68 ✗ void Traits::Free(const Api& hip, uint64_t address) {
69 ✗ hip.Warn(hip.hipFree(reinterpret_cast<void*>(address)), "hipFree");
70 ✗ }
71
72 ✗ hipMemPool_t Traits::CreatePool(const HipDevice& device, uint64_t keptBytes) {
73 ✗ const Api& hip = device.Functions();
74 ✗ int pools = 0;
75 ✗ hip.Check(
76 ✗ hip.hipDeviceGetAttribute(&pools, hipDeviceAttributeMemoryPoolsSupported, device.Ordinal()),
77 "hipDeviceGetAttribute");
78
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (pools == 0) {
79 return nullptr;
80 }
81
82 ✗ hipMemPoolProps properties{};
83 ✗ properties.allocType = hipMemAllocationTypePinned;
84 ✗ properties.location = {hipMemLocationTypeDevice, device.Ordinal()};
85 ✗ hipMemPool_t pool = nullptr;
86 ✗ hip.Check(hip.hipMemPoolCreate(&pool, &properties), "hipMemPoolCreate");
87
88 ✗ const hipError_t set =
89 ✗ hip.hipMemPoolSetAttribute(pool, hipMemPoolAttrReleaseThreshold, &keptBytes);
90
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (set != hipSuccess) {
91 ✗ DestroyPool(hip, pool);
92 ✗ hip.Check(set, "hipMemPoolSetAttribute");
93 }
94 ✗ return pool;
95 }
96
97 ✗ void Traits::DestroyPool(const Api& hip, hipMemPool_t pool) {
98 ✗ hip.Warn(hip.hipMemPoolDestroy(pool), "hipMemPoolDestroy");
99 ✗ }
100
101 ✗ uint64_t Traits::AllocateFrom(const Api& hip, hipMemPool_t pool, uint64_t bytes,
102 hipStream_t stream) {
103 ✗ void* pointer = nullptr;
104 ✗ hip.Check(hip.hipMallocFromPoolAsync(&pointer, bytes, pool, stream), "hipMallocFromPoolAsync");
105 ✗ return reinterpret_cast<uint64_t>(pointer);
106 }
107
108 ✗ void Traits::FreeTo(const Api& hip, uint64_t address, hipStream_t stream) {
109 ✗ hip.Warn(hip.hipFreeAsync(reinterpret_cast<void*>(address), stream), "hipFreeAsync");
110 ✗ }
111
112 ✗ void Traits::SynchronizeStream(const Api& hip, hipStream_t stream) {
113 ✗ hip.Check(hip.hipStreamSynchronize(stream), "hipStreamSynchronize");
114 ✗ }
115
116 ✗ void* Traits::AllocateHost(const Api& hip, uint64_t bytes, bool writeCombined) {
117 ✗ void* host = nullptr;
118
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ const unsigned flags = hipHostMallocMapped | (writeCombined ? hipHostMallocWriteCombined : 0);
119 ✗ hip.Check(hip.hipHostMalloc(&host, bytes, flags), "hipHostMalloc");
120 ✗ return host;
121 }
122
123 ✗ uint64_t Traits::DeviceAddress(const Api& hip, void* host) {
124 ✗ void* pointer = nullptr;
125 ✗ hip.Check(hip.hipHostGetDevicePointer(&pointer, host, 0), "hipHostGetDevicePointer");
126 ✗ return reinterpret_cast<uint64_t>(pointer);
127 }
128
129 ✗ void Traits::FreeHost(const Api& hip, void* host) {
130 ✗ hip.Warn(hip.hipHostFree(host), "hipHostFree");
131 ✗ }
132
133 ✗ hipStream_t Traits::CreateStream(const Api& hip) {
134 ✗ hipStream_t stream = nullptr;
135 ✗ hip.Check(hip.hipStreamCreateWithFlags(&stream, hipStreamNonBlocking),
136 "hipStreamCreateWithFlags");
137 ✗ return stream;
138 }
139
140 ✗ void Traits::DestroyStream(const Api& hip, hipStream_t stream) {
141 ✗ hip.Warn(hip.hipStreamDestroy(stream), "hipStreamDestroy");
142 ✗ }
143
144 ✗ hipEvent_t Traits::CreateEvent(const Api& hip, bool timed) {
145 ✗ hipEvent_t event = nullptr;
146
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ hip.Check(hip.hipEventCreateWithFlags(&event, timed ? 0 : hipEventDisableTiming),
147 "hipEventCreateWithFlags");
148 ✗ return event;
149 }
150
151 ✗ void Traits::DestroyEvent(const Api& hip, hipEvent_t event) {
152 ✗ hip.Warn(hip.hipEventDestroy(event), "hipEventDestroy");
153 ✗ }
154
155 ✗ void Traits::Record(const Api& hip, hipEvent_t event, hipStream_t stream, bool node) {
156
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ hip.Check(hip.hipEventRecordWithFlags(event, stream, node ? hipEventRecordExternal : 0),
157 "hipEventRecordWithFlags");
158 ✗ }
159
160 ✗ void Traits::WaitEvent(const Api& hip, hipStream_t stream, hipEvent_t event) {
161 ✗ hip.Check(hip.hipStreamWaitEvent(stream, event, 0), "hipStreamWaitEvent");
162 ✗ }
163
164 ✗ hipStream_t Traits::ApplicationStream(void* handle) {
165 ✗ return static_cast<hipStream_t>(handle);
166 }
167
168 ✗ uint64_t Traits::ElapsedNs(const Api& hip, hipEvent_t from, hipEvent_t to) {
169 ✗ float milliseconds = 0;
170 ✗ hip.Check(hip.hipEventElapsedTime(&milliseconds, from, to), "hipEventElapsedTime");
171 ✗ return uint64_t(double(milliseconds) * 1e6);
172 }
173
174 ✗ bool Traits::Query(const Api& hip, hipEvent_t event) {
175 ✗ const hipError_t error = hip.hipEventQuery(event);
176
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (error == hipErrorNotReady) {
177 return false;
178 }
179 ✗ hip.Check(error, "hipEventQuery");
180 ✗ return true;
181 }
182
183 ✗ void Traits::Synchronize(const Api& hip, hipEvent_t event) {
184 ✗ hip.Check(hip.hipEventSynchronize(event), "hipEventSynchronize");
185 ✗ }
186
187 ✗ void Traits::Copy(const Api& hip, uint64_t dst, uint64_t src, uint64_t bytes, hipStream_t stream) {
188 ✗ hip.Check(hip.hipMemcpyAsync(reinterpret_cast<void*>(dst), reinterpret_cast<void*>(src), bytes,
189 hipMemcpyDefault, stream),
190 "hipMemcpyAsync");
191 ✗ }
192
193 ✗ void Traits::Fill(const Api& hip, uint64_t dst, uint32_t value, uint64_t bytes,
194 hipStream_t stream) {
195 ✗ hip.Check(hip.hipMemsetD32Async(reinterpret_cast<void*>(dst), int(value), bytes / 4, stream),
196 "hipMemsetD32Async");
197 ✗ }
198
199 ✗ void Traits::Upload(const Api& hip, uint64_t dst, const void* host, uint64_t bytes,
200 hipStream_t stream) {
201 ✗ hip.Check(hip.hipMemcpyHtoDAsync(reinterpret_cast<void*>(dst), host, bytes, stream),
202 "hipMemcpyHtoDAsync");
203 ✗ }
204
205 ✗ void* Traits::CreateTextureArray(const Api& hip, const ArrayDesc& desc, uint32_t levels) {
206 ✗ void* array = nullptr;
207 ✗ hip.Check(hip.hipMipmappedArrayCreate(&array, &desc, levels), "hipMipmappedArrayCreate");
208 ✗ return array;
209 }
210
211 ✗ void Traits::DestroyTextureArray(const Api& hip, void* array) {
212 ✗ hip.Warn(hip.hipMipmappedArrayDestroy(array), "hipMipmappedArrayDestroy");
213 ✗ }
214
215 ✗ void* Traits::TextureLevel(const Api& hip, void* array, uint32_t level) {
216 ✗ void* found = nullptr;
217 ✗ hip.Check(hip.hipMipmappedArrayGetLevel(&found, array, level), "hipMipmappedArrayGetLevel");
218 ✗ return found;
219 }
220
221 ✗ uint64_t Traits::CreateTextureObject(const Api& hip, const ResourceDesc& resource,
222 const SamplerDesc& sampler) {
223 ✗ void* object = nullptr;
224 ✗ hip.Check(hip.hipTexObjectCreate(&object, &resource, &sampler, nullptr), "hipTexObjectCreate");
225 ✗ return reinterpret_cast<uint64_t>(object);
226 }
227
228 ✗ void Traits::DestroyTextureObject(const Api& hip, uint64_t object) {
229 ✗ hip.Warn(hip.hipTexObjectDestroy(reinterpret_cast<void*>(object)), "hipTexObjectDestroy");
230 ✗ }
231
232 ✗ void Traits::CopyToArray(const Api& hip, void* level, uint64_t source, uint64_t rowBytes,
233 uint32_t rows, hipStream_t stream) {
234 ✗ const hip_Memcpy2D copy{.srcMemoryType = hipMemoryTypeDevice,
235 .srcDevice = reinterpret_cast<void*>(source),
236 .srcPitch = rowBytes,
237 .dstMemoryType = hipMemoryTypeArray,
238 .dstArray = level,
239 .WidthInBytes = rowBytes,
240 ✗ .Height = rows};
241 ✗ hip.Check(hip.hipMemcpyParam2DAsync(&copy, stream), "hipMemcpyParam2DAsync");
242 ✗ }
243
244 ✗ void Traits::BeginCapture(const Api& hip, hipStream_t stream) {
245 ✗ hip.Check(hip.hipStreamBeginCapture(stream, hipStreamCaptureModeThreadLocal),
246 "hipStreamBeginCapture");
247 ✗ }
248
249 ✗ hipGraphExec_t Traits::EndCapture(const Api& hip, hipStream_t stream) {
250 ✗ hipGraph_t graph = nullptr;
251 ✗ hip.Check(hip.hipStreamEndCapture(stream, &graph), "hipStreamEndCapture");
252 ✗ hipGraphExec_t exec = nullptr;
253 ✗ const auto instantiated = hip.hipGraphInstantiateWithFlags(&exec, graph, 0);
254 ✗ hip.Warn(hip.hipGraphDestroy(graph), "hipGraphDestroy");
255 ✗ hip.Check(instantiated, "hipGraphInstantiateWithFlags");
256 ✗ return exec;
257 }
258
259 ✗ void Traits::LaunchGraph(const Api& hip, hipGraphExec_t graph, hipStream_t stream) {
260 ✗ hip.Check(hip.hipGraphLaunch(graph, stream), "hipGraphLaunch");
261 ✗ }
262
263 ✗ void Traits::DestroyGraph(const Api& hip, hipGraphExec_t graph) {
264 ✗ hip.Warn(hip.hipGraphExecDestroy(graph), "hipGraphExecDestroy");
265 ✗ }
266
267 ✗ const compute::RuntimeCompiler& Traits::RuntimeCompiler() {
268 ✗ return GetHiprtc();
269 }
270
271 ✗ std::string_view Traits::Prelude() {
272 ✗ return {reinterpret_cast<const char*>(chrono_batch_render_hip_prelude),
273 ✗ chrono_batch_render_hip_prelude_size};
274 }
275
276 ✗ std::vector<std::string> Traits::CompileOptions(const std::string& architecture, Size3 groupSize,
277 uint32_t, bool lineInfo) {
278 ✗ const uint32_t threads = groupSize.x * groupSize.y * groupSize.z;
279 ✗ std::vector<std::string> options = {"--offload-arch=" + architecture, "-O3", "-std=c++17",
280 ✗ "--gpu-max-threads-per-block=" + std::to_string(threads),
281
0/6
✗ Branch 1 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 9 not taken.
✗ Branch 10 not taken.
 ✗ "-DSLANG_CUDA_ENABLE_FAST_MATH=1"};
282 ✗ if (lineInfo) {
283
0/1
✗ Branch 1 not taken.
 ✗ options.emplace_back("-gline-tables-only");
284 }
285 ✗ return options;
286 ✗ }
287
288 ✗ hipModule_t Traits::LoadModule(const Api& hip, const std::string& binary) {
289 ✗ hipModule_t module = nullptr;
290 ✗ hip.Check(hip.hipModuleLoadData(&module, binary.data()), "hipModuleLoadData");
291 ✗ return module;
292 }
293
294 ✗ void Traits::UnloadModule(const Api& hip, hipModule_t module) {
295 ✗ hip.Warn(hip.hipModuleUnload(module), "hipModuleUnload");
296 ✗ }
297
298 ✗ std::pair<uint64_t, size_t> Traits::Symbol(const Api& hip, hipModule_t module, const char* name) {
299 ✗ hipDeviceptr_t pointer = nullptr;
300 ✗ size_t bytes = 0;
301 ✗ const hipError_t pending = hip.hipPeekAtLastError();
302 ✗ const hipError_t found = hip.hipModuleGetGlobal(&pointer, &bytes, module, name);
303
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (found == hipErrorNotFound) {
304
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (pending == hipSuccess) {
305 ✗ hip.hipGetLastError();
306 }
307 ✗ return {0, 0};
308 }
309 ✗ hip.Check(found, "hipModuleGetGlobal");
310 ✗ return {reinterpret_cast<uint64_t>(pointer), bytes};
311 }
312
313 ✗ void Traits::CopyToDevice(const Api& hip, uint64_t dst, const void* host, uint64_t bytes) {
314 ✗ hip.Check(hip.hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(dst), host, bytes),
315 "hipMemcpyHtoD");
316 ✗ }
317
318 ✗ hipFunction_t Traits::GetFunction(const Api& hip, hipModule_t module, const char* name) {
319 ✗ hipFunction_t function = nullptr;
320 ✗ hip.Check(hip.hipModuleGetFunction(&function, module, name), "hipModuleGetFunction");
321 ✗ return function;
322 }
323
324 ✗ KernelStats Traits::Stats(const Api& hip, hipFunction_t function) {
325 ✗ const auto attribute = [&](int which) {
326 ✗ int value = 0;
327 ✗ hip.Check(hip.hipFuncGetAttribute(&value, which, function), "hipFuncGetAttribute");
328 ✗ return uint32_t(value);
329 ✗ };
330 ✗ return {.registers = attribute(HIP_FUNC_ATTRIBUTE_NUM_REGS),
331 ✗ .spillBytes = attribute(HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES),
332 ✗ .sharedBytes = attribute(HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES)};
333 ✗ }
334
335 ✗ void Traits::Launch(const Api& hip, hipFunction_t function, Size3 groups, Size3 threads,
336 void** extra, hipStream_t stream) {
337 ✗ hip.Check(hip.hipModuleLaunchKernel(function, groups.x, groups.y, groups.z, threads.x,
338 threads.y, threads.z, 0, stream, nullptr, extra),
339 "hipModuleLaunchKernel");
340 ✗ }
341
342 }
343