11779645611

Committed 11 Nov 2024 01:37PM UTC coverage: 95.165% (-0.002%) from 95.167%

Build # 11779645611

Build Type

Pull #303

github

Committed by

PeterTh

Commit Message

Attempt workaround for suspected DPCPP thread safety bug

Pull Request Pull Request #303: Introduce per-device dispatch workers in the SYCL backend

Run Details

3032 of 3432 branches covered (88.34%)

Branch coverage included in aggregate %.

68 of 69 new or added lines in 6 files covered. (98.55%)

3 existing lines in 1 file now uncovered.

6751 of 6848 relevant lines covered (98.58%)

1276938.92 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

92.27

/src/backend/sycl_backend.cc

#include "backend/sycl_backend.h"

#include "closure_hydrator.h"
#include "dense_map.h"
#include "nd_memory.h"
#include "system_info.h"
#include "thread_queue.h"
#include "types.h"

#include <fmt/ranges.h>

namespace {
void wait_synchronously_for(const celerity::detail::async_event& ev) {
        while(!ev.is_complete()) {
                std::this_thread::yield();
        }
}
} // namespace

namespace celerity::detail::sycl_backend_detail {

bool sycl_event::is_complete() { return m_last.get_info<sycl::info::event::command_execution_status>() == sycl::info::event_command_status::complete; }

std::optional<std::chrono::nanoseconds> sycl_event::get_native_execution_time() {
        if(!m_first.has_value()) return std::nullopt; // avoid the cost of throwing + catching a sycl exception by when profiling is disabled
        return std::chrono::nanoseconds(m_last.get_profiling_info<sycl::info::event_profiling::command_end>() //
                                        - m_first->get_profiling_info<sycl::info::event_profiling::command_start>());
}

bool delayed_async_event::is_complete() {
        std::lock_guard lock{m_data->mutex};
        if(!m_data->event.has_value()) return false;
        return m_data->event->is_complete();
}

void* delayed_async_event::get_result() {
        std::lock_guard lock{m_data->mutex};
        assert(m_data->event.has_value());
        return m_data->event->get_result();
}

std::optional<std::chrono::nanoseconds> delayed_async_event::get_native_execution_time() {
        std::lock_guard lock{m_data->mutex};
        assert(m_data->event.has_value());
        return m_data->event->get_native_execution_time();
}

void flush(sycl::queue& queue) {
#if CELERITY_WORKAROUND(ACPP)
        // AdaptiveCpp does not guarantee that command groups are actually scheduled until an explicit await operation, which we cannot insert without
        // blocking the executor loop (see https://github.com/AdaptiveCpp/AdaptiveCpp/issues/599). Instead, we explicitly flush the queue to be able to continue
        // using our polling-based approach.
        queue.get_context().AdaptiveCpp_runtime()->dag().flush_async();
#else
        (void)queue;
#endif
}

// LCOV_EXCL_START
void report_errors(const sycl::exception_list& errors) {
        if(errors.size() == 0) return;

        std::vector<std::string> what;
        for(const auto& e : errors) {
                try {
                        std::rethrow_exception(e);
                } catch(sycl::exception& e) { //
                        what.push_back(e.what());
                } catch(std::exception& e) { //
                        what.push_back(e.what());
                } catch(...) { //
                        what.push_back("unknown exception");
                }
        }

        // Errors usually manifest on calls to sycl::event::get_info(), not their actual origin, and therefore will contain many duplicates
        std::sort(what.begin(), what.end());
        what.erase(std::unique(what.begin(), what.end()), what.end());

        utils::panic("asynchronous SYCL errors:\n\t{}", fmt::join(what, "\n\t"));
}
// LCOV_EXCL_STOP

} // namespace celerity::detail::sycl_backend_detail

namespace celerity::detail {

struct sycl_backend::impl {
        struct device_state {
                sycl::device sycl_device;
                sycl::context sycl_context;
                std::vector<sycl::queue> queues;
                std::optional<detail::thread_queue> worker_thread;

                device_state() = default;
                explicit device_state(const sycl::device& dev) : sycl_device(dev), sycl_context(sycl_device) {}
        };

        struct host_state {
                sycl::context sycl_context;
                thread_queue alloc_queue;
                std::vector<thread_queue> queues; // TODO naming vs alloc_queue?

                // pass devices to ensure the sycl_context receives the correct platform
                explicit host_state(const std::vector<sycl::device>& all_devices, bool enable_profiling)
                    // DPC++ requires exactly one CUDA device here, but for allocation the sycl_context mostly means "platform".
                    // - TODO assert that all devices belong to the same platform + backend here
                    // - TODO test Celerity on a (SimSYCL) system without GPUs
                    : sycl_context(all_devices.at(0)), //
                      alloc_queue("cy-alloc", enable_profiling) {}
        };

