device/slang/compiler.cpp - GCC Code Coverage Report

device/slang/compiler.cpp

Line	Branch	Exec	Source
1			// SPDX-License-Identifier: BSD-3-Clause
2			#include "chrono_batch_render/device/slang/compiler.h"
3			#include "chrono_batch_render/device/slang/reflection.h"
4
5			#include <algorithm>
6			#include <charconv>
7			#include <cstddef>
8			#include <iterator>
9			#include <optional>
10			#include <regex>
11			#include <set>
12			#include <vector>
13
14			#include <fmt/format.h>
15
16			#include <slang-com-ptr.h>
17			#include <slang.h>
18
19			#include "chrono_batch_render/device/slang/debug.h"
20			#include "chrono_batch_render/util/disk_cache.h"
21			#include "chrono_batch_render/util/environment.h"
22			#include "chrono_batch_render/util/error.h"
23			#include "chrono_batch_render/util/log.h"
24			#include "chrono_batch_render/util/text.h"
25
26			namespace chrono::batch_render::device {
27
28			namespace {
29
30			using Slang::ComPtr;
31
32		15	std::string Text(slang::IBlob* diagnostics) {
33	2/2 ✓ Branch 0 taken 7 times. ✓ Branch 1 taken 8 times.	15	if (diagnostics == nullptr) {
34		7	return {};
35			}
36		8	return {static_cast<const char*>(diagnostics->getBufferPointer()),
37		8	diagnostics->getBufferSize()};
38			}
39
40		8	[[noreturn]] void Fail(std::string_view what, std::string_view label, slang::IBlob* diagnostics) {
41	2/2 ✓ Branch 2 taken 8 times. ✓ Branch 5 taken 8 times.	16	throw Error("{} of program \"{}\" failed:\n{}", what, label, Text(diagnostics));
42			}
43
44		56	slang::IGlobalSession& GlobalSession() {
45		6	thread_local ComPtr<slang::IGlobalSession> global = [] {
46	1/1 ✓ Branch 1 taken 6 times.	6	ComPtr<slang::IGlobalSession> created;
47	2/3 ✓ Branch 1 taken 6 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 6 times.	6	if (SLANG_FAILED(slang::createGlobalSession(created.writeRef()))) {
48	0/2 ✗ Branch 2 not taken. ✗ Branch 5 not taken.	✗	throw Error("Slang failed to start");
49			}
50	1/1 ✓ Branch 1 taken 6 times.	6	const std::string prelude = fmt::format("#include \"{}\"\n", kPreludeHeader);
51	1/1 ✓ Branch 1 taken 6 times.	6	created->setLanguagePrelude(SLANG_SOURCE_LANGUAGE_CUDA, prelude.c_str());
52		6	return created;
53	2/2 ✓ Branch 0 taken 6 times. ✓ Branch 1 taken 50 times.	62	}();
54		56	return *global;
55			}
56
57		147	slang::CompilerOptionEntry Enable(slang::CompilerOptionName name) {
58		147	slang::CompilerOptionEntry entry{};
59		147	entry.name = name;
60		147	entry.value.kind = slang::CompilerOptionValueKind::Int;
61		147	entry.value.intValue0 = 1;
62		147	return entry;
63			}
64
65			constexpr std::string_view kExtension = ".slang";
66
67			constexpr int kEntryFormat = 5;
68
69		84896	std::string FilePath(std::string_view path) {
70		84896	std::string file(path);
71	4/4 ✓ Branch 0 taken 22700 times. ✓ Branch 1 taken 62196 times. ✓ Branch 2 taken 226 times. ✓ Branch 3 taken 22474 times.	84896	if (file.ends_with(kExtension) && file.find('/') == std::string::npos) {
72		452	std::replace(file.begin(), file.end() - ptrdiff_t(kExtension.size()), '.', '/');
73			}
74		84896	return file;
75			}
76
77		327573	std::optional<std::string_view> FindEmbedded(const KernelLibrary& library, std::string_view file) {
78	2/2 ✓ Branch 0 taken 7656179 times. ✓ Branch 1 taken 291403 times.	7947582	for (const EmbeddedKernel& embedded : library.files) {
79	2/2 ✓ Branch 0 taken 36170 times. ✓ Branch 1 taken 7620009 times.	7656179	if (file == embedded.path) {
80		36170	return std::string_view(reinterpret_cast<const char*>(embedded.bytes), embedded.size);
81			}
82			}
83	2/2 ✓ Branch 0 taken 270212 times. ✓ Branch 1 taken 263129 times.	533341	for (const KernelLibrary* used : library.uses) {
84	2/2 ✓ Branch 1 taken 28274 times. ✓ Branch 2 taken 241938 times.	270212	if (const std::optional<std::string_view> found = FindEmbedded(*used, file)) {
85		28274	return found;
86			}
87			}
88		263129	return std::nullopt;
89			}
90
91		57820	std::optional<std::string_view> FindFile(const KernelLibrary* library, const SourceFiles& sources,
92			std::string_view path) {
93		57820	const std::string file = FilePath(path);
94	2/2 ✓ Branch 1 taken 459 times. ✓ Branch 2 taken 57361 times.	57820	if (const auto found = sources.find(file); found != sources.end()) {
95		459	return std::optional<std::string_view>(found->second);
96			}
97	2/2 ✓ Branch 0 taken 62 times. ✓ Branch 1 taken 57299 times.	57423	return FindEmbedded(library != nullptr ? *library : kDeviceKernels, file);
98		57820	}
99
100		27076	std::string Normal(std::string_view path) {
101		27076	std::vector<std::string_view> segments;
102	2/2 ✓ Branch 0 taken 133299 times. ✓ Branch 1 taken 27076 times.	160375	while (!path.empty()) {
103		133299	const size_t slash = path.find('/');
104	1/1 ✓ Branch 1 taken 133299 times.	133299	const std::string_view segment = path.substr(0, slash);
105	1/4 ✗ Branch 0 not taken. ✓ Branch 1 taken 133299 times. ✗ Branch 2 not taken. ✗ Branch 3 not taken.	133299	if (segment == ".." && !segments.empty()) {
106		✗	segments.pop_back();
107	3/4 ✓ Branch 0 taken 131978 times. ✓ Branch 1 taken 1321 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 131978 times.	133299	} else if (segment != "." && !segment.empty()) {
108	1/1 ✓ Branch 1 taken 131978 times.	131978	segments.push_back(segment);
109			}
110	3/3 ✓ Branch 0 taken 27076 times. ✓ Branch 1 taken 106223 times. ✓ Branch 3 taken 106223 times.	133299	path = slash == std::string_view::npos ? std::string_view() : path.substr(slash + 1);
111			}
112		27076	std::string normal;
113	2/2 ✓ Branch 0 taken 131978 times. ✓ Branch 1 taken 27076 times.	159054	for (const std::string_view segment : segments) {
114	4/4 ✓ Branch 0 taken 104902 times. ✓ Branch 1 taken 27076 times. ✓ Branch 3 taken 131978 times. ✓ Branch 6 taken 131978 times.	236880	normal.append(normal.empty() ? "" : "/").append(segment);
115			}
116		27076	return normal;
117		27076	}
118
119			constexpr std::string_view kBinaryExtension = ".slang-module";
120
121			class ModuleFiles final : public ISlangFileSystem {
122			public:
123		49	ModuleFiles(const KernelLibrary* library, const SourceFiles& sources,
124			const DiskCache* modules = nullptr, std::string key = {})
125		49	: m_library(library), m_sources(sources), m_modules(modules), m_key(std::move(key)) {}
126
127		98	SLANG_NO_THROW SlangResult SLANG_MCALL queryInterface(const SlangUUID& uuid,
128			void** object) override {
129		98	*object = castAs(uuid);
130	1/2 ✓ Branch 0 taken 98 times. ✗ Branch 1 not taken.	98	return *object != nullptr ? SLANG_OK : SLANG_E_NO_INTERFACE;
131			}
132
133		98	SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override {
134		98	return 1;
135			}
136
137		98	SLANG_NO_THROW uint32_t SLANG_MCALL release() override {
138		98	return 1;
139			}
140
141		147	SLANG_NO_THROW void* SLANG_MCALL castAs(const SlangUUID& uuid) override {
142	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 147 times.	147	const bool known = uuid == ISlangUnknown::getTypeGuid() \|\|
143	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 147 times.	147	uuid == ISlangCastable::getTypeGuid() \|\|
144	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 147 times.	147	uuid == ISlangFileSystem::getTypeGuid();
145		147	return known ? static_cast<ISlangFileSystem*>(this) : nullptr;
146			}
147
148		16521	SLANG_NO_THROW SlangResult SLANG_MCALL loadFile(const char* path, ISlangBlob** blob) override {
149	1/1 ✓ Branch 2 taken 16521 times.	16521	std::string file = FilePath(Normal(path));
150	2/2 ✓ Branch 0 taken 3396 times. ✓ Branch 1 taken 13125 times.	16521	if (file.ends_with(kBinaryExtension)) {
151	1/1 ✓ Branch 1 taken 3396 times.	3396	return LoadBinary(file, blob);
152			}
153	1/1 ✓ Branch 1 taken 13125 times.	13125	const std::optional<std::string_view> found = FindFile(m_library, m_sources, file);
154	2/2 ✓ Branch 0 taken 2929 times. ✓ Branch 1 taken 10196 times.	13125	if (!found) {
155			return SLANG_E_NOT_FOUND;
156			}
157	1/1 ✓ Branch 2 taken 2929 times.	2929	Record(std::move(file), *found);
158	1/1 ✓ Branch 1 taken 2929 times.	2929	*blob = slang_createBlob(found->data(), found->size());
159		2929	return SLANG_OK;
160		16521	}
161
162		7	std::vector<std::pair<std::string, std::string>> Read() const {
163		7	return m_read;
164			}
165
166		7	void RecordDependencies(slang::ISession& session) {
167	2/2 ✓ Branch 1 taken 216 times. ✓ Branch 2 taken 7 times.	223	for (SlangInt index = 0; index < session.getLoadedModuleCount(); ++index) {
168		216	slang::IModule* module = session.getLoadedModule(index);
169	2/2 ✓ Branch 2 taken 9137 times. ✓ Branch 3 taken 216 times.	9353	for (SlangInt32 dependency = 0; dependency < module->getDependencyFileCount();
170			++dependency) {
171	1/1 ✓ Branch 3 taken 9137 times.	9137	std::string file = FilePath(Normal(module->getDependencyFilePath(dependency)));
172	1/1 ✓ Branch 1 taken 9137 times.	9137	if (const std::optional<std::string_view> found =
173	2/3 ✓ Branch 1 taken 9137 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 9137 times.	9137	FindFile(m_library, m_sources, file)) {
174	0/1 ✗ Branch 2 not taken.	✗	Record(std::move(file), *found);
175			}
176		9137	}
177			}
178		7	}
179
180		40	void StoreBinaries(slang::ISession& session) const {
181	1/2 ✓ Branch 0 taken 40 times. ✗ Branch 1 not taken.	40	if (m_modules == nullptr) {
182			return;
183			}
184	2/2 ✓ Branch 1 taken 1418 times. ✓ Branch 2 taken 40 times.	1458	for (SlangInt index = 0; index < session.getLoadedModuleCount(); ++index) {
185		1418	slang::IModule* module = session.getLoadedModule(index);
186		1418	const char* path = module->getFilePath();
187	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1418 times.	1418	if (path == nullptr) {
188		1259	continue;
189			}
190	1/1 ✓ Branch 2 taken 1418 times.	1418	const std::string file = FilePath(Normal(path));
191	5/6 ✓ Branch 0 taken 230 times. ✓ Branch 1 taken 1188 times. ✓ Branch 3 taken 230 times. ✗ Branch 4 not taken. ✓ Branch 5 taken 64 times. ✓ Branch 6 taken 166 times.	1648	if (!file.ends_with(kExtension) \|\| m_served.contains(file) \|\|
192		230	m_sources.contains(file)) {
193		1252	continue;
194			}
195	1/1 ✓ Branch 1 taken 166 times.	166	const std::optional<std::string_view> found = FindFile(m_library, m_sources, file);
196	2/2 ✓ Branch 0 taken 159 times. ✓ Branch 1 taken 7 times.	166	ComPtr<ISlangBlob> binary;
197	4/5 ✓ Branch 0 taken 159 times. ✓ Branch 1 taken 7 times. ✓ Branch 3 taken 159 times. ✓ Branch 5 taken 159 times. ✗ Branch 6 not taken.	166	if (!found \|\| SLANG_FAILED(module->serialize(binary.writeRef()))) {
198	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	continue;
199			}
200	4/4 ✓ Branch 1 taken 159 times. ✓ Branch 4 taken 159 times. ✓ Branch 7 taken 159 times. ✓ Branch 10 taken 159 times.	159	m_modules->Store(
201	2/5 ✓ Branch 1 taken 159 times. ✓ Branch 3 taken 159 times. ✗ Branch 4 not taken. ✗ Branch 5 not taken. ✗ Branch 6 not taken.	318	BinaryKey(file, *found),
202	2/2 ✓ Branch 1 taken 159 times. ✓ Branch 4 taken 159 times.	159	{static_cast<const char*>(binary->getBufferPointer()), binary->getBufferSize()});
203		1425	}
204			}
205
206			private:
207		2929	void Record(std::string file, std::string_view source) {
208	1/2 ✓ Branch 0 taken 2929 times. ✗ Branch 1 not taken.	2929	if (std::ranges::find(m_read, file, &std::pair<std::string, std::string>::first) ==
209	1/2 ✓ Branch 1 taken 2929 times. ✗ Branch 2 not taken.	2929	m_read.end()) {
210	1/1 ✓ Branch 2 taken 2929 times.	5858	m_read.emplace_back(std::move(file), Digest(source));
211			}
212		2929	}
213
214		1668	std::string BinaryKey(std::string_view file, std::string_view source) const {
215	1/1 ✓ Branch 2 taken 1668 times.	3336	return fmt::format("{} module {} {}", m_key, file, Digest(source));
216			}
217
218		3396	SlangResult LoadBinary(const std::string& file, ISlangBlob** blob) {
219	1/2 ✓ Branch 0 taken 3396 times. ✗ Branch 1 not taken.	3396	if (m_modules == nullptr) {
220			return SLANG_E_NOT_FOUND;
221			}
222		3396	const std::string source =
223	2/2 ✓ Branch 2 taken 3396 times. ✓ Branch 5 taken 3396 times.	6792	file.substr(0, file.size() - kBinaryExtension.size()) + std::string(kExtension);
224	2/2 ✓ Branch 1 taken 3359 times. ✓ Branch 2 taken 37 times.	3396	if (m_sources.contains(source)) {
225			return SLANG_E_NOT_FOUND;
226			}
227	1/1 ✓ Branch 1 taken 3359 times.	3359	const std::optional<std::string_view> found = FindFile(m_library, m_sources, source);
228	2/2 ✓ Branch 0 taken 1509 times. ✓ Branch 1 taken 1850 times.	3359	if (!found) {
229			return SLANG_E_NOT_FOUND;
230			}
231	2/2 ✓ Branch 1 taken 1509 times. ✓ Branch 4 taken 1509 times.	1509	const std::optional<std::string> binary = m_modules->Load(BinaryKey(source, *found));
232	2/2 ✓ Branch 0 taken 1310 times. ✓ Branch 1 taken 199 times.	1509	if (!binary) {
233			return SLANG_E_NOT_FOUND;
234			}
235	1/1 ✓ Branch 1 taken 1310 times.	1310	m_served.insert(source);
236	1/1 ✓ Branch 1 taken 1310 times.	1310	*blob = slang_createBlob(binary->data(), binary->size());
237		1310	return SLANG_OK;
238		4905	}
239
240			const KernelLibrary* m_library;
241			const SourceFiles& m_sources;
242			const DiskCache* m_modules;
243			std::string m_key;
244			std::vector<std::pair<std::string, std::string>> m_read;
245			std::set<std::string> m_served;
246			};
247
248		7	void Reflect(slang::EntryPointReflection& reflected, slang::ProgramLayout& layout,
249			std::string_view entry, std::string_view label, KernelCode& kernel) {
250		7	SlangUInt groupSize[3] = {};
251		7	reflected.getComputeThreadGroupSize(3, groupSize);
252	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	kernel.shape.groupSize = {uint32_t(std::max<SlangUInt>(groupSize[0], 1)),
253	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	uint32_t(std::max<SlangUInt>(groupSize[1], 1)),
254	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	uint32_t(std::max<SlangUInt>(groupSize[2], 1))};
255	1/2 ✓ Branch 1 taken 7 times. ✗ Branch 2 not taken.	7	if (slang::FunctionReflection* function = reflected.getFunction()) {
256		7	int count = 0;
257		7	slang::Attribute* limit = function->findAttributeByName(&GlobalSession(), "MaxRegisters");
258	3/4 ✓ Branch 0 taken 2 times. ✓ Branch 1 taken 5 times. ✓ Branch 2 taken 2 times. ✗ Branch 3 not taken.	9	if (limit != nullptr && SLANG_SUCCEEDED(limit->getArgumentValueInt(0, &count)) &&
259	1/2 ✓ Branch 0 taken 2 times. ✗ Branch 1 not taken.	2	count > 0) {
260		2	kernel.shape.maxRegisters = uint32_t(count);
261			}
262			}
263
264		7	int uniforms = 0;
265	2/2 ✓ Branch 1 taken 7 times. ✓ Branch 2 taken 7 times.	14	for (unsigned i = 0; i < reflected.getParameterCount(); ++i) {
266		7	slang::VariableLayoutReflection* parameter = reflected.getParameterByIndex(i);
267		7	const char* semantic = parameter->getSemanticName();
268	1/2 ✓ Branch 0 taken 7 times. ✗ Branch 1 not taken.	7	if (semantic == nullptr) {
269		7	++uniforms;
270		7	kernel.shape.paramBytes = uint32_t(parameter->getTypeLayout()->getStride());
271		✗	} else if (const std::string name = Lowercase(semantic);
272	0/4 ✗ Branch 0 not taken. ✗ Branch 1 not taken. ✗ Branch 2 not taken. ✗ Branch 3 not taken.	✗	name == "sv_dispatchthreadid" \|\| name == "sv_groupid") {
273		✗	kernel.shape.foldable = false;
274		✗	}
275			}
276	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	if (uniforms > 1) {
277		✗	throw Error("kernel {} of program \"{}\" takes {} uniform parameters; kernels take one "
278			"struct",
279	0/1 ✗ Branch 2 not taken.	✗	entry, label, uniforms);
280			}
281
282	1/2 ✗ Branch 1 not taken. ✓ Branch 2 taken 7 times.	7	for (SlangUInt i = 0; i < layout.getHashedStringCount(); ++i) {
283		✗	size_t length = 0;
284		✗	const char* text = layout.getHashedString(i, &length);
285		✗	kernel.strings.emplace_back(text, length);
286			}
287		7	}
288
289		✗	void AdaptCudaSource(Backend backend, std::string_view entry, std::string_view label,
290			std::string& code) {
291	0/5 ✗ Branch 0 not taken. ✗ Branch 1 not taken. ✗ Branch 3 not taken. ✗ Branch 4 not taken. ✗ Branch 6 not taken.	✗	static const std::regex kernels(R"((extern "C" __global__ void \w+$[^)]$\s\{))");
292	0/2 ✗ Branch 1 not taken. ✗ Branch 2 not taken.	✗	if (!std::regex_search(code, kernels)) {
293		✗	throw Error("Slang no longer declares kernel {} of program \"{}\" as the CUDA and HIP "
294			"backends expect; update AdaptCudaSource for this Slang version",
295	0/1 ✗ Branch 2 not taken.	✗	entry, label);
296			}
297		✗	code = std::regex_replace(code, kernels,
298		✗	"$1\n if (chrono_batch_render_outside_dispatch()) { return; }");
299	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (backend == Backend::Hip) {
300	0/5 ✗ Branch 0 not taken. ✗ Branch 1 not taken. ✗ Branch 3 not taken. ✗ Branch 4 not taken. ✗ Branch 6 not taken.	✗	static const std::regex globals(R"(extern "C" (__constant__ [^;]*;))");
301		✗	const bool reached = code.find("SLANG_globalParams") != std::string::npos;
302	0/4 ✗ Branch 0 not taken. ✗ Branch 1 not taken. ✗ Branch 3 not taken. ✗ Branch 4 not taken.	✗	if (reached && !std::regex_search(code, globals)) {
303		✗	throw Error("Slang no longer declares the globals of kernel {} of program \"{}\" as "
304			"the HIP backend expects; update AdaptCudaSource for this Slang version",
305	0/1 ✗ Branch 2 not taken.	✗	entry, label);
306			}
307		✗	code = std::regex_replace(code, globals, R"(extern "C" { $1 })");
308			}
309		✗	}
310
311		49	std::string SessionKey(Backend backend, Traversal traversal, Sampling sampling, bool debugInfo,
312			bool narrowWaves) {
313	1/1 ✓ Branch 1 taken 49 times.	49	return fmt::format("slang {} {}{}{}{}{} {} {} {}", spGetBuildTagString(), BackendName(backend),
314	1/1 ✓ Branch 1 taken 49 times.	49	traversal == Traversal::Hardware ? " with hardware traversal" : "",
315	1/2 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken.	49	sampling == Sampling::Hardware ? " with hardware textures" : "",
316	4/6 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 49 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 45 times. ✓ Branch 5 taken 4 times.	196	debugInfo ? " with debug info" : "", narrowWaves ? " with narrow waves" : "",
317	1/1 ✓ Branch 1 taken 49 times.	98	Environment("CHRONO_BATCH_RENDER_EXP_OPT"),
318	1/1 ✓ Branch 1 taken 49 times.	98	Environment("CHRONO_BATCH_RENDER_EXP_MINOPT"),
319	1/1 ✓ Branch 2 taken 49 times.	147	Environment("CHRONO_BATCH_RENDER_EXP_SKIPVAL"));
320			}
321
322		49	ComPtr<slang::ISession> CreateSession(slang::IGlobalSession& global, Backend backend,
323			Traversal traversal, Sampling sampling, bool debugInfo,
324			ModuleFiles& files, std::string_view label,
325			bool narrowWaves) {
326		49	slang::TargetDesc target{};
327		49	std::vector<slang::CompilerOptionEntry> options;
328	1/4 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken. ✗ Branch 2 not taken. ✗ Branch 3 not taken.	49	switch (backend) {
329		49	case Backend::Vulkan:
330		49	target.format = SLANG_SPIRV;
331	1/1 ✓ Branch 1 taken 49 times.	49	target.profile = global.findProfile("spirv_1_6");
332	1/1 ✓ Branch 1 taken 49 times.	49	options.push_back(Enable(slang::CompilerOptionName::GLSLForceScalarLayout));
333	1/1 ✓ Branch 1 taken 49 times.	49	options.push_back(Enable(slang::CompilerOptionName::VulkanUseEntryPointName));
334		49	break;
335		✗	case Backend::Cuda:
336		✗	case Backend::Hip:
337		✗	target.format = SLANG_CUDA_SOURCE;
338		✗	break;
339		✗	case Backend::Metal:
340		✗	case Backend::Cpu:
341	0/2 ✗ Branch 2 not taken. ✗ Branch 5 not taken.	✗	throw UnsupportedError("{} has no Slang target yet", BackendName(backend));
342			}
343	1/1 ✓ Branch 1 taken 49 times.	49	options.push_back(Enable(slang::CompilerOptionName::UseUpToDateBinaryModule));
344	2/3 ✓ Branch 1 taken 49 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 49 times.	49	if (const std::string level = Environment("CHRONO_BATCH_RENDER_EXP_OPT"); !level.empty()) {
345		✗	slang::CompilerOptionEntry entry = Enable(slang::CompilerOptionName::Optimization);
346	0/1 ✗ Branch 1 not taken.	✗	entry.value.intValue0 = std::stoi(level);
347	0/1 ✗ Branch 1 not taken.	✗	options.push_back(entry);
348		✗	}
349	2/3 ✓ Branch 1 taken 49 times. ✗ Branch 4 not taken. ✓ Branch 5 taken 49 times.	49	if (Environment("CHRONO_BATCH_RENDER_EXP_MINOPT") == "1") {
350	0/1 ✗ Branch 1 not taken.	✗	options.push_back(Enable(slang::CompilerOptionName::MinimumSlangOptimization));
351			}
352	2/3 ✓ Branch 1 taken 49 times. ✗ Branch 4 not taken. ✓ Branch 5 taken 49 times.	49	if (Environment("CHRONO_BATCH_RENDER_EXP_SKIPVAL") == "1") {
353	0/1 ✗ Branch 1 not taken.	✗	options.push_back(Enable(slang::CompilerOptionName::SkipSPIRVValidation));
354			}
355	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 49 times.	49	if (debugInfo) {
356		✗	slang::CompilerOptionEntry lines = Enable(slang::CompilerOptionName::DebugInformation);
357		✗	lines.value.intValue0 = SLANG_DEBUG_INFO_LEVEL_STANDARD;
358	0/1 ✗ Branch 1 not taken.	✗	options.push_back(lines);
359			}
360	1/2 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken.	49	target.compilerOptionEntries = options.data();
361	1/2 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken.	49	target.compilerOptionEntryCount = uint32_t(options.size());
362
363		294	const auto flag = [](bool on) { return on ? "1" : "0"; };
364		49	const slang::PreprocessorMacroDesc macros[] = {
365	1/2 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken.	49	{"CHRONO_BATCH_RENDER_HIP", flag(backend == Backend::Hip)},
366		98	{"CHRONO_BATCH_RENDER_VULKAN", flag(backend == Backend::Vulkan)},
367		98	{"CHRONO_BATCH_RENDER_HARDWARE_TRAVERSAL", flag(traversal == Traversal::Hardware)},
368			{"CHRONO_BATCH_RENDER_HARDWARE_TEXTURES",
369		98	flag(sampling == Sampling::Hardware && backend == Backend::Vulkan)},
370			{"CHRONO_BATCH_RENDER_OBJECT_TEXTURES",
371		98	flag(sampling == Sampling::Hardware && backend != Backend::Vulkan)},
372		98	{"CHRONO_BATCH_RENDER_NARROW_WAVES", flag(narrowWaves)},
373	6/10 ✗ Branch 0 not taken. ✓ Branch 1 taken 49 times. ✓ Branch 2 taken 49 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 45 times. ✓ Branch 5 taken 4 times. ✓ Branch 6 taken 49 times. ✗ Branch 7 not taken. ✓ Branch 8 taken 49 times. ✗ Branch 9 not taken.	192	};
374
375		49	slang::SessionDesc desc{};
376		49	desc.targets = &target;
377		49	desc.targetCount = 1;
378		49	desc.fileSystem = &files;
379		49	const char* root = ".";
380		49	desc.searchPaths = &root;
381		49	desc.searchPathCount = 1;
382		49	desc.preprocessorMacros = macros;
383		49	desc.preprocessorMacroCount = uint32_t(std::size(macros));
384	1/1 ✓ Branch 1 taken 49 times.	49	ComPtr<slang::ISession> session;
385	2/3 ✓ Branch 1 taken 49 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 49 times.	49	if (SLANG_FAILED(global.createSession(desc, session.writeRef()))) {
386	0/1 ✗ Branch 2 not taken.	✗	throw Error("Slang failed to start a session for program \"{}\"", label);
387			}
388		49	return session;
389		49	}
390
391		9	struct KernelParts {
392			std::vector<slang::IComponentType*> components;
393			ComPtr<slang::IEntryPoint> entryPoint;
394			};
395
396		9	KernelParts LoadParts(slang::ISession& session, std::span<const std::string> modules,
397			std::string_view entry, std::string_view label,
398			ComPtr<slang::IBlob>& diagnostics) {
399		9	KernelParts loaded;
400	2/2 ✓ Branch 0 taken 9 times. ✓ Branch 1 taken 8 times.	17	for (const std::string& name : modules) {
401	2/3 ✗ Branch 0 not taken. ✓ Branch 1 taken 9 times. ✓ Branch 3 taken 9 times.	9	slang::IModule* module = session.loadModule(name.c_str(), diagnostics.writeRef());
402	2/2 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 8 times.	9	if (module == nullptr) {
403	1/1 ✓ Branch 1 taken 1 time.	2	Fail(fmt::format("Compiling module {}", name), label, diagnostics);
404			}
405	1/1 ✓ Branch 1 taken 8 times.	8	loaded.components.push_back(module);
406	1/2 ✓ Branch 0 taken 8 times. ✗ Branch 1 not taken.	8	if (!loaded.entryPoint) {
407	2/2 ✓ Branch 1 taken 8 times. ✓ Branch 4 taken 8 times.	16	module->findEntryPointByName(std::string(entry).c_str(), loaded.entryPoint.writeRef());
408			}
409			}
410	2/2 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 7 times.	8	if (!loaded.entryPoint) {
411	1/1 ✓ Branch 2 taken 1 time.	1	throw Error("program \"{}\" has no compute entry point {}", label, entry);
412			}
413	1/1 ✓ Branch 1 taken 7 times.	7	loaded.components.push_back(loaded.entryPoint);
414
415		7	const std::string named =
416	2/2 ✓ Branch 1 taken 7 times. ✓ Branch 4 taken 7 times.	7	fmt::format("export static const uint kKernelHash = {};", SlangStringHash(entry));
417	3/3 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 6 times. ✓ Branch 3 taken 7 times.	8	slang::IModule* kernelHash = session.loadModuleFromSourceString(
418			"chrono_batch_render_kernel_hash", "chrono_batch_render_kernel_hash.slang", named.c_str(),
419			diagnostics.writeRef());
420	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	if (kernelHash == nullptr) {
421		✗	Fail("Naming the kernel", label, diagnostics);
422			}
423	1/1 ✓ Branch 1 taken 7 times.	7	loaded.components.push_back(kernelHash);
424		7	return loaded;
425		9	}
426
427		7	struct LinkedKernel {
428			ComPtr<slang::IComponentType> composite;
429			ComPtr<slang::IComponentType> linked;
430			ComPtr<slang::IBlob> code;
431			};
432
433		7	LinkedKernel Link(slang::ISession& session, std::span<slang::IComponentType* const> components,
434			std::string_view entry, std::string_view label,
435			ComPtr<slang::IBlob>& diagnostics) {
436	1/1 ✓ Branch 1 taken 7 times.	7	LinkedKernel kernel;
437	1/1 ✓ Branch 1 taken 7 times.	7	const std::string what = fmt::format("Compiling kernel {}", entry);
438	3/5 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 7 times. ✓ Branch 5 taken 7 times.	7	if (SLANG_FAILED(session.createCompositeComponentType(
439			components.data(), SlangInt(components.size()), kernel.composite.writeRef(),
440	1/2 ✓ Branch 0 taken 7 times. ✗ Branch 1 not taken.	7	diagnostics.writeRef())) \|\|
441	5/9 ✓ Branch 0 taken 7 times. ✗ Branch 1 not taken. ✗ Branch 2 not taken. ✓ Branch 3 taken 7 times. ✗ Branch 4 not taken. ✓ Branch 5 taken 7 times. ✓ Branch 7 taken 7 times. ✗ Branch 9 not taken. ✓ Branch 10 taken 7 times.	14	SLANG_FAILED(kernel.composite->link(kernel.linked.writeRef(), diagnostics.writeRef())) \|\|
442	3/5 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 7 times. ✓ Branch 5 taken 7 times.	7	SLANG_FAILED(kernel.linked->getEntryPointCode(0, 0, kernel.code.writeRef(),
443			diagnostics.writeRef()))) {
444		✗	Fail(what, label, diagnostics);
445			}
446	2/3 ✓ Branch 1 taken 7 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 7 times.	7	if (const std::string warnings = Text(diagnostics); !warnings.empty()) {
447	0/1 ✗ Branch 1 not taken.	✗	Log(LogLevel::Debug, "{} of program \"{}\": {}", what, label, warnings);
448		✗	}
449		7	return kernel;
450		7	}
451
452		52	std::optional<FieldType> ReflectedType(slang::TypeReflection* type) {
453		52	const slang::TypeReflection::Kind kind = type->getKind();
454	3/4 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 51 times. ✓ Branch 2 taken 1 time. ✗ Branch 3 not taken.	53	if (kind == slang::TypeReflection::Kind::Struct &&
455	1/2 ✓ Branch 1 taken 1 time. ✗ Branch 2 not taken.	1	std::string_view(type->getName()) == "SpectralValue") {
456		1	return FieldType::Spectral;
457			}
458		51	const SlangScalarType scalar = SlangScalarType(type->getScalarType());
459	2/2 ✓ Branch 0 taken 33 times. ✓ Branch 1 taken 18 times.	51	if (kind == slang::TypeReflection::Kind::Scalar) {
460	3/4 ✓ Branch 0 taken 8 times. ✓ Branch 1 taken 4 times. ✓ Branch 2 taken 21 times. ✗ Branch 3 not taken.	33	switch (scalar) {
461		8	case SLANG_SCALAR_TYPE_BOOL:
462		8	return FieldType::Bool;
463		4	case SLANG_SCALAR_TYPE_INT32:
464		4	return FieldType::Int;
465		21	case SLANG_SCALAR_TYPE_FLOAT32:
466		21	return FieldType::Float;
467		✗	default:
468		✗	return std::nullopt;
469			}
470			}
471	2/2 ✓ Branch 0 taken 17 times. ✓ Branch 1 taken 1 time.	18	if (kind == slang::TypeReflection::Kind::Vector && scalar == SLANG_SCALAR_TYPE_FLOAT32) {
472	1/4 ✗ Branch 1 not taken. ✓ Branch 2 taken 17 times. ✗ Branch 3 not taken. ✗ Branch 4 not taken.	17	switch (type->getElementCount()) {
473		✗	case 2:
474		✗	return FieldType::Float2;
475		17	case 3:
476		17	return FieldType::Float3;
477		✗	case 4:
478		✗	return FieldType::Float4;
479		✗	default:
480		✗	return std::nullopt;
481			}
482			}
483		1	return std::nullopt;
484			}
485
486		49	uint32_t FieldBytes(FieldType type) {
487	3/5 ✓ Branch 0 taken 15 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 1 time. ✓ Branch 3 taken 33 times. ✗ Branch 4 not taken.	49	switch (type) {
488			case FieldType::Float2:
489			return 8;
490		15	case FieldType::Float3:
491		15	return 12;
492		✗	case FieldType::Float4:
493		✗	return 16;
494		1	case FieldType::Spectral:
495		1	return 12;
496		33	default:
497		33	return 4;
498			}
499			}
500
501		67	bool IsShadeMethod(slang::DeclReflection* decl) {
502	2/2 ✓ Branch 1 taken 3 times. ✓ Branch 2 taken 64 times.	67	if (decl->getKind() != slang::DeclReflection::Kind::Generic) {
503			return false;
504			}
505		3	slang::DeclReflection* inner = decl->asGeneric()->getInnerDecl();
506	3/6 ✓ Branch 0 taken 3 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 3 times. ✗ Branch 3 not taken. ✗ Branch 4 not taken. ✓ Branch 5 taken 3 times.	9	return inner != nullptr && inner->getKind() == slang::DeclReflection::Kind::Func &&
507	1/2 ✗ Branch 1 not taken. ✓ Branch 2 taken 3 times.	3	std::string_view(inner->getName()) == "Shade";
508			}
509
510		54	bool IsMaterialMethod(slang::DeclReflection* decl) {
511	2/2 ✓ Branch 1 taken 21 times. ✓ Branch 2 taken 33 times.	54	if (decl->getKind() != slang::DeclReflection::Kind::Func) {
512			return false;
513			}
514		21	slang::FunctionReflection* function = decl->asFunction();
515	3/4 ✓ Branch 0 taken 21 times. ✗ Branch 1 not taken. ✓ Branch 3 taken 17 times. ✓ Branch 4 taken 4 times.	38	return function != nullptr && std::string_view(function->getName()) == "Material" &&
516	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 17 times.	17	function->getParameterCount() == 2;
517			}
518
519		50	ReflectedField Field(slang::IGlobalSession& global, slang::VariableReflection* variable,
520			std::string_view node, std::string_view label) {
521		50	const std::optional<FieldType> type = ReflectedType(variable->getType());
522	2/2 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 49 times.	50	if (!type) {
523		1	throw Error("{}'s input {} in {} has type {}, but node inputs are bool, int, "
524			"float, float2, float3 or float4",
525	4/4 ✓ Branch 2 taken 1 time. ✓ Branch 5 taken 1 time. ✓ Branch 8 taken 1 time. ✓ Branch 11 taken 1 time.	1	node, variable->getName(), label, variable->getType()->getName());
526			}
527	1/1 ✓ Branch 3 taken 49 times.	49	ReflectedField field{.name = variable->getName(), .type = *type};
528	1/1 ✓ Branch 1 taken 49 times.	49	field.fallback.assign(FieldBytes(*type), std::byte{0});
529	1/1 ✓ Branch 1 taken 49 times.	49	ComPtr<ISlangBlob> fallback;
530	1/1 ✓ Branch 1 taken 49 times.	49	const SlangResult read = variable->getDefaultValueBlob(fallback.writeRef());
531	3/4 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 2 times. ✓ Branch 3 taken 47 times.	49	const bool authored = SLANG_SUCCEEDED(read) && fallback;
532	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 49 times.	49	if (read == SLANG_E_NOT_AVAILABLE) {
533	0/1 ✗ Branch 1 not taken.	✗	Log(LogLevel::Warning,
534			"{}'s input {} in {} has a default Chrono Batch Render can't read, so it starts at "
535			"zero; write it "
536			"as a literal, such as {{ 1, 1, float3(0.5) }} for an RGB SpectralValue",
537	0/3 ✗ Branch 1 not taken. ✗ Branch 3 not taken. ✗ Branch 4 not taken.	✗	node, variable->getName(), label);
538			}
539	6/7 ✓ Branch 0 taken 47 times. ✓ Branch 1 taken 2 times. ✓ Branch 2 taken 1 time. ✓ Branch 3 taken 46 times. ✓ Branch 5 taken 1 time. ✓ Branch 7 taken 1 time. ✗ Branch 8 not taken.	49	if (authored && *type == FieldType::Spectral && fallback->getBufferSize() >= 20) {
540		1	struct Packed {
541			uint32_t kind;
542			float scale;
543			float data[3];
544			} packed;
545
546	1/1 ✓ Branch 1 taken 1 time.	1	std::memcpy(&packed, fallback->getBufferPointer(), sizeof(packed));
547		1	float rgb[3];
548	2/2 ✓ Branch 0 taken 3 times. ✓ Branch 1 taken 1 time.	4	for (int i = 0; i < 3; ++i) {
549	1/2 ✓ Branch 0 taken 3 times. ✗ Branch 1 not taken.	3	rgb[i] = packed.kind == 1 ? packed.data[i] * packed.scale : packed.scale;
550			}
551		1	std::memcpy(field.fallback.data(), rgb, sizeof(rgb));
552	4/5 ✓ Branch 0 taken 46 times. ✓ Branch 1 taken 2 times. ✓ Branch 3 taken 46 times. ✓ Branch 5 taken 46 times. ✗ Branch 6 not taken.	48	} else if (authored && fallback->getBufferSize() >= field.fallback.size()) {
553	1/1 ✓ Branch 1 taken 46 times.	46	std::memcpy(field.fallback.data(), fallback->getBufferPointer(), field.fallback.size());
554			}
555	3/4 ✓ Branch 0 taken 49 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 1 time. ✓ Branch 3 taken 48 times.	49	field.color = field.color \|\| *type == FieldType::Spectral;
556		49	SlangSession* session = &global;
557	1/1 ✓ Branch 1 taken 49 times.	49	field.color = variable->findUserAttributeByName(session, "Color") != nullptr;
558	3/3 ✓ Branch 1 taken 49 times. ✓ Branch 3 taken 10 times. ✓ Branch 4 taken 39 times.	49	if (slang::Attribute* input = variable->findUserAttributeByName(session, "Input")) {
559	1/1 ✓ Branch 1 taken 10 times.	10	input->getArgumentValueFloat(0, &field.minimum);
560	1/1 ✓ Branch 1 taken 10 times.	10	input->getArgumentValueFloat(1, &field.maximum);
561			}
562	2/3 ✓ Branch 1 taken 49 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 49 times.	49	if (slang::Attribute* named = variable->findUserAttributeByName(session, "DisplayName")) {
563		✗	size_t size = 0;
564	0/3 ✗ Branch 1 not taken. ✗ Branch 3 not taken. ✗ Branch 4 not taken.	✗	if (const char* text = named->getArgumentValueString(0, &size)) {
565	0/1 ✗ Branch 1 not taken.	✗	field.displayName.assign(text, size);
566			}
567			}
568	2/2 ✓ Branch 0 taken 47 times. ✓ Branch 1 taken 2 times.	49	return field;
569		49	}
570
571		6	std::optional<ReflectedNode> ReflectPattern(slang::IGlobalSession& global,
572			slang::DeclReflection* decl, std::string_view label) {
573		6	slang::FunctionReflection* function = decl->asFunction();
574	3/4 ✓ Branch 0 taken 6 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 4 times. ✓ Branch 3 taken 2 times.	12	if (function == nullptr \|\| function->findUserAttributeByName(&global, "Pattern") == nullptr) {
575		4	return std::nullopt;
576			}
577	1/1 ✓ Branch 3 taken 2 times.	2	ReflectedNode node{.name = function->getName(), .role = NodeRole::Pattern};
578	2/2 ✓ Branch 1 taken 2 times. ✓ Branch 4 taken 2 times.	2	const std::optional<FieldType> result = ReflectedType(function->getReturnType());
579	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	2	if (!result) {
580		✗	throw Error("pattern {} in {} returns {}, but patterns return bool, int, float, float2, "
581			"float3 or float4",
582	0/3 ✗ Branch 2 not taken. ✗ Branch 5 not taken. ✗ Branch 8 not taken.	✗	node.name, label, function->getReturnType()->getName());
583			}
584		2	node.result = *result;
585	3/3 ✓ Branch 1 taken 6 times. ✓ Branch 3 taken 4 times. ✓ Branch 4 taken 2 times.	6	for (unsigned i = 0; i < function->getParameterCount(); ++i) {
586	1/1 ✓ Branch 1 taken 4 times.	4	slang::VariableReflection* parameter = function->getParameterByIndex(i);
587	4/4 ✓ Branch 1 taken 4 times. ✓ Branch 4 taken 4 times. ✓ Branch 6 taken 1 time. ✓ Branch 7 taken 3 times.	4	const std::string_view type = parameter->getType()->getName();
588	2/2 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 3 times.	4	if (type == "ShadeContext") {
589	1/1 ✓ Branch 1 taken 1 time.	1	node.parameters.push_back({.source = ParameterSource::Context});
590	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 3 times.	3	} else if (type == "SurfacePoint") {
591	0/1 ✗ Branch 1 not taken.	✗	node.parameters.push_back({.source = ParameterSource::Point});
592			} else {
593	1/1 ✓ Branch 1 taken 3 times.	3	node.parameters.push_back(
594	1/1 ✓ Branch 1 taken 3 times.	3	{.source = ParameterSource::Input, .input = uint32_t(node.fields.size())});
595	1/1 ✓ Branch 1 taken 3 times.	6	node.fields.push_back(Field(global, parameter, node.name, label));
596			}
597			}
598		2	return node;
599		2	}
600
601		24	std::optional<ReflectedNode> ReflectNode(slang::IGlobalSession& global, slang::DeclReflection* decl,
602			std::string_view label) {
603	1/2 ✓ Branch 1 taken 24 times. ✗ Branch 2 not taken.	24	if (decl->getType() != nullptr &&
604	2/2 ✓ Branch 1 taken 2 times. ✓ Branch 2 taken 22 times.	24	decl->getType()->findUserAttributeByName("ClosurePrimitive") != nullptr) {
605		2	return ReflectedNode{.name = decl->getName(), .role = NodeRole::Primitive};
606			}
607		22	bool material = false;
608		22	bool surface = false;
609	2/2 ✓ Branch 1 taken 69 times. ✓ Branch 2 taken 22 times.	91	for (slang::DeclReflection* child : decl->getChildren()) {
610	4/4 ✓ Branch 0 taken 54 times. ✓ Branch 1 taken 15 times. ✓ Branch 3 taken 17 times. ✓ Branch 4 taken 37 times.	69	material = material \|\| IsMaterialMethod(child);
611	4/4 ✓ Branch 0 taken 67 times. ✓ Branch 1 taken 2 times. ✓ Branch 3 taken 3 times. ✓ Branch 4 taken 64 times.	72	surface = surface \|\| IsShadeMethod(child);
612			}
613	2/2 ✓ Branch 0 taken 2 times. ✓ Branch 1 taken 20 times.	22	if (!material && !surface) {
614		2	return std::nullopt;
615			}
616	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 20 times.	20	if (material && surface) {
617		✗	throw Error("{} in {} has both Material and Shade; a node fills SurfaceInputs or builds a "
618			"closure, not both",
619	0/2 ✗ Branch 2 not taken. ✗ Branch 5 not taken.	✗	decl->getName(), label);
620			}
621		20	ReflectedNode node{.name = decl->getName(),
622	3/3 ✓ Branch 1 taken 17 times. ✓ Branch 2 taken 3 times. ✓ Branch 4 taken 20 times.	37	.role = surface ? NodeRole::Surface : NodeRole::Material};
623	3/3 ✓ Branch 1 taken 20 times. ✓ Branch 3 taken 63 times. ✓ Branch 4 taken 19 times.	82	for (slang::DeclReflection* child : decl->getChildren()) {
624	3/3 ✓ Branch 1 taken 63 times. ✓ Branch 3 taken 35 times. ✓ Branch 4 taken 28 times.	63	if (child->getKind() != slang::DeclReflection::Kind::Variable) {
625		35	continue;
626			}
627	1/1 ✓ Branch 1 taken 28 times.	28	slang::VariableReflection* variable = child->asVariable();
628	2/3 ✓ Branch 1 taken 28 times. ✓ Branch 3 taken 28 times. ✗ Branch 4 not taken.	28	if (variable->findModifier(slang::Modifier::Static) != nullptr \|\|
629	1/2 ✓ Branch 0 taken 28 times. ✗ Branch 1 not taken.	28	variable->findModifier(slang::Modifier::Const) != nullptr) {
630		✗	continue;
631			}
632	1/1 ✓ Branch 1 taken 27 times.	55	node.fields.push_back(Field(global, variable, node.name, label));
633			}
634		19	return node;
635		20	}
636
637			}
638
639			namespace {
640
641		37	std::string ShortName(slang::TypeReflection* type) {
642	1/1 ✓ Branch 1 taken 37 times.	37	ComPtr<ISlangBlob> full;
643		37	std::string name =
644	2/3 ✓ Branch 1 taken 37 times. ✓ Branch 3 taken 37 times. ✗ Branch 4 not taken.	74	SLANG_SUCCEEDED(type->getFullName(full.writeRef())) && full
645	2/3 ✓ Branch 0 taken 37 times. ✗ Branch 1 not taken. ✓ Branch 3 taken 37 times.	74	? std::string(static_cast<const char*>(full->getBufferPointer()), full->getBufferSize())
646	3/7 ✓ Branch 1 taken 37 times. ✓ Branch 4 taken 37 times. ✓ Branch 7 taken 37 times. ✗ Branch 10 not taken. ✗ Branch 13 not taken. ✗ Branch 15 not taken. ✗ Branch 16 not taken.	37	: std::string(type->getName());
647	2/4 ✓ Branch 0 taken 37 times. ✗ Branch 1 not taken. ✗ Branch 2 not taken. ✓ Branch 3 taken 37 times.	37	while (!name.empty() && name.back() == '\0') {
648		✗	name.pop_back();
649			}
650		37	const size_t open = name.find('<');
651		37	const size_t scope = name.rfind('.', open);
652	2/4 ✓ Branch 0 taken 37 times. ✗ Branch 1 not taken. ✓ Branch 3 taken 37 times. ✗ Branch 6 not taken.	37	return scope == std::string::npos ? name : name.substr(scope + 1);
653		74	}
654
655		39	std::optional<ReflectedSensorPipelineNode> ReflectSensorPipelineNode(slang::IGlobalSession& global,
656			slang::ProgramLayout& layout,
657			slang::DeclReflection* decl,
658			std::string_view label) {
659		39	slang::TypeReflection* type = decl->getType();
660		39	slang::TypeReflection* point = layout.findTypeByName("IPointNode");
661		39	slang::TypeReflection* stencil = layout.findTypeByName("IStencilNode");
662	3/4 ✓ Branch 0 taken 16 times. ✓ Branch 1 taken 23 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 16 times.	39	if (type == nullptr \|\| point == nullptr \|\| stencil == nullptr) {
663		23	return std::nullopt;
664			}
665		16	const bool isPoint = layout.isSubType(type, point);
666		16	const bool isStencil = layout.isSubType(type, stencil);
667	2/2 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 15 times.	16	if (!isPoint && !isStencil) {
668		1	return std::nullopt;
669			}
670		15	ReflectedSensorPipelineNode node{.name = decl->getName(),
671	2/2 ✓ Branch 0 taken 10 times. ✓ Branch 1 taken 5 times.	15	.kind = isStencil ? SensorPipelineNodeKind::Stencil
672	3/4 ✓ Branch 1 taken 10 times. ✓ Branch 2 taken 5 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 15 times.	25	: SensorPipelineNodeKind::Point};
673	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 15 times.	15	if (isPoint && isStencil) {
674	0/1 ✗ Branch 2 not taken.	✗	throw Error("{} in {} is both a point and a stencil node; pick one", node.name, label);
675			}
676	3/3 ✓ Branch 1 taken 15 times. ✓ Branch 3 taken 4 times. ✓ Branch 4 taken 11 times.	15	if (slang::Attribute* named = type->findUserAttributeByName("Node")) {
677		4	size_t size = 0;
678	2/3 ✓ Branch 1 taken 4 times. ✓ Branch 3 taken 4 times. ✗ Branch 4 not taken.	4	if (const char* text = named->getArgumentValueString(0, &size)) {
679	1/1 ✓ Branch 1 taken 4 times.	4	node.displayName.assign(text, size);
680			}
681			}
682		15	slang::FunctionReflection* apply = nullptr;
683	3/3 ✓ Branch 1 taken 15 times. ✓ Branch 3 taken 59 times. ✓ Branch 4 taken 15 times.	74	for (slang::DeclReflection* child : decl->getChildren()) {
684	5/5 ✓ Branch 1 taken 59 times. ✓ Branch 3 taken 21 times. ✓ Branch 4 taken 38 times. ✓ Branch 5 taken 14 times. ✓ Branch 6 taken 7 times.	80	if (child->getKind() == slang::DeclReflection::Kind::Func &&
685	3/3 ✓ Branch 1 taken 21 times. ✓ Branch 3 taken 14 times. ✓ Branch 4 taken 7 times.	21	std::string_view(child->getName()) == "Apply") {
686	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 14 times.	14	if (apply != nullptr) {
687		✗	throw Error("{} in {} has more than one Apply; a node has exactly one", node.name,
688	0/1 ✗ Branch 2 not taken.	✗	label);
689			}
690	1/1 ✓ Branch 1 taken 14 times.	14	apply = child->asFunction();
691		40	continue;
692			}
693	3/3 ✓ Branch 1 taken 45 times. ✓ Branch 3 taken 22 times. ✓ Branch 4 taken 23 times.	45	if (child->getKind() != slang::DeclReflection::Kind::Variable) {
694		22	continue;
695			}
696	1/1 ✓ Branch 1 taken 23 times.	23	slang::VariableReflection* variable = child->asVariable();
697	3/3 ✓ Branch 1 taken 23 times. ✓ Branch 3 taken 19 times. ✓ Branch 4 taken 4 times.	23	const bool constant = variable->findModifier(slang::Modifier::Static) != nullptr \|\|
698	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 19 times.	19	variable->findModifier(slang::Modifier::Const) != nullptr;
699	2/3 ✓ Branch 1 taken 4 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 4 times.	4	if (constant && std::string_view(variable->getName()) == "kRadius") {
700	1/1 ✓ Branch 1 taken 4 times.	4	ComPtr<ISlangBlob> value;
701	4/7 ✓ Branch 1 taken 4 times. ✓ Branch 3 taken 4 times. ✗ Branch 4 not taken. ✓ Branch 5 taken 4 times. ✗ Branch 6 not taken. ✓ Branch 7 taken 4 times. ✗ Branch 8 not taken.	8	if (SLANG_SUCCEEDED(variable->getDefaultValueBlob(value.writeRef())) && value &&
702	1/1 ✓ Branch 1 taken 4 times.	4	value->getBufferSize() >= sizeof(int32_t)) {
703		4	int32_t radius = 0;
704	2/3 ✓ Branch 1 taken 4 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 4 times.	4	std::memcpy(&radius, value->getBufferPointer(), sizeof(radius));
705	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 4 times.	4	node.radius = uint32_t(std::max(radius, 0));
706			}
707	1/2 ✓ Branch 0 taken 4 times. ✗ Branch 1 not taken.	4	continue;
708		4	}
709	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 19 times.	19	if (constant) {
710		✗	continue;
711			}
712	1/1 ✓ Branch 1 taken 19 times.	19	ReflectedField field = Field(global, variable, node.name, label);
713	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 19 times.	19	if (field.type == FieldType::Spectral) {
714		✗	throw Error("{}'s input {} in {} is a SpectralValue, which sensor pipeline nodes "
715			"don't take",
716	0/1 ✗ Branch 2 not taken.	✗	node.name, field.name, label);
717			}
718	1/1 ✓ Branch 1 taken 19 times.	19	node.fields.push_back(std::move(field));
719		19	}
720	2/2 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 14 times.	15	if (apply == nullptr) {
721		1	throw Error("{} in {} is a pipeline node, so it needs an Apply method that takes its "
722			"signals and a PointContext and returns a signal",
723	1/1 ✓ Branch 2 taken 1 time.	1	node.name, label);
724			}
725	4/4 ✓ Branch 0 taken 5 times. ✓ Branch 1 taken 9 times. ✓ Branch 2 taken 1 time. ✓ Branch 3 taken 4 times.	14	if (isStencil && node.radius == 0) {
726		1	throw Error("{} in {} is a stencil node, so it needs a static const int kRadius above "
727			"zero",
728	1/1 ✓ Branch 2 taken 1 time.	1	node.name, label);
729			}
730	2/2 ✓ Branch 1 taken 13 times. ✓ Branch 4 taken 13 times.	13	node.result = ShortName(apply->getReturnType());
731	3/3 ✓ Branch 1 taken 37 times. ✓ Branch 3 taken 24 times. ✓ Branch 4 taken 13 times.	37	for (unsigned index = 0; index < apply->getParameterCount(); ++index) {
732	1/1 ✓ Branch 1 taken 24 times.	24	slang::VariableReflection* parameter = apply->getParameterByIndex(index);
733	2/2 ✓ Branch 1 taken 24 times. ✓ Branch 4 taken 24 times.	24	std::string parameterType = ShortName(parameter->getType());
734	2/2 ✓ Branch 0 taken 11 times. ✓ Branch 1 taken 13 times.	24	if (parameterType == "PointContext") {
735	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 11 times.	11	if (node.context != kNoContext) {
736	0/1 ✗ Branch 2 not taken.	✗	throw Error("{}'s Apply in {} takes more than one PointContext", node.name, label);
737			}
738		11	node.context = index;
739		11	continue;
740			}
741	2/2 ✓ Branch 1 taken 4 times. ✓ Branch 2 taken 9 times.	13	if (parameterType.starts_with("Neighborhood<")) {
742	2/4 ✓ Branch 0 taken 4 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 4 times. ✗ Branch 3 not taken.	4	if (!isStencil \|\| !node.inputs.empty()) {
743		✗	throw Error("{}'s Apply in {} takes a Neighborhood, which only the first signal "
744			"of a stencil node may be",
745	0/1 ✗ Branch 2 not taken.	✗	node.name, label);
746			}
747	1/1 ✓ Branch 1 taken 4 times.	4	parameterType = parameterType.substr(13, parameterType.size() - 14);
748		4	const size_t scope = parameterType.rfind('.');
749	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 4 times.	4	if (scope != std::string::npos) {
750	0/1 ✗ Branch 1 not taken.	✗	parameterType = parameterType.substr(scope + 1);
751			}
752	1/4 ✗ Branch 0 not taken. ✓ Branch 1 taken 9 times. ✗ Branch 2 not taken. ✗ Branch 3 not taken.	9	} else if (isStencil && node.inputs.empty()) {
753		✗	throw Error("{} in {} is a stencil node, so its Apply takes a Neighborhood of its "
754			"first signal",
755	0/1 ✗ Branch 2 not taken.	✗	node.name, label);
756			}
757	1/1 ✓ Branch 1 taken 13 times.	26	node.inputs.push_back({parameter->getName(), std::move(parameterType)});
758		24	}
759		13	return node;
760		41	}
761
762			}
763
764		40	ReflectedNodes ReflectNodes(Backend backend, const KernelLibrary* library,
765			const SourceFiles& sources, std::string_view module,
766			std::string_view label, const DiskCache* modules) {
767		40	slang::IGlobalSession& global = GlobalSession();
768		40	ModuleFiles files(library, sources, modules,
769		40	SessionKey(backend, Traversal::Software, Sampling::Software, false, false));
770		40	ComPtr<slang::ISession> session = CreateSession(global, backend, Traversal::Software,
771	1/1 ✓ Branch 1 taken 40 times.	40	Sampling::Software, false, files, label, false);
772	1/1 ✓ Branch 1 taken 40 times.	40	ComPtr<slang::IBlob> diagnostics;
773	1/1 ✓ Branch 1 taken 40 times.	40	slang::IModule* loaded =
774	2/2 ✓ Branch 1 taken 40 times. ✓ Branch 4 taken 40 times.	40	session->loadModule(std::string(module).c_str(), diagnostics.writeRef());
775	2/2 ✓ Branch 0 taken 7 times. ✓ Branch 1 taken 33 times.	40	if (loaded == nullptr) {
776	1/1 ✓ Branch 1 taken 7 times.	14	Fail(fmt::format("Loading {}", label), label, diagnostics);
777			}
778	3/3 ✓ Branch 0 taken 1 time. ✓ Branch 1 taken 32 times. ✓ Branch 3 taken 33 times.	34	slang::ProgramLayout* layout = loaded->getLayout(0, diagnostics.writeRef());
779	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 33 times.	33	if (layout == nullptr) {
780	2/3 ✗ Branch 1 not taken. ✓ Branch 4 taken 7 times. ✓ Branch 5 taken 3 times.	10	Fail(fmt::format("Reflecting {}", label), label, diagnostics);
781			}
782	1/1 ✓ Branch 1 taken 33 times.	33	files.StoreBinaries(*session);
783		33	ReflectedNodes nodes;
784	4/4 ✓ Branch 1 taken 33 times. ✓ Branch 4 taken 33 times. ✓ Branch 6 taken 120 times. ✓ Branch 7 taken 30 times.	150	for (slang::DeclReflection* decl : loaded->getModuleReflection()->getChildren()) {
785	3/3 ✓ Branch 1 taken 120 times. ✓ Branch 3 taken 39 times. ✓ Branch 4 taken 81 times.	120	if (decl->getKind() == slang::DeclReflection::Kind::Struct) {
786		76	if (std::optional<ReflectedSensorPipelineNode> node =
787	3/3 ✓ Branch 1 taken 37 times. ✓ Branch 3 taken 13 times. ✓ Branch 4 taken 24 times.	39	ReflectSensorPipelineNode(global, *layout, decl, label)) {
788	1/1 ✓ Branch 1 taken 13 times.	36	nodes.pipeline.push_back(std::move(*node));
789	3/3 ✓ Branch 1 taken 23 times. ✓ Branch 3 taken 21 times. ✓ Branch 4 taken 2 times.	24	} else if (std::optional<ReflectedNode> material = ReflectNode(global, decl, label)) {
790	1/1 ✓ Branch 1 taken 21 times.	23	nodes.material.push_back(std::move(*material));
791		60	}
792	3/3 ✓ Branch 1 taken 81 times. ✓ Branch 3 taken 6 times. ✓ Branch 4 taken 75 times.	81	} else if (decl->getKind() == slang::DeclReflection::Kind::Func) {
793	3/3 ✓ Branch 1 taken 6 times. ✓ Branch 3 taken 2 times. ✓ Branch 4 taken 4 times.	6	if (std::optional<ReflectedNode> pattern = ReflectPattern(global, decl, label)) {
794	1/1 ✓ Branch 1 taken 2 times.	6	nodes.material.push_back(std::move(*pattern));
795		6	}
796			}
797			}
798	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 30 times.	60	return nodes;
799		83	}
800
801			std::vector<ReflectedSensorPipelineNode>
802		✗	ReflectSensorPipelineNodes(Backend backend, const KernelLibrary* library,
803			const SourceFiles& sources, std::string_view module,
804			std::string_view label, const DiskCache* modules) {
805		✗	return ReflectNodes(backend, library, sources, module, label, modules).pipeline;
806			}
807
808		2	std::vector<ReflectedNode> ReflectMaterialNodes(Backend backend, const KernelLibrary* library,
809			const SourceFiles& sources, std::string_view module,
810			std::string_view label, const DiskCache* modules) {
811		2	return ReflectNodes(backend, library, sources, module, label, modules).material;
812			}
813
814			namespace {
815
816		9	KernelCode CompileWithSlang(Backend backend, Traversal traversal, Sampling sampling, bool debugInfo,
817			const KernelLibrary* library, const SourceFiles& sources,
818			std::span<const std::string> modules, std::string_view entry,
819			std::string_view label, bool narrowWaves, const DiskCache& binaries) {
820		9	slang::IGlobalSession& global = GlobalSession();
821		9	ModuleFiles files(library, sources, &binaries,
822		9	SessionKey(backend, traversal, sampling, debugInfo, narrowWaves));
823		9	ComPtr<slang::ISession> session =
824	1/1 ✓ Branch 1 taken 9 times.	9	CreateSession(global, backend, traversal, sampling, debugInfo, files, label, narrowWaves);
825
826	1/1 ✓ Branch 1 taken 7 times.	9	ComPtr<slang::IBlob> diagnostics;
827	1/1 ✓ Branch 1 taken 7 times.	9	const KernelParts parts = LoadParts(*session, modules, entry, label, diagnostics);
828	1/1 ✓ Branch 1 taken 7 times.	7	const LinkedKernel linked = Link(*session, parts.components, entry, label, diagnostics);
829
830	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	slang::EntryPointReflection* reflected =
831	3/4 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times. ✓ Branch 3 taken 7 times. ✓ Branch 6 taken 7 times.	7	linked.linked->getLayout(0, diagnostics.writeRef())->getEntryPointByIndex(0);
832	2/3 ✓ Branch 1 taken 7 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 7 times.	7	if (reflected->getStage() != SLANG_STAGE_COMPUTE) {
833	0/1 ✗ Branch 2 not taken.	✗	throw Error("entry point {} of program \"{}\" isn't a compute shader", entry, label);
834			}
835	1/1 ✓ Branch 1 taken 7 times.	7	KernelCode kernel{.code = std::string(static_cast<const char*>(linked.code->getBufferPointer()),
836	4/5 ✓ Branch 1 taken 7 times. ✓ Branch 4 taken 7 times. ✓ Branch 7 taken 7 times. ✗ Branch 9 not taken. ✓ Branch 10 taken 7 times.	7	linked.code->getBufferSize())};
837	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	if (backend == Backend::Cuda \|\| backend == Backend::Hip) {
838	0/1 ✗ Branch 1 not taken.	✗	AdaptCudaSource(backend, entry, label, kernel.code);
839			}
840	3/4 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times. ✓ Branch 3 taken 7 times. ✓ Branch 6 taken 7 times.	7	Reflect(reflected, linked.linked->getLayout(0, diagnostics.writeRef()), entry, label, kernel);
841	1/1 ✓ Branch 1 taken 7 times.	7	files.RecordDependencies(*session);
842	1/1 ✓ Branch 1 taken 7 times.	7	files.StoreBinaries(*session);
843	1/1 ✓ Branch 1 taken 7 times.	7	kernel.files = files.Read();
844		7	return kernel;
845		18	}
846
847			}
848
849			namespace {
850
851		586	void AppendField(std::string& entry, std::string_view text) {
852		1172	entry += fmt::format("{} ", text.size());
853		586	entry += text;
854		586	entry += '\n';
855		586	}
856
857		72320	template <class T> bool ReadNumber(std::string_view entry, size_t& at, char after, T& value) {
858		72320	const char* end = entry.data() + entry.size();
859		72320	const auto [next, error] = std::from_chars(entry.data() + at, end, value);
860	6/12 bool chrono::batch_render::device::(anonymous namespace)::ReadNumber<int>(std::basic_string_view<char, std::char_traits<char> >, unsigned long&, char, int&): ✓ Branch 0 taken 1117 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 1117 times. ✗ Branch 3 not taken. bool chrono::batch_render::device::(anonymous namespace)::ReadNumber<unsigned int>(std::basic_string_view<char, std::char_traits<char> >, unsigned long&, char, unsigned int&): ✓ Branch 0 taken 5585 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 5585 times. ✗ Branch 3 not taken. bool chrono::batch_render::device::(anonymous namespace)::ReadNumber<unsigned long>(std::basic_string_view<char, std::char_traits<char> >, unsigned long&, char, unsigned long&): ✓ Branch 0 taken 65618 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 65618 times. ✗ Branch 3 not taken.	72320	if (error != std::errc() \|\| next == end \|\| *next != after) {
861			return false;
862			}
863		72320	at = size_t(next - entry.data()) + 1;
864		72320	return true;
865			}
866
867		63384	std::optional<std::string_view> ReadField(std::string_view entry, size_t& at) {
868		63384	size_t size = 0;
869	2/4 ✓ Branch 1 taken 63384 times. ✗ Branch 2 not taken. ✗ Branch 3 not taken. ✓ Branch 4 taken 63384 times.	63384	if (!ReadNumber(entry, at, ' ', size) \|\| at + size >= entry.size()) {
870		✗	return std::nullopt;
871			}
872		63384	const std::string_view field = entry.substr(at, size);
873		63384	at += size + 1;
874		63384	return field;
875			}
876
877		1117	bool Unchanged(const KernelLibrary* library, const SourceFiles& sources,
878			std::span<const std::pair<std::string, std::string>> files) {
879		1117	return std::ranges::all_of(files, [&](const auto& file) {
880		31692	const std::optional<std::string_view> source = FindFile(library, sources, file.first);
881	2/4 ✓ Branch 0 taken 31692 times. ✗ Branch 1 not taken. ✓ Branch 3 taken 31692 times. ✗ Branch 4 not taken.	63384	return source && Digest(*source) == file.second;
882		1117	});
883			}
884
885			}
886
887		7	std::string KernelCode::Encode() const {
888		7	std::string entry = fmt::format(
889		7	"{} {} {} {} {} {} {} {}\n", shape.groupSize.x, shape.groupSize.y, shape.groupSize.z,
890	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	shape.paramBytes, shape.maxRegisters, shape.foldable ? 1 : 0, strings.size(), files.size());
891	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 7 times.	7	for (const std::string& text : strings) {
892	0/1 ✗ Branch 1 not taken.	✗	AppendField(entry, text);
893			}
894	3/3 ✓ Branch 1 taken 293 times. ✓ Branch 3 taken 293 times. ✓ Branch 4 taken 7 times.	300	for (const auto& [path, digest] : files) {
895	1/1 ✓ Branch 1 taken 293 times.	293	AppendField(entry, path);
896	1/1 ✓ Branch 1 taken 293 times.	293	AppendField(entry, digest);
897			}
898	1/1 ✓ Branch 1 taken 7 times.	7	return entry + code;
899		7	}
900
901		1117	std::optional<KernelCode> KernelCode::Decode(std::string_view entry) {
902		1117	KernelCode kernel;
903		1117	size_t strings = 0;
904		1117	size_t files = 0;
905		1117	int foldable = 0;
906		1117	KernelShape& shape = kernel.shape;
907		1117	size_t at = 0;
908	1/2 ✓ Branch 1 taken 1117 times. ✗ Branch 2 not taken.	2234	if (!ReadNumber(entry, at, ' ', shape.groupSize.x) \|\|
909	1/2 ✓ Branch 1 taken 1117 times. ✗ Branch 2 not taken.	2234	!ReadNumber(entry, at, ' ', shape.groupSize.y) \|\|
910	1/2 ✓ Branch 1 taken 1117 times. ✗ Branch 2 not taken.	2234	!ReadNumber(entry, at, ' ', shape.groupSize.z) \|\|
911	1/2 ✓ Branch 1 taken 1117 times. ✗ Branch 2 not taken.	2234	!ReadNumber(entry, at, ' ', shape.paramBytes) \|\|
912	2/4 ✓ Branch 2 taken 1117 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 1117 times. ✗ Branch 5 not taken.	3351	!ReadNumber(entry, at, ' ', shape.maxRegisters) \|\| !ReadNumber(entry, at, ' ', foldable) \|\|
913	2/4 ✓ Branch 0 taken 1117 times. ✗ Branch 1 not taken. ✗ Branch 4 not taken. ✓ Branch 5 taken 1117 times.	3351	!ReadNumber(entry, at, ' ', strings) \|\| !ReadNumber(entry, at, '\n', files)) {
914		✗	return std::nullopt;
915			}
916		1117	shape.foldable = foldable != 0;
917
918	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1117 times.	1117	for (size_t i = 0; i < strings; ++i) {
919	0/1 ✗ Branch 1 not taken.	✗	const std::optional<std::string_view> text = ReadField(entry, at);
920	0/2 ✗ Branch 0 not taken. ✗ Branch 1 not taken.	✗	if (!text) {
921		✗	return std::nullopt;
922			}
923	0/1 ✗ Branch 1 not taken.	✗	kernel.strings.emplace_back(*text);
924			}
925	2/2 ✓ Branch 0 taken 31692 times. ✓ Branch 1 taken 1117 times.	32809	for (size_t i = 0; i < files; ++i) {
926	1/1 ✓ Branch 1 taken 31692 times.	31692	const std::optional<std::string_view> path = ReadField(entry, at);
927	2/3 ✓ Branch 0 taken 31692 times. ✗ Branch 1 not taken. ✓ Branch 3 taken 31692 times.	31692	const std::optional<std::string_view> digest = path ? ReadField(entry, at) : std::nullopt;
928	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 31692 times.	31692	if (!digest) {
929		✗	return std::nullopt;
930			}
931	1/1 ✓ Branch 1 taken 31692 times.	31692	kernel.files.emplace_back(path, digest);
932			}
933
934	1/1 ✓ Branch 1 taken 1117 times.	1117	kernel.code = entry.substr(at);
935		1117	return kernel;
936		1117	}
937
938		341	bool HasModule(const KernelLibrary* library, const SourceFiles& sources, std::string_view name) {
939	1/1 ✓ Branch 2 taken 341 times.	341	return FindFile(library, sources, ModulePath(name)).has_value();
940			}
941
942		1126	KernelCode Compile(Backend backend, Traversal traversal, Sampling sampling, bool debugInfo,
943			const KernelLibrary* library, const SourceFiles& sources,
944			std::span<const std::string> modules, std::string_view entry,
945			std::string_view label, const DiskCache& cache, bool narrowWaves) {
946		1126	uint64_t sourcesHash = kFnvBasis;
947	2/2 ✓ Branch 1 taken 552 times. ✓ Branch 2 taken 1126 times.	1678	for (const auto& [path, source] : sources) {
948		552	sourcesHash =
949	1/1 ✓ Branch 2 taken 552 times.	552	Fnv1a(fmt::format("{} {} {} ", path.size(), path, source.size()), sourcesHash);
950		552	sourcesHash = Fnv1a(source, sourcesHash);
951			}
952		1126	const std::string key = fmt::format(
953	1/1 ✓ Branch 1 taken 1126 times.	1126	"slang {} entry {} {}{}{}{}{} {} of {} from {:016x}", spGetBuildTagString(), kEntryFormat,
954	1/1 ✓ Branch 1 taken 1126 times.	1126	BackendName(backend), traversal == Traversal::Hardware ? " with hardware traversal" : "",
955	1/2 ✓ Branch 0 taken 1126 times. ✗ Branch 1 not taken.	1126	sampling == Sampling::Hardware ? " with hardware textures" : "",
956	4/6 ✓ Branch 0 taken 1126 times. ✗ Branch 1 not taken. ✓ Branch 2 taken 1126 times. ✗ Branch 3 not taken. ✓ Branch 4 taken 8 times. ✓ Branch 5 taken 1118 times.	4504	debugInfo ? " with debug info" : "", narrowWaves ? " with narrow waves" : "", entry,
957	3/3 ✓ Branch 2 taken 1126 times. ✓ Branch 5 taken 1126 times. ✓ Branch 9 taken 1126 times.	3378	Join({modules.begin(), modules.end()}), sourcesHash);
958
959		1126	std::optional<KernelCode> kernel;
960	3/3 ✓ Branch 1 taken 1126 times. ✓ Branch 3 taken 1117 times. ✓ Branch 4 taken 9 times.	1126	if (const std::optional<std::string> cached = cache.Load(key)) {
961	1/1 ✓ Branch 1 taken 1117 times.	2234	kernel = KernelCode::Decode(*cached);
962	3/5 ✓ Branch 0 taken 1117 times. ✗ Branch 1 not taken. ✓ Branch 3 taken 1117 times. ✗ Branch 5 not taken. ✓ Branch 6 taken 1117 times.	1117	if (kernel && !Unchanged(library, sources, kernel->files)) {
963		✗	kernel.reset();
964			}
965		✗	}
966	2/2 ✓ Branch 0 taken 1117 times. ✓ Branch 1 taken 9 times.	1126	if (kernel) {
967		1117	kernel->cached = true;
968			} else {
969	1/1 ✓ Branch 1 taken 7 times.	16	kernel = CompileWithSlang(backend, traversal, sampling, debugInfo, library, sources,
970		7	modules, entry, label, narrowWaves, cache);
971	2/2 ✓ Branch 1 taken 7 times. ✓ Branch 4 taken 7 times.	16	cache.Store(key, kernel->Encode());
972			}
973
974		1124	kernel->debugInfo = debugInfo;
975		1124	return *std::move(kernel);
976		1126	}
977
978			}
979