device/cuda/device.cpp - GCC Code Coverage Report

device/cuda/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/cuda/device.h"
9			#include "chrono_batch_render/device/queue.h"
10
11			extern "C" {
12			extern const unsigned char chrono_batch_render_cuda_prelude[];
13			extern const std::size_t chrono_batch_render_cuda_prelude_size;
14			}
15
16			namespace chrono::batch_render::device::cuda {
17
18		✗	ContextScope::ContextScope(const Api& cuda, CUcontext context) : m_cuda(cuda) {
19		✗	cuda.Check(cuda.cuCtxPushCurrent(context), "cuCtxPushCurrent");
20		✗	}
21
22		✗	ContextScope::~ContextScope() {
23		✗	CUcontext previous = nullptr;
24		✗	m_cuda.Warn(m_cuda.cuCtxPopCurrent(&previous), "cuCtxPopCurrent");
25		✗	}
26
27		✗	CudaDevice::CudaDevice(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_cuda(GetApi()) {
29	0/2 ✗ Branch 1 not taken. ✗ Branch 4 not taken.	✗	m_cuda.Check(m_cuda.cuDeviceGet(&m_device, int(adapter.ordinal)), "cuDeviceGet");
30	0/2 ✗ Branch 1 not taken. ✗ Branch 4 not taken.	✗	m_cuda.Check(m_cuda.cuDevicePrimaryCtxRetain(&m_context, m_device), "cuDevicePrimaryCtxRetain");
31		✗	try {
32	0/1 ✗ Branch 1 not taken.	✗	m_queue = std::make_unique<compute::ComputeQueue<Traits>>(*this);
33		✗	} catch (...) {
34	0/2 ✗ Branch 1 not taken. ✗ Branch 4 not taken.	✗	m_cuda.Warn(m_cuda.cuDevicePrimaryCtxRelease(m_device), "cuDevicePrimaryCtxRelease");
35		✗	throw;
36		✗	}
37		✗	}
38
39		✗	CudaDevice::~CudaDevice() {
40		✗	m_modules.Clear();
41		✗	m_cuda.Warn(m_cuda.cuDevicePrimaryCtxRelease(m_device), "cuDevicePrimaryCtxRelease");
42		✗	}
43
44		✗	Ref<Buffer> CudaDevice::NewBuffer(const BufferDesc& desc) {
45		✗	return compute::ComputeBuffer<Traits>::Create(*this, desc);
46			}
47
48		✗	Ref<Texture> CudaDevice::NewTexture(const TextureDesc& desc) {
49	0/1 ✗ Branch 2 not taken.	✗	return Ref<Texture>(new compute::ComputeTexture<Traits>(*this, desc));
50			}
51
52		✗	void CudaDevice::ExternalCopy(BufferRange dst, uint64_t source, Stream producer) {
53		✗	compute::ExternalCopy<Traits>(*this, dst, source, producer);
54		✗	}
55
56		✗	std::unique_ptr<Kernel> CudaDevice::NewKernel() {
57		✗	return std::make_unique<compute::ComputeKernel<Traits>>(*this);
58			}
59
60		✗	std::vector<MemoryHeap> CudaDevice::Heaps() const {
61		✗	const ContextScope scope = Traits::Enter(*this);
62		✗	size_t free = 0;
63		✗	size_t total = 0;
64	0/3 ✗ Branch 1 not taken. ✗ Branch 3 not taken. ✗ Branch 4 not taken.	✗	if (Functions().cuMemGetInfo(&free, &total) != 0) {
65		✗	return {};
66			}
67	0/1 ✗ Branch 1 not taken.	✗	return {{.usage = total - free, .budget = total, .size = total, .deviceLocal = true}};
68		✗	}
69
70		✗	ContextScope Traits::Enter(const CudaDevice& device) {
71		✗	return ContextScope(device.Functions(), device.Context());
72			}
73
74		✗	uint64_t Traits::Allocate(const Api& cuda, uint64_t bytes) {
75		✗	CUdeviceptr pointer = 0;
76		✗	cuda.Check(cuda.cuMemAlloc(&pointer, bytes), "cuMemAlloc");
77		✗	return pointer;
78			}
79
80		✗	void Traits::Free(const Api& cuda, uint64_t address) {
81		✗	cuda.Warn(cuda.cuMemFree(address), "cuMemFree");
82		✗	}
83
84		✗	CUmemoryPool Traits::CreatePool(const CudaDevice& device, uint64_t keptBytes) {
85		✗	const Api& cuda = device.Functions();
86		✗	int pools = 0;
87		✗	cuda.Check(cuda.cuDeviceGetAttribute(&pools, CU_DEVICE_ATTRIBUTE_MEMORY_POOLS_SUPPORTED,
88			device.Handle()),
89			"cuDeviceGetAttribute");
90	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (pools == 0) {
91			return nullptr;
92			}
93
94		✗	CUmemPoolProps properties{};
95		✗	properties.allocType = CU_MEM_ALLOCATION_TYPE_PINNED;
96		✗	properties.location = {CU_MEM_LOCATION_TYPE_DEVICE, device.Handle()};
97		✗	CUmemoryPool pool = nullptr;
98		✗	cuda.Check(cuda.cuMemPoolCreate(&pool, &properties), "cuMemPoolCreate");
99
100		✗	const CUresult set =
101		✗	cuda.cuMemPoolSetAttribute(pool, CU_MEMPOOL_ATTR_RELEASE_THRESHOLD, &keptBytes);
102	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (set != CUDA_SUCCESS) {
103		✗	DestroyPool(cuda, pool);
104		✗	cuda.Check(set, "cuMemPoolSetAttribute");
105			}
106		✗	return pool;
107			}
108
109		✗	void Traits::DestroyPool(const Api& cuda, CUmemoryPool pool) {
110		✗	cuda.Warn(cuda.cuMemPoolDestroy(pool), "cuMemPoolDestroy");
111		✗	}
112
113		✗	uint64_t Traits::AllocateFrom(const Api& cuda, CUmemoryPool pool, uint64_t bytes, CUstream stream) {
114		✗	CUdeviceptr pointer = 0;
115		✗	cuda.Check(cuda.cuMemAllocFromPoolAsync(&pointer, bytes, pool, stream),
116			"cuMemAllocFromPoolAsync");
117		✗	return pointer;
118			}
119
120		✗	void Traits::FreeTo(const Api& cuda, uint64_t address, CUstream stream) {
121		✗	cuda.Warn(cuda.cuMemFreeAsync(address, stream), "cuMemFreeAsync");
122		✗	}
123
124		✗	void Traits::SynchronizeStream(const Api& cuda, CUstream stream) {
125		✗	cuda.Check(cuda.cuStreamSynchronize(stream), "cuStreamSynchronize");
126		✗	}
127
128		✗	void* Traits::AllocateHost(const Api& cuda, uint64_t bytes, bool writeCombined) {
129		✗	void* host = nullptr;
130		✗	const unsigned flags =
131	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	CU_MEMHOSTALLOC_DEVICEMAP \| (writeCombined ? CU_MEMHOSTALLOC_WRITECOMBINED : 0);
132		✗	cuda.Check(cuda.cuMemHostAlloc(&host, bytes, flags), "cuMemHostAlloc");
133		✗	return host;
134			}
135
136		✗	uint64_t Traits::DeviceAddress(const Api& cuda, void* host) {
137		✗	CUdeviceptr pointer = 0;
138		✗	cuda.Check(cuda.cuMemHostGetDevicePointer(&pointer, host, 0), "cuMemHostGetDevicePointer");
139		✗	return pointer;
140			}
141
142		✗	void Traits::FreeHost(const Api& cuda, void* host) {
143		✗	cuda.Warn(cuda.cuMemFreeHost(host), "cuMemFreeHost");
144		✗	}
145
146		✗	CUstream Traits::CreateStream(const Api& cuda) {
147		✗	CUstream stream = nullptr;
148		✗	cuda.Check(cuda.cuStreamCreate(&stream, CU_STREAM_NON_BLOCKING), "cuStreamCreate");
149		✗	return stream;
150			}
151
152		✗	void Traits::DestroyStream(const Api& cuda, CUstream stream) {
153		✗	cuda.Warn(cuda.cuStreamDestroy(stream), "cuStreamDestroy");
154		✗	}
155
156		✗	CUevent Traits::CreateEvent(const Api& cuda, bool timed) {
157		✗	CUevent event = nullptr;
158	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	cuda.Check(cuda.cuEventCreate(&event, timed ? 0 : CU_EVENT_DISABLE_TIMING), "cuEventCreate");
159		✗	return event;
160			}
161
162		✗	void Traits::DestroyEvent(const Api& cuda, CUevent event) {
163		✗	cuda.Warn(cuda.cuEventDestroy(event), "cuEventDestroy");
164		✗	}
165
166		✗	void Traits::Record(const Api& cuda, CUevent event, CUstream stream, bool node) {
167	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	cuda.Check(cuda.cuEventRecordWithFlags(event, stream, node ? CU_EVENT_RECORD_EXTERNAL : 0),
168			"cuEventRecordWithFlags");
169		✗	}
170
171		✗	void Traits::WaitEvent(const Api& cuda, CUstream stream, CUevent event) {
172		✗	cuda.Check(cuda.cuStreamWaitEvent(stream, event, 0), "cuStreamWaitEvent");
173		✗	}
174
175		✗	CUstream Traits::ApplicationStream(void* handle) {
176		✗	return static_cast<CUstream>(handle);
177			}
178
179		✗	uint64_t Traits::ElapsedNs(const Api& cuda, CUevent from, CUevent to) {
180		✗	float milliseconds = 0;
181		✗	cuda.Check(cuda.cuEventElapsedTime(&milliseconds, from, to), "cuEventElapsedTime");
182		✗	return uint64_t(double(milliseconds) * 1e6);
183			}
184
185		✗	bool Traits::Query(const Api& cuda, CUevent event) {
186		✗	const CUresult result = cuda.cuEventQuery(event);
187	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (result == CUDA_ERROR_NOT_READY) {
188			return false;
189			}
190		✗	cuda.Check(result, "cuEventQuery");
191		✗	return true;
192			}
193
194		✗	void Traits::Synchronize(const Api& cuda, CUevent event) {
195		✗	cuda.Check(cuda.cuEventSynchronize(event), "cuEventSynchronize");
196		✗	}
197
198		✗	void Traits::Copy(const Api& cuda, uint64_t dst, uint64_t src, uint64_t bytes, CUstream stream) {
199		✗	cuda.Check(cuda.cuMemcpyAsync(dst, src, bytes, stream), "cuMemcpyAsync");
200		✗	}
201
202		✗	void Traits::Fill(const Api& cuda, uint64_t dst, uint32_t value, uint64_t bytes, CUstream stream) {
203		✗	cuda.Check(cuda.cuMemsetD32Async(dst, value, bytes / 4, stream), "cuMemsetD32Async");
204		✗	}
205
206		✗	void Traits::Upload(const Api& cuda, uint64_t dst, const void* host, uint64_t bytes,
207			CUstream stream) {
208		✗	cuda.Check(cuda.cuMemcpyHtoDAsync(dst, host, bytes, stream), "cuMemcpyHtoDAsync");
209		✗	}
210
211		✗	void* Traits::CreateTextureArray(const Api& cuda, const ArrayDesc& desc, uint32_t levels) {
212		✗	void* array = nullptr;
213		✗	cuda.Check(cuda.cuMipmappedArrayCreate(&array, &desc, levels), "cuMipmappedArrayCreate");
214		✗	return array;
215			}
216
217		✗	void Traits::DestroyTextureArray(const Api& cuda, void* array) {
218		✗	cuda.Warn(cuda.cuMipmappedArrayDestroy(array), "cuMipmappedArrayDestroy");
219		✗	}
220
221		✗	void* Traits::TextureLevel(const Api& cuda, void* array, uint32_t level) {
222		✗	void* found = nullptr;
223		✗	cuda.Check(cuda.cuMipmappedArrayGetLevel(&found, array, level), "cuMipmappedArrayGetLevel");
224		✗	return found;
225			}
226
227		✗	uint64_t Traits::CreateTextureObject(const Api& cuda, const ResourceDesc& resource,
228			const SamplerDesc& sampler) {
229		✗	unsigned long long object = 0;
230		✗	cuda.Check(cuda.cuTexObjectCreate(&object, &resource, &sampler, nullptr), "cuTexObjectCreate");
231		✗	return object;
232			}
233
234		✗	void Traits::DestroyTextureObject(const Api& cuda, uint64_t object) {
235		✗	cuda.Warn(cuda.cuTexObjectDestroy(object), "cuTexObjectDestroy");
236		✗	}
237
238		✗	void Traits::CopyToArray(const Api& cuda, void* level, uint64_t source, uint64_t rowBytes,
239			uint32_t rows, CUstream stream) {
240		✗	const CUDA_MEMCPY2D copy{.srcMemoryType = CU_MEMORYTYPE_DEVICE,
241			.srcDevice = source,
242			.srcPitch = rowBytes,
243			.dstMemoryType = CU_MEMORYTYPE_ARRAY,
244			.dstArray = level,
245			.WidthInBytes = rowBytes,
246		✗	.Height = rows};
247		✗	cuda.Check(cuda.cuMemcpy2DAsync(&copy, stream), "cuMemcpy2DAsync");
248		✗	}
249
250		✗	void Traits::BeginCapture(const Api& cuda, CUstream stream) {
251		✗	cuda.Check(cuda.cuStreamBeginCapture(stream, CU_STREAM_CAPTURE_MODE_THREAD_LOCAL),
252			"cuStreamBeginCapture");
253		✗	}
254
255		✗	CUgraphExec Traits::EndCapture(const Api& cuda, CUstream stream) {
256		✗	CUgraph graph = nullptr;
257		✗	cuda.Check(cuda.cuStreamEndCapture(stream, &graph), "cuStreamEndCapture");
258		✗	CUgraphExec exec = nullptr;
259		✗	const auto instantiated = cuda.cuGraphInstantiateWithFlags(&exec, graph, 0);
260		✗	cuda.Warn(cuda.cuGraphDestroy(graph), "cuGraphDestroy");
261		✗	cuda.Check(instantiated, "cuGraphInstantiateWithFlags");
262		✗	return exec;
263			}
264
265		✗	void Traits::LaunchGraph(const Api& cuda, CUgraphExec graph, CUstream stream) {
266		✗	cuda.Check(cuda.cuGraphLaunch(graph, stream), "cuGraphLaunch");
267		✗	}
268
269		✗	void Traits::DestroyGraph(const Api& cuda, CUgraphExec graph) {
270		✗	cuda.Warn(cuda.cuGraphExecDestroy(graph), "cuGraphExecDestroy");
271		✗	}
272
273		✗	const compute::RuntimeCompiler& Traits::RuntimeCompiler() {
274		✗	return GetNvrtc();
275			}
276
277		✗	std::string_view Traits::Prelude() {
278		✗	return {reinterpret_cast<const char*>(chrono_batch_render_cuda_prelude),
279		✗	chrono_batch_render_cuda_prelude_size};
280			}
281
282		✗	std::vector<std::string> Traits::CompileOptions(const std::string& architecture, Size3,
283			uint32_t maxRegisters, bool lineInfo) {
284		✗	std::vector<std::string> options = {
285			"--gpu-architecture=" + architecture, "--std=c++17", "-default-device",
286	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", "-prec-div=false", "-prec-sqrt=false"};
287		✗	if (maxRegisters > 0) {
288	0/1 ✗ Branch 2 not taken.	✗	options.push_back("-maxrregcount=" + std::to_string(maxRegisters));
289			}
290	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (lineInfo) {
291	0/1 ✗ Branch 1 not taken.	✗	options.emplace_back("-lineinfo");
292			}
293		✗	return options;
294		✗	}
295
296		✗	CUmodule Traits::LoadModule(const Api& cuda, const std::string& binary) {
297		✗	CUmodule module = nullptr;
298		✗	cuda.Check(cuda.cuModuleLoadData(&module, binary.data()), "cuModuleLoadData");
299		✗	return module;
300			}
301
302		✗	void Traits::UnloadModule(const Api& cuda, CUmodule module) {
303		✗	cuda.Warn(cuda.cuModuleUnload(module), "cuModuleUnload");
304		✗	}
305
306		✗	std::pair<uint64_t, size_t> Traits::Symbol(const Api& cuda, CUmodule module, const char* name) {
307		✗	CUdeviceptr pointer = 0;
308		✗	size_t bytes = 0;
309		✗	const CUresult found = cuda.cuModuleGetGlobal(&pointer, &bytes, module, name);
310	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (found == CUDA_ERROR_NOT_FOUND) {
311		✗	return {0, 0};
312			}
313		✗	cuda.Check(found, "cuModuleGetGlobal");
314		✗	return {pointer, bytes};
315			}
316
317		✗	void Traits::CopyToDevice(const Api& cuda, uint64_t dst, const void* host, uint64_t bytes) {
318		✗	cuda.Check(cuda.cuMemcpyHtoD(dst, host, bytes), "cuMemcpyHtoD");
319		✗	}
320
321		✗	CUfunction Traits::GetFunction(const Api& cuda, CUmodule module, const char* name) {
322		✗	CUfunction function = nullptr;
323		✗	cuda.Check(cuda.cuModuleGetFunction(&function, module, name), "cuModuleGetFunction");
324		✗	return function;
325			}
326
327		✗	KernelStats Traits::Stats(const Api& cuda, CUfunction function) {
328		✗	const auto attribute = [&](int which) {
329		✗	int value = 0;
330		✗	cuda.Check(cuda.cuFuncGetAttribute(&value, which, function), "cuFuncGetAttribute");
331		✗	return uint32_t(value);
332		✗	};
333		✗	return {.registers = attribute(CU_FUNC_ATTRIBUTE_NUM_REGS),
334		✗	.spillBytes = attribute(CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES),
335		✗	.sharedBytes = attribute(CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES)};
336		✗	}
337
338		✗	void Traits::Launch(const Api& cuda, CUfunction function, Size3 groups, Size3 threads, void** extra,
339			CUstream stream) {
340		✗	cuda.Check(cuda.cuLaunchKernel(function, groups.x, groups.y, groups.z, threads.x, threads.y,
341			threads.z, 0, stream, nullptr, extra),
342			"cuLaunchKernel");
343		✗	}
344
345			}
346