        system_info system;
        dense_map<device_id, device_state> devices; // thread-safe for read access (not resized after construction)
        host_state host;
        using configuration = sycl_backend::configuration;
        configuration config;

        impl(const std::vector<sycl::device>& devices, const configuration& config)
            : devices(devices.begin(), devices.end()), host(devices, config.profiling), config(config) //
        {
                // For now, we assume distinct memories per device. TODO some targets, (OpenMP emulated devices), might deviate from that.
                system.devices.resize(devices.size());
                system.memories.resize(2 + devices.size()); //  user + host + device memories
                system.memories[user_memory_id].copy_peers.set(user_memory_id);
                system.memories[host_memory_id].copy_peers.set(host_memory_id);
                system.memories[host_memory_id].copy_peers.set(user_memory_id);
                system.memories[user_memory_id].copy_peers.set(host_memory_id);
                for(device_id did = 0; did < devices.size(); ++did) {
                        const memory_id mid = first_device_memory_id + did;
                        system.devices[did].native_memory = mid;
                        system.memories[mid].copy_peers.set(mid);
                        system.memories[mid].copy_peers.set(host_memory_id);
                        system.memories[host_memory_id].copy_peers.set(mid);
                        // device-to-device copy capabilities are added in cuda_backend constructor
                }
        }

        template <typename F>
        async_event submit_alloc(F&& f) {
#if CELERITY_WORKAROUND(SIMSYCL)
                // SimSYCL is not thread safe => skip alloc_queue and complete allocations in executor thread.
                if constexpr(std::is_void_v<std::invoke_result_t<F>>) {
                        return f(), make_complete_event();
                } else {
                        return make_complete_event(f());
                }
#else
                return host.alloc_queue.submit(std::forward<F>(f));
#endif
        }

        thread_queue& get_host_queue(const size_t lane) {
                assert(lane <= host.queues.size());
                if(lane == host.queues.size()) { host.queues.emplace_back(fmt::format("cy-host-{}", lane), config.profiling); }
                return host.queues[lane];
        }

        sycl::queue& get_device_queue(const device_id did, const size_t lane) {
                auto& device = devices[did];
                assert(lane <= device.queues.size());
                if(lane == device.queues.size()) {
                        const auto properties = config.profiling ? sycl::property_list{sycl::property::queue::enable_profiling{}, sycl::property::queue::in_order{}}
                                                                 : sycl::property_list{sycl::property::queue::in_order{}};
                        device.queues.emplace_back(device.sycl_device, sycl::async_handler(sycl_backend_detail::report_errors), properties);
                }
                return device.queues[lane];
        }
};

sycl_backend::sycl_backend(const std::vector<sycl::device>& devices, const configuration& config) : m_impl(new impl(devices, config)) {
        // Initialize a worker thread with hydrator for each device, if they are enabled
        if(m_impl->config.per_device_worker_threads) {
                for(device_id did = 0; did < m_impl->system.devices.size(); ++did) {
                        m_impl->devices[did].worker_thread.emplace(fmt::format("cy-be-worker-{}", did.value), m_impl->config.profiling);
                        wait_synchronously_for(m_impl->devices[did].worker_thread->submit([did] {
                                closure_hydrator::make_available();
                                // TODO: Set BE worker thread affinity
                        }));
                }
        }
}

sycl_backend::~sycl_backend() {
        // If we are using worker threads, tear down their hydrators before they are destroyed
        if(m_impl->config.per_device_worker_threads) {
                for(auto& device : m_impl->devices) {
                        wait_synchronously_for(device.worker_thread->submit([] { closure_hydrator::teardown(); }));
                }
        }
}

const system_info& sycl_backend::get_system_info() const { return m_impl->system; }

void sycl_backend::init() {
        CELERITY_DETAIL_TRACY_ZONE_SCOPED("sycl::init", Orange2);

        // Instantiate the first in-order queue on each device. At least for CUDA systems this will perform device initialization, which can take > 100 ms / device.
        for(device_id did = 0; did < m_impl->system.devices.size(); ++did) {
                (void)m_impl->get_device_queue(did, 0 /* lane */);
        }
        // TODO: Set executor thread affinity
}

void* sycl_backend::debug_alloc(const size_t size) {
        const auto ptr = sycl::malloc_host(size, m_impl->host.sycl_context);
#if CELERITY_DETAIL_ENABLE_DEBUG
        memset(ptr, static_cast<int>(sycl_backend_detail::uninitialized_memory_pattern), size);
#endif
        return ptr;
}

void sycl_backend::debug_free(void* const ptr) { sycl::free(ptr, m_impl->host.sycl_context); }

async_event sycl_backend::enqueue_host_alloc(const size_t size, const size_t alignment) {
        return m_impl->submit_alloc([this, size, alignment] {
                const auto ptr = sycl::aligned_alloc_host(alignment, size, m_impl->host.sycl_context);
#if CELERITY_DETAIL_ENABLE_DEBUG
                memset(ptr, static_cast<int>(sycl_backend_detail::uninitialized_memory_pattern), size);
#endif
                return ptr;
        });
}

async_event sycl_backend::enqueue_device_alloc(const device_id device, const size_t size, const size_t alignment) {
        return m_impl->submit_alloc([this, device, size, alignment] {
                auto& d = m_impl->devices[device];
                const auto ptr = sycl::aligned_alloc_device(alignment, size, d.sycl_device, d.sycl_context);
#if CELERITY_DETAIL_ENABLE_DEBUG
                sycl::queue(d.sycl_context, d.sycl_device, sycl::async_handler(sycl_backend_detail::report_errors), sycl::property::queue::in_order{})
                    .fill(ptr, sycl_backend_detail::uninitialized_memory_pattern, size)
                    .wait_and_throw();
#endif
                return ptr;
        });
}

async_event sycl_backend::enqueue_host_free(void* const ptr) {
        return m_impl->submit_alloc([this, ptr] { sycl::free(ptr, m_impl->host.sycl_context); });
}

async_event sycl_backend::enqueue_device_free(const device_id device, void* const ptr) {
        return m_impl->submit_alloc([this, device, ptr] { sycl::free(ptr, m_impl->devices[device].sycl_context); });
}

async_event sycl_backend::enqueue_host_task(size_t host_lane, const host_task_launcher& launcher, std::vector<closure_hydrator::accessor_info> accessor_infos,
    const box<3>& execution_range, const communicator* collective_comm) //
{
        auto& hydrator = closure_hydrator::get_instance();
        hydrator.arm(target::host_task, std::move(accessor_infos));
        auto launch_hydrated = hydrator.hydrate<target::host_task>(launcher);
        return m_impl->get_host_queue(host_lane).submit([=, launch_hydrated = std::move(launch_hydrated)] { launch_hydrated(execution_range, collective_comm); });
}

async_event sycl_backend::enqueue_device_kernel(const device_id device, const size_t lane, const device_kernel_launcher& launch,
    std::vector<closure_hydrator::accessor_info> accessor_infos, const box<3>& execution_range, const std::vector<void*>& reduction_ptrs) //
{
        return enqueue_device_work(device, [=, this] {
                auto& queue = m_impl->get_device_queue(device, lane);
                CELERITY_DETAIL_TRACY_ZONE_SCOPED("sycl::submit", Orange2);
                auto event = queue.submit([&](sycl::handler& sycl_cgh) {
                        auto& hydrator = closure_hydrator::get_instance();
                        hydrator.arm(target::device, std::move(accessor_infos));
                        const auto launch_hydrated = hydrator.hydrate<target::device>(sycl_cgh, launch);
                        launch_hydrated(sycl_cgh, execution_range, reduction_ptrs);
                });
                sycl_backend_detail::flush(queue);
                return make_async_event<sycl_backend_detail::sycl_event>(std::move(event), m_impl->config.profiling);
        });
}

async_event sycl_backend::enqueue_host_copy(size_t host_lane, const void* const source_base, void* const dest_base, const region_layout& source_layout,
    const region_layout& dest_layout, const region<3>& copy_region, const size_t elem_size) //
{
        return m_impl->get_host_queue(host_lane).submit([=] { nd_copy_host(source_base, dest_base, source_layout, dest_layout, copy_region, elem_size); });
}

void sycl_backend::check_async_errors() {
        for(size_t i = 0; i < m_impl->devices.size(); ++i) {
                // throw_asynchronous should be thread safe, however, in dpcpp it crashes sporadically if not called from the right thread
                // as such, we need to perform this operation on the device worker thread
                auto& device = m_impl->devices[i];
                wait_synchronously_for(enqueue_device_work(i, [&]() {
                        for(auto& queue : device.queues) {
                                queue.throw_asynchronous();
                        }
                        return make_complete_event();
                }));
        }
}

sycl::queue& sycl_backend::get_device_queue(device_id device, size_t lane) { return m_impl->get_device_queue(device, lane); }

system_info& sycl_backend::get_system_info() { return m_impl->system; }

async_event celerity::detail::sycl_backend::enqueue_device_work(device_id device, std::function<async_event()> work) {
        // Basic case: no per-device worker threads
        if(!m_impl->config.per_device_worker_threads) { return work(); }

        assert(device < m_impl->devices.size());
        auto& worker = m_impl->devices[device].worker_thread;
        assert(worker.has_value());

        sycl_backend_detail::delayed_async_event::shared_data_block async_event_data = std::make_shared<sycl_backend_detail::delayed_async_event::data_block>();
        auto async_event = make_async_event<sycl_backend_detail::delayed_async_event>(async_event_data);

        [[maybe_unused]] auto ev = worker->submit([work, async_event_data] {
                auto event = work();
                std::lock_guard lock{async_event_data->mutex};
                async_event_data->event = std::move(event);
        });
        return async_event;
}

bool sycl_backend::is_profiling_enabled() const { return m_impl->config.profiling; }

std::vector<sycl_backend_type> sycl_backend_enumerator::compatible_backends(const sycl::device& device) const {
        std::vector<backend_type> backends{backend_type::generic};
#if CELERITY_WORKAROUND(ACPP) && defined(SYCL_EXT_HIPSYCL_BACKEND_CUDA)
        if(device.get_backend() == sycl::backend::cuda) { backends.push_back(sycl_backend_type::cuda); }
#elif CELERITY_WORKAROUND(DPCPP)
        if(device.get_backend() == sycl::backend::ext_oneapi_cuda) { backends.push_back(sycl_backend_type::cuda); }
#endif
        assert(std::is_sorted(backends.begin(), backends.end()));
        return backends;
}

std::vector<sycl_backend_type> sycl_backend_enumerator::available_backends() const {
        std::vector<backend_type> backends{backend_type::generic};
#if CELERITY_DETAIL_BACKEND_CUDA_ENABLED
        backends.push_back(sycl_backend_type::cuda);
#endif
        assert(std::is_sorted(backends.begin(), backends.end()));
        return backends;
}

bool sycl_backend_enumerator::is_specialized(backend_type type) const {
        switch(type) {
        case backend_type::generic: return false;
        case backend_type::cuda: return true;
        default: utils::unreachable(); // LCOV_EXCL_LINE
        }
}

int sycl_backend_enumerator::get_priority(backend_type type) const {
        switch(type) {
        case backend_type::generic: return 0;
        case backend_type::cuda: return 1;
        default: utils::unreachable(); // LCOV_EXCL_LINE
        }
}

} // namespace celerity::detail

namespace celerity::detail {

std::unique_ptr<backend> make_sycl_backend(const sycl_backend_type type, const std::vector<sycl::device>& devices, const sycl_backend::configuration& config) {
        assert(std::all_of(
            devices.begin(), devices.end(), [=](const sycl::device& d) { return utils::contains(sycl_backend_enumerator{}.compatible_backends(d), type); }));

        switch(type) {
        case sycl_backend_type::generic: //
                return std::make_unique<sycl_generic_backend>(devices, config);

        case sycl_backend_type::cuda:
#if CELERITY_DETAIL_BACKEND_CUDA_ENABLED
                return std::make_unique<sycl_cuda_backend>(devices, config);
#else
                utils::panic("CUDA backend has not been compiled");
#endif
        }
        utils::unreachable(); // LCOV_EXCL_LINE
}

} // namespace celerity::detail

1	#include "backend/sycl_backend.h"
2
3	#include "closure_hydrator.h"
4	#include "dense_map.h"
5	#include "nd_memory.h"
6	#include "system_info.h"
7	#include "thread_queue.h"
8	#include "types.h"
9
10	#include <fmt/ranges.h>
11
12	namespace {
13	void wait_synchronously_for(const celerity::detail::async_event& ev) {	23,277✔
14	while(!ev.is_complete()) {	50,323✔
15	std::this_thread::yield();	27,046✔
16	}
17	}	23,277✔
18	} // namespace
19
20	namespace celerity::detail::sycl_backend_detail {
21
22	bool sycl_event::is_complete() { return m_last.get_info<sycl::info::event::command_execution_status>() == sycl::info::event_command_status::complete; }	3,810✔
23
24	std::optional<std::chrono::nanoseconds> sycl_event::get_native_execution_time() {	832✔
25	if(!m_first.has_value()) return std::nullopt; // avoid the cost of throwing + catching a sycl exception by when profiling is disabled	832✔
26	return std::chrono::nanoseconds(m_last.get_profiling_info<sycl::info::event_profiling::command_end>() //	4✔
27	- m_first->get_profiling_info<sycl::info::event_profiling::command_start>());	6✔
28	}
29
30	bool delayed_async_event::is_complete() {	374,900✔
31	std::lock_guard lock{m_data->mutex};	374,900✔
32	if(!m_data->event.has_value()) return false;	374,900✔
33	return m_data->event->is_complete();	823✔
34	}	374,900✔
35
36	void* delayed_async_event::get_result() {	823✔
37	std::lock_guard lock{m_data->mutex};	823✔
38	assert(m_data->event.has_value());	823✔
39	return m_data->event->get_result();	1,646✔
40	}	823✔
41
42	std::optional<std::chrono::nanoseconds> delayed_async_event::get_native_execution_time() {	4✔
43	std::lock_guard lock{m_data->mutex};	4✔
44	assert(m_data->event.has_value());	4✔
45	return m_data->event->get_native_execution_time();	8✔
46	}	4✔
47
48	void flush(sycl::queue& queue) {	3,810✔
49	#if CELERITY_WORKAROUND(ACPP)
50	// AdaptiveCpp does not guarantee that command groups are actually scheduled until an explicit await operation, which we cannot insert without
51	// blocking the executor loop (see https://github.com/AdaptiveCpp/AdaptiveCpp/issues/599). Instead, we explicitly flush the queue to be able to continue
52	// using our polling-based approach.
53	queue.get_context().AdaptiveCpp_runtime()->dag().flush_async();
54	#else
55	(void)queue;
56	#endif
57	}	3,810✔
58
59	// LCOV_EXCL_START
60	void report_errors(const sycl::exception_list& errors) {
61	if(errors.size() == 0) return;
62
63	std::vector<std::string> what;
64	for(const auto& e : errors) {
65	try {
66	std::rethrow_exception(e);
67	} catch(sycl::exception& e) { //
68	what.push_back(e.what());
69	} catch(std::exception& e) { //
70	what.push_back(e.what());
71	} catch(...) { //
72	what.push_back("unknown exception");
73	}
74	}
75
76	// Errors usually manifest on calls to sycl::event::get_info(), not their actual origin, and therefore will contain many duplicates
77	std::sort(what.begin(), what.end());
78	what.erase(std::unique(what.begin(), what.end()), what.end());
79
80	utils::panic("asynchronous SYCL errors:\n\t{}", fmt::join(what, "\n\t"));
81	}
82	// LCOV_EXCL_STOP
83
84	} // namespace celerity::detail::sycl_backend_detail
85
86	namespace celerity::detail {
87
88	struct sycl_backend::impl {
89	struct device_state {
90	sycl::device sycl_device;
91	sycl::context sycl_context;
92	std::vector<sycl::queue> queues;
93	std::optional<detail::thread_queue> worker_thread;
94
95	device_state() = default;
96	explicit device_state(const sycl::device& dev) : sycl_device(dev), sycl_context(sycl_device) {}	564✔
97	};
98
99	struct host_state {
100	sycl::context sycl_context;
101	thread_queue alloc_queue;
102	std::vector<thread_queue> queues; // TODO naming vs alloc_queue?
103
104	// pass devices to ensure the sycl_context receives the correct platform
105	explicit host_state(const std::vector<sycl::device>& all_devices, bool enable_profiling)	239✔
106	// DPC++ requires exactly one CUDA device here, but for allocation the sycl_context mostly means "platform".
107	// - TODO assert that all devices belong to the same platform + backend here
108	// - TODO test Celerity on a (SimSYCL) system without GPUs
109	: sycl_context(all_devices.at(0)), //	239✔
110	alloc_queue("cy-alloc", enable_profiling) {}	717✔
111	};
112
113	system_info system;
114	dense_map<device_id, device_state> devices; // thread-safe for read access (not resized after construction)
115	host_state host;
116	using configuration = sycl_backend::configuration;
117	configuration config;
118
119	impl(const std::vector<sycl::device>& devices, const configuration& config)	239✔
120	: devices(devices.begin(), devices.end()), host(devices, config.profiling), config(config) //	239✔
121	{
122	// For now, we assume distinct memories per device. TODO some targets, (OpenMP emulated devices), might deviate from that.
123	system.devices.resize(devices.size());	239✔
124	system.memories.resize(2 + devices.size()); // user + host + device memories	239✔
125	system.memories[user_memory_id].copy_peers.set(user_memory_id);	239✔
126	system.memories[host_memory_id].copy_peers.set(host_memory_id);	239✔
127	system.memories[host_memory_id].copy_peers.set(user_memory_id);	239✔
128	system.memories[user_memory_id].copy_peers.set(host_memory_id);	239✔
129	for(device_id did = 0; did < devices.size(); ++did) {	803✔
130	const memory_id mid = first_device_memory_id + did;	564✔
131	system.devices[did].native_memory = mid;	564✔
132	system.memories[mid].copy_peers.set(mid);	564✔
133	system.memories[mid].copy_peers.set(host_memory_id);	564✔
134	system.memories[host_memory_id].copy_peers.set(mid);	564✔
135	// device-to-device copy capabilities are added in cuda_backend constructor
136	}
137	}	239✔
138
139	template <typename F>
140	async_event submit_alloc(F&& f) {	2,212✔
141	#if CELERITY_WORKAROUND(SIMSYCL)
142	// SimSYCL is not thread safe => skip alloc_queue and complete allocations in executor thread.
143	if constexpr(std::is_void_v<std::invoke_result_t<F>>) {
144	return f(), make_complete_event();	1,106✔
145	} else {
146	return make_complete_event(f());	1,106✔
147	}
148	#else
149	return host.alloc_queue.submit(std::forward<F>(f));
150	#endif
151	}
152
153	thread_queue& get_host_queue(const size_t lane) {	1,646✔
154	assert(lane <= host.queues.size());	1,646✔
155	if(lane == host.queues.size()) { host.queues.emplace_back(fmt::format("cy-host-{}", lane), config.profiling); }	1,809✔
156	return host.queues[lane];	1,646✔
157	}
158
159	sycl::queue& get_device_queue(const device_id did, const size_t lane) {	4,279✔
160	auto& device = devices[did];	4,279✔
161	assert(lane <= device.queues.size());	4,279✔
162	if(lane == device.queues.size()) {	4,279✔
163	const auto properties = config.profiling ? sycl::property_list{sycl::property::queue::enable_profiling{}, sycl::property::queue::in_order{}}	481✔
164	: sycl::property_list{sycl::property::queue::in_order{}};	481✔
165	device.queues.emplace_back(device.sycl_device, sycl::async_handler(sycl_backend_detail::report_errors), properties);	481✔
166	}	481✔
167	return device.queues[lane];	4,279✔
168	}
169	};
170
171	sycl_backend::sycl_backend(const std::vector<sycl::device>& devices, const configuration& config) : m_impl(new impl(devices, config)) {	239!
172	// Initialize a worker thread with hydrator for each device, if they are enabled
173	if(m_impl->config.per_device_worker_threads) {	239✔
174	for(device_id did = 0; did < m_impl->system.devices.size(); ++did) {	90✔
175	m_impl->devices[did].worker_thread.emplace(fmt::format("cy-be-worker-{}", did.value), m_impl->config.profiling);	144✔
176	wait_synchronously_for(m_impl->devices[did].worker_thread->submit([did] {	72✔
177	closure_hydrator::make_available();	72✔
178	// TODO: Set BE worker thread affinity
179	}));	72✔
180	}
181	}
182	}	239✔
183
184	sycl_backend::~sycl_backend() {	239✔
185	// If we are using worker threads, tear down their hydrators before they are destroyed
186	if(m_impl->config.per_device_worker_threads) {	239✔
187	for(auto& device : m_impl->devices) {	90✔
188	wait_synchronously_for(device.worker_thread->submit([] { closure_hydrator::teardown(); }));	144✔
189	}
190	}
191	}	239✔
192
193	const system_info& sycl_backend::get_system_info() const { return m_impl->system; }	433✔
194
195	void sycl_backend::init() {	213✔
196	CELERITY_DETAIL_TRACY_ZONE_SCOPED("sycl::init", Orange2);
197
198	// Instantiate the first in-order queue on each device. At least for CUDA systems this will perform device initialization, which can take > 100 ms / device.
199	for(device_id did = 0; did < m_impl->system.devices.size(); ++did) {	682✔
200	(void)m_impl->get_device_queue(did, 0 /* lane */);	469✔
201	}
202	// TODO: Set executor thread affinity
203	}	213✔
204
205	void* sycl_backend::debug_alloc(const size_t size) {	1,934✔
206	const auto ptr = sycl::malloc_host(size, m_impl->host.sycl_context);	1,934✔
207	#if CELERITY_DETAIL_ENABLE_DEBUG
208	memset(ptr, static_cast<int>(sycl_backend_detail::uninitialized_memory_pattern), size);	1,934✔
209	#endif
210	return ptr;	1,934✔
211	}
212
213	void sycl_backend::debug_free(void* const ptr) { sycl::free(ptr, m_impl->host.sycl_context); }	1,934✔
214
215	async_event sycl_backend::enqueue_host_alloc(const size_t size, const size_t alignment) {	550✔
216	return m_impl->submit_alloc([this, size, alignment] {	550✔
217	const auto ptr = sycl::aligned_alloc_host(alignment, size, m_impl->host.sycl_context);	550✔
218	#if CELERITY_DETAIL_ENABLE_DEBUG
219	memset(ptr, static_cast<int>(sycl_backend_detail::uninitialized_memory_pattern), size);	550✔
220	#endif
221	return ptr;	550✔
222	});	1,100✔
223	}
224
225	async_event sycl_backend::enqueue_device_alloc(const device_id device, const size_t size, const size_t alignment) {	556✔
226	return m_impl->submit_alloc([this, device, size, alignment] {	556✔
227	auto& d = m_impl->devices[device];	556✔
228	const auto ptr = sycl::aligned_alloc_device(alignment, size, d.sycl_device, d.sycl_context);	556✔
229	#if CELERITY_DETAIL_ENABLE_DEBUG
230	sycl::queue(d.sycl_context, d.sycl_device, sycl::async_handler(sycl_backend_detail::report_errors), sycl::property::queue::in_order{})	1,112✔
231	.fill(ptr, sycl_backend_detail::uninitialized_memory_pattern, size)	2,224✔
232	.wait_and_throw();	1,112✔
233	#endif
234	return ptr;	556✔
235	});	1,112✔
236	}
237
238	async_event sycl_backend::enqueue_host_free(void* const ptr) {	550✔
239	return m_impl->submit_alloc([this, ptr] { sycl::free(ptr, m_impl->host.sycl_context); });	1,650✔
240	}
241
242	async_event sycl_backend::enqueue_device_free(const device_id device, void* const ptr) {	556✔
243	return m_impl->submit_alloc([this, device, ptr] { sycl::free(ptr, m_impl->devices[device].sycl_context); });	1,668✔
244	}
245
246	async_event sycl_backend::enqueue_host_task(size_t host_lane, const host_task_launcher& launcher, std::vector<closure_hydrator::accessor_info> accessor_infos,	1,182✔
247	const box<3>& execution_range, const communicator* collective_comm) //
248	{
249	auto& hydrator = closure_hydrator::get_instance();	1,182✔
250	hydrator.arm(target::host_task, std::move(accessor_infos));	1,182✔
251	auto launch_hydrated = hydrator.hydrate<target::host_task>(launcher);	1,182✔
252	return m_impl->get_host_queue(host_lane).submit([=, launch_hydrated = std::move(launch_hydrated)] { launch_hydrated(execution_range, collective_comm); });	3,546!
253	}	1,182✔
254
255	async_event sycl_backend::enqueue_device_kernel(const device_id device, const size_t lane, const device_kernel_launcher& launch,	863✔
256	std::vector<closure_hydrator::accessor_info> accessor_infos, const box<3>& execution_range, const std::vector<void*>& reduction_ptrs) //
257	{
258	return enqueue_device_work(device, [=, this] {	1,726✔
259	auto& queue = m_impl->get_device_queue(device, lane);	863✔
260	CELERITY_DETAIL_TRACY_ZONE_SCOPED("sycl::submit", Orange2);
261	auto event = queue.submit([&](sycl::handler& sycl_cgh) {	863✔
262	auto& hydrator = closure_hydrator::get_instance();	863✔
263	hydrator.arm(target::device, std::move(accessor_infos));	863✔
264	const auto launch_hydrated = hydrator.hydrate<target::device>(sycl_cgh, launch);	863✔
265	launch_hydrated(sycl_cgh, execution_range, reduction_ptrs);	863✔
266	});	2,589✔
267	sycl_backend_detail::flush(queue);	863✔
268	return make_async_event<sycl_backend_detail::sycl_event>(std::move(event), m_impl->config.profiling);	1,726✔
269	});	3,452✔
270	}
271
272	async_event sycl_backend::enqueue_host_copy(size_t host_lane, const void* const source_base, void* const dest_base, const region_layout& source_layout,	464✔
273	const region_layout& dest_layout, const region<3>& copy_region, const size_t elem_size) //
274	{
275	return m_impl->get_host_queue(host_lane).submit([=] { nd_copy_host(source_base, dest_base, source_layout, dest_layout, copy_region, elem_size); });	928!
276	}
277
278	void sycl_backend::check_async_errors() {	15,494✔
279	for(size_t i = 0; i < m_impl->devices.size(); ++i) {	38,627✔
280	// throw_asynchronous should be thread safe, however, in dpcpp it crashes sporadically if not called from the right thread
281	// as such, we need to perform this operation on the device worker thread
282	auto& device = m_impl->devices[i];	23,133✔
283	wait_synchronously_for(enqueue_device_work(i, [&]() {	23,133✔
284	for(auto& queue : device.queues) {	46,266✔
285	queue.throw_asynchronous();	23,133✔
286	}
287	return make_complete_event();	23,133✔
288	}));
289	}
290	}	15,494✔
291
292	sycl::queue& sycl_backend::get_device_queue(device_id device, size_t lane) { return m_impl->get_device_queue(device, lane); }	2,947✔
293
NEW 294	system_info& sycl_backend::get_system_info() { return m_impl->system; }	×
295
296	async_event celerity::detail::sycl_backend::enqueue_device_work(device_id device, std::function<async_event()> work) {	26,943✔
297	// Basic case: no per-device worker threads
298	if(!m_impl->config.per_device_worker_threads) { return work(); }	26,943✔
299
300	assert(device < m_impl->devices.size());	823✔
301	auto& worker = m_impl->devices[device].worker_thread;	823✔
302	assert(worker.has_value());	823✔
303
304	sycl_backend_detail::delayed_async_event::shared_data_block async_event_data = std::make_shared<sycl_backend_detail::delayed_async_event::data_block>();	823✔
305	auto async_event = make_async_event<sycl_backend_detail::delayed_async_event>(async_event_data);	823✔
306
307	[[maybe_unused]] auto ev = worker->submit([work, async_event_data] {	2,469!
308	auto event = work();	823✔
309	std::lock_guard lock{async_event_data->mutex};	823✔
310	async_event_data->event = std::move(event);	823✔
311	});	3,292✔
312	return async_event;	823✔
313	}	823✔
314
315	bool sycl_backend::is_profiling_enabled() const { return m_impl->config.profiling; }	2,947✔
316
317	std::vector<sycl_backend_type> sycl_backend_enumerator::compatible_backends(const sycl::device& device) const {	1,128✔
318	std::vector<backend_type> backends{backend_type::generic};	3,384✔
319	#if CELERITY_WORKAROUND(ACPP) && defined(SYCL_EXT_HIPSYCL_BACKEND_CUDA)
320	if(device.get_backend() == sycl::backend::cuda) { backends.push_back(sycl_backend_type::cuda); }
321	#elif CELERITY_WORKAROUND(DPCPP)
322	if(device.get_backend() == sycl::backend::ext_oneapi_cuda) { backends.push_back(sycl_backend_type::cuda); }
323	#endif
324	assert(std::is_sorted(backends.begin(), backends.end()));	1,128✔
325	return backends;	1,128✔
UNCOV 326	}	×
327
328	std::vector<sycl_backend_type> sycl_backend_enumerator::available_backends() const {	229✔
329	std::vector<backend_type> backends{backend_type::generic};	687✔
330	#if CELERITY_DETAIL_BACKEND_CUDA_ENABLED
331	backends.push_back(sycl_backend_type::cuda);
332	#endif
333	assert(std::is_sorted(backends.begin(), backends.end()));	229✔
334	return backends;	229✔
UNCOV 335	}	×
336
337	bool sycl_backend_enumerator::is_specialized(backend_type type) const {	221✔
338	switch(type) {	221!
339	case backend_type::generic: return false;	220✔
340	case backend_type::cuda: return true;	1✔
341	default: utils::unreachable(); // LCOV_EXCL_LINE
342	}
343	}
344
345	int sycl_backend_enumerator::get_priority(backend_type type) const {	2✔
346	switch(type) {	2!
347	case backend_type::generic: return 0;	1✔
348	case backend_type::cuda: return 1;	1✔
349	default: utils::unreachable(); // LCOV_EXCL_LINE
350	}
351	}
352
353	} // namespace celerity::detail
354
355	namespace celerity::detail {
356
357	std::unique_ptr<backend> make_sycl_backend(const sycl_backend_type type, const std::vector<sycl::device>& devices, const sycl_backend::configuration& config) {	239✔
358	assert(std::all_of(	803✔
359	devices.begin(), devices.end(), [=](const sycl::device& d) { return utils::contains(sycl_backend_enumerator{}.compatible_backends(d), type); }));
360
361	switch(type) {	239!
362	case sycl_backend_type::generic: //	239✔
363	return std::make_unique<sycl_generic_backend>(devices, config);	239✔
364
365	case sycl_backend_type::cuda:	×
366	#if CELERITY_DETAIL_BACKEND_CUDA_ENABLED
367	return std::make_unique<sycl_cuda_backend>(devices, config);
368	#else
UNCOV 369	utils::panic("CUDA backend has not been compiled");	×
370	#endif
371	}
372	utils::unreachable(); // LCOV_EXCL_LINE
373	}
374
375	} // namespace celerity::detail

celerity / celerity-runtime / 11779645611

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous