diff --git a/.github/workflows/celerity_ci.yml b/.github/workflows/celerity_ci.yml
index af222e21c..87fe7672f 100644
--- a/.github/workflows/celerity_ci.yml
+++ b/.github/workflows/celerity_ci.yml
@@ -165,14 +165,17 @@ jobs:
         timeout-minutes: 5
         working-directory: ${{ env.build-dir }}
         run: ${{ env.container-workspace }}/ci/run-unit-tests.sh
-      - name: Run integration tests
+      - name: Run examples
         timeout-minutes: 5
         # We build examples twice, but only run the installed version (which probably has more failure modes)
         working-directory: ${{ env.examples-build-dir }}
-        run: ${{ env.container-workspace }}/ci/run-integration-tests.sh /data/Lenna.png 1 2 4
+        run: ${{ env.container-workspace }}/ci/run-examples.sh /data/Lenna.png 1 2 4
       - name: Run system tests
         working-directory: ${{ env.build-dir }}
         run: ${{ env.container-workspace }}/ci/run-system-tests.sh 2 4
+      - name: Run integration tests
+        working-driectory: ${{ env.build-dir }}
+        run: ${{ env.container-workspace }}/test/integration/run-integration-tests.py . ${{ matrix.platform }}
       - name: Upload stack traces (if any)
         if: always()
         uses: actions/upload-artifact@v2
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 388f4a3b1..d450ad21c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -262,6 +262,9 @@ target_link_libraries(celerity_runtime PUBLIC
   ${SYCL_LIB}
 )
 
+add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/src/backend)
+target_link_libraries(celerity_runtime PUBLIC celerity_backends)
+
 # For debug builds, we set the CELERITY_DETAIL_ENABLE_DEBUG preprocessor flag,
 # which allows Celerity to control debug functionality within headers regardless
 # of a user target's build type. (This flag is not intended to be modified by
@@ -362,7 +365,7 @@ install(
   DESTINATION include/celerity/vendor
 )
 install(
-  TARGETS celerity_runtime
+  TARGETS celerity_runtime celerity_backends
   EXPORT install_exports
   LIBRARY DESTINATION lib
   ARCHIVE DESTINATION lib
diff --git a/ci/run-integration-tests.sh b/ci/run-examples.sh
similarity index 100%
rename from ci/run-integration-tests.sh
rename to ci/run-examples.sh
diff --git a/cmake/celerity-config.cmake.in b/cmake/celerity-config.cmake.in
index f2d3daf75..4851dad28 100644
--- a/cmake/celerity-config.cmake.in
+++ b/cmake/celerity-config.cmake.in
@@ -21,6 +21,9 @@ find_dependency(Threads REQUIRED)
 find_dependency(fmt REQUIRED)
 find_dependency(spdlog REQUIRED)
 find_dependency(small_vector REQUIRED)
+if(@CELERITY_ENABLE_CUDA_BACKEND@)
+  find_dependency(CUDAToolkit REQUIRED)
+endif()
 
 if(CELERITY_SYCL_IMPL STREQUAL "hipSYCL")
   if(NOT DEFINED HIPSYCL_TARGETS AND NOT "@HIPSYCL_TARGETS@" STREQUAL "")
diff --git a/include/backend/backend.h b/include/backend/backend.h
new file mode 100644
index 000000000..eae91b500
--- /dev/null
+++ b/include/backend/backend.h
@@ -0,0 +1,56 @@
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+#include "backend/generic_backend.h"
+#include "backend/traits.h"
+#include "backend/type.h"
+
+// NOTE: These should not leak any symbols from the backend library (i.e. don't include it in the header)
+#if CELERITY_DETAIL_BACKEND_CUDA_ENABLED
+#include "backend/cuda_backend.h"
+#endif
+
+namespace celerity::detail::backend_detail {
+template <template <backend::type> typename Template, typename Callback>
+auto specialize_for_backend(backend::type type, Callback cb) {
+	switch(type) {
+	case backend::type::cuda: return cb(Template<backend::type::cuda>{});
+	case backend::type::generic: return cb(Template<backend::type::generic>{});
+	case backend::type::unknown: [[fallthrough]];
+	default: return cb(Template<backend::type::unknown>{});
+	}
+}
+} // namespace celerity::detail::backend_detail
+
+namespace celerity::detail::backend {
+
+/**
+ * Returns the detected backend type for this SYCL device.
+ *
+ * Returns either a specialized backend or 'unknown', never 'generic'.
+ */
+type get_type(const sycl::device& device);
+
+/**
+ * Returns the effective backend type for this SYCL device, depending on the detected
+ * backend type and which backend modules have been compiled.
+ *
+ * Returns either a specialized backend or 'generic', never 'unknown'.
+ */
+type get_effective_type(const sycl::device& device);
+
+inline std::string_view get_name(type type) {
+	return backend_detail::specialize_for_backend<backend_detail::name>(type, [](auto op) { return decltype(op)::value; });
+}
+
+template <int Dims>
+void memcpy_strided_device(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<Dims>& source_range,
+    const sycl::id<Dims>& source_offset, const sycl::range<Dims>& target_range, const sycl::id<Dims>& target_offset, const sycl::range<Dims>& copy_range) {
+	backend_detail::specialize_for_backend<backend_detail::backend_operations>(get_effective_type(queue.get_device()), [&](auto op) {
+		decltype(op)::memcpy_strided_device(
+		    queue, source_base_ptr, target_base_ptr, elem_size, source_range, source_offset, target_range, target_offset, copy_range);
+	});
+}
+
+} // namespace celerity::detail::backend
\ No newline at end of file
diff --git a/include/backend/cuda_backend.h b/include/backend/cuda_backend.h
new file mode 100644
index 000000000..1ac62ac55
--- /dev/null
+++ b/include/backend/cuda_backend.h
@@ -0,0 +1,27 @@
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+#include "backend/operations.h"
+#include "backend/type.h"
+
+namespace celerity::detail::backend_detail {
+
+void memcpy_strided_device_cuda(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<1>& source_range,
+    const sycl::id<1>& source_offset, const sycl::range<1>& target_range, const sycl::id<1>& target_offset, const sycl::range<1>& copy_range);
+
+void memcpy_strided_device_cuda(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<2>& source_range,
+    const sycl::id<2>& source_offset, const sycl::range<2>& target_range, const sycl::id<2>& target_offset, const sycl::range<2>& copy_range);
+
+void memcpy_strided_device_cuda(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<3>& source_range,
+    const sycl::id<3>& source_offset, const sycl::range<3>& target_range, const sycl::id<3>& target_offset, const sycl::range<3>& copy_range);
+
+template <>
+struct backend_operations<backend::type::cuda> {
+	template <typename... Args>
+	static void memcpy_strided_device(Args&&... args) {
+		memcpy_strided_device_cuda(args...);
+	}
+};
+
+} // namespace celerity::detail::backend_detail
diff --git a/include/backend/generic_backend.h b/include/backend/generic_backend.h
new file mode 100644
index 000000000..3436efac1
--- /dev/null
+++ b/include/backend/generic_backend.h
@@ -0,0 +1,27 @@
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+#include "backend/operations.h"
+#include "backend/type.h"
+
+namespace celerity::detail::backend_detail {
+
+void memcpy_strided_device_generic(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<1>& source_range,
+    const sycl::id<1>& source_offset, const sycl::range<1>& target_range, const sycl::id<1>& target_offset, const sycl::range<1>& copy_range);
+
+void memcpy_strided_device_generic(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<2>& source_range,
+    const sycl::id<2>& source_offset, const sycl::range<2>& target_range, const sycl::id<2>& target_offset, const sycl::range<2>& copy_range);
+
+void memcpy_strided_device_generic(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<3>& source_range,
+    const sycl::id<3>& source_offset, const sycl::range<3>& target_range, const sycl::id<3>& target_offset, const sycl::range<3>& copy_range);
+
+template <>
+struct backend_operations<backend::type::generic> {
+	template <typename... Args>
+	static void memcpy_strided_device(Args&&... args) {
+		memcpy_strided_device_generic(args...);
+	}
+};
+
+} // namespace celerity::detail::backend_detail
diff --git a/include/backend/operations.h b/include/backend/operations.h
new file mode 100644
index 000000000..aa1e25552
--- /dev/null
+++ b/include/backend/operations.h
@@ -0,0 +1,17 @@
+#pragma once
+
+#include <stdexcept>
+
+#include "backend/type.h"
+
+namespace celerity::detail::backend_detail {
+
+template <backend::type Type>
+struct backend_operations {
+	template <typename... Args>
+	static void memcpy_strided_device(Args&&... args) {
+		throw std::runtime_error{"Invalid backend"};
+	}
+};
+
+} // namespace celerity::detail::backend_detail
\ No newline at end of file
diff --git a/include/backend/traits.h b/include/backend/traits.h
new file mode 100644
index 000000000..29d0fc588
--- /dev/null
+++ b/include/backend/traits.h
@@ -0,0 +1,41 @@
+#pragma once
+
+#include <type_traits>
+
+#include "backend/type.h"
+
+namespace celerity::detail::backend_detail {
+
+template <backend::type Type>
+struct is_enabled : public std::false_type {};
+
+template <backend::type Type>
+constexpr bool is_enabled_v = is_enabled<Type>::value;
+
+template <backend::type Type>
+struct name {
+	static constexpr const char* value = "(unknown)";
+};
+
+template <backend::type Type>
+constexpr const char* name_v = name<Type>::value;
+
+template <>
+struct is_enabled<backend::type::generic> : public std::true_type {};
+
+template <>
+struct name<backend::type::generic> {
+	static constexpr const char* value = "generic";
+};
+
+#if CELERITY_DETAIL_BACKEND_CUDA_ENABLED
+template <>
+struct is_enabled<backend::type::cuda> : public std::true_type {};
+#endif
+
+template <>
+struct name<backend::type::cuda> {
+	static constexpr const char* value = "CUDA";
+};
+
+} // namespace celerity::detail::backend_detail
\ No newline at end of file
diff --git a/include/backend/type.h b/include/backend/type.h
new file mode 100644
index 000000000..b1c03468c
--- /dev/null
+++ b/include/backend/type.h
@@ -0,0 +1,5 @@
+#pragma once
+
+namespace celerity::detail::backend {
+enum class type { generic, cuda, unknown };
+} // namespace celerity::detail::backend
diff --git a/include/buffer_storage.h b/include/buffer_storage.h
index 72b789abe..e1504d5dc 100644
--- a/include/buffer_storage.h
+++ b/include/buffer_storage.h
@@ -6,6 +6,7 @@
 
 #include <CL/sycl.hpp>
 
+#include "backend/backend.h"
 #include "payload.h"
 #include "ranges.h"
 #include "workaround.h"
@@ -25,19 +26,6 @@ namespace detail {
 	    const cl::sycl::id<3>& source_offset, const cl::sycl::range<3>& target_range, const cl::sycl::id<3>& target_offset,
 	    const cl::sycl::range<3>& copy_range);
 
-	// SYCL 2020 doesn't include any strided overloads for memcpy, so much like on the host, we have to roll our own.
-	void memcpy_strided_device(cl::sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size,
-	    const cl::sycl::range<1>& source_range, const cl::sycl::id<1>& source_offset, const cl::sycl::range<1>& target_range,
-	    const cl::sycl::id<1>& target_offset, const cl::sycl::range<1>& copy_range);
-
-	void memcpy_strided_device(cl::sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size,
-	    const cl::sycl::range<2>& source_range, const cl::sycl::id<2>& source_offset, const cl::sycl::range<2>& target_range,
-	    const cl::sycl::id<2>& target_offset, const cl::sycl::range<2>& copy_range);
-
-	void memcpy_strided_device(cl::sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size,
-	    const cl::sycl::range<3>& source_range, const cl::sycl::id<3>& source_offset, const cl::sycl::range<3>& target_range,
-	    const cl::sycl::id<3>& target_offset, const cl::sycl::range<3>& copy_range);
-
 	void linearize_subrange(const void* source_base_ptr, void* target_ptr, size_t elem_size, const range<3>& source_range, const subrange<3>& copy_sr);
 
 	template <typename DataT, int Dims>
@@ -155,8 +143,8 @@ namespace detail {
 			assert_copy_is_in_range(range_cast<3>(m_device_buf.get_range()), sr.range, sr.offset, id<3>{}, sr.range);
 
 			// TODO: Ideally we'd make this non-blocking and return some sort of async handle that can be waited upon
-			memcpy_strided_device(m_owning_queue, m_device_buf.get_pointer(), out_linearized, sizeof(DataT), m_device_buf.get_range(), id_cast<Dims>(sr.offset),
-			    range_cast<Dims>(sr.range), id<Dims>{}, range_cast<Dims>(sr.range));
+			backend::memcpy_strided_device(m_owning_queue, m_device_buf.get_pointer(), out_linearized, sizeof(DataT), m_device_buf.get_range(),
+			    id_cast<Dims>(sr.offset), range_cast<Dims>(sr.range), id<Dims>{}, range_cast<Dims>(sr.range));
 		}
 
 		void set_data(const subrange<3>& sr, const void* in_linearized) override {
@@ -165,7 +153,7 @@ namespace detail {
 			assert_copy_is_in_range(sr.range, range_cast<3>(m_device_buf.get_range()), id<3>{}, sr.offset, sr.range);
 
 			// TODO: Ideally we'd make this non-blocking and return some sort of async handle that can be waited upon
-			memcpy_strided_device(m_owning_queue, in_linearized, m_device_buf.get_pointer(), sizeof(DataT), range_cast<Dims>(sr.range), id<Dims>{},
+			backend::memcpy_strided_device(m_owning_queue, in_linearized, m_device_buf.get_pointer(), sizeof(DataT), range_cast<Dims>(sr.range), id<Dims>{},
 			    m_device_buf.get_range(), id_cast<Dims>(sr.offset), range_cast<Dims>(sr.range));
 		}
 
@@ -222,7 +210,7 @@ namespace detail {
 
 		if(source.get_type() == buffer_type::device_buffer) {
 			auto& device_source = dynamic_cast<const device_buffer_storage<DataT, Dims>&>(source);
-			memcpy_strided_device(m_owning_queue, device_source.m_device_buf.get_pointer(), m_device_buf.get_pointer(), sizeof(DataT),
+			backend::memcpy_strided_device(m_owning_queue, device_source.m_device_buf.get_pointer(), m_device_buf.get_pointer(), sizeof(DataT),
 			    device_source.m_device_buf.get_range(), id_cast<Dims>(source_offset), m_device_buf.get_range(), id_cast<Dims>(target_offset),
 			    range_cast<Dims>(copy_range));
 		}
diff --git a/include/device_queue.h b/include/device_queue.h
index 6f6f7f6df..6a5c3af33 100644
--- a/include/device_queue.h
+++ b/include/device_queue.h
@@ -7,6 +7,7 @@
 
 #include <CL/sycl.hpp>
 
+#include "backend/backend.h"
 #include "config.h"
 #include "log.h"
 #include "workaround.h"
@@ -245,6 +246,21 @@ namespace detail {
 		const auto device_name = device.template get_info<sycl::info::device::name>();
 		CELERITY_INFO("Using platform '{}', device '{}' ({})", platform_name, device_name, how_selected);
 
+		if constexpr(std::is_same_v<DeviceT, sycl::device>) {
+			if(backend::get_effective_type(device) == backend::type::generic) {
+				if(backend::get_type(device) == backend::type::unknown) {
+					CELERITY_WARN("No backend specialization available for selected platform '{}', falling back to generic. Performance may be degraded.",
+					    device.get_platform().template get_info<sycl::info::platform::name>());
+				} else {
+					CELERITY_WARN("Selected platform '{}' is compatible with specialized {} backend, but it has not been compiled.",
+					    device.get_platform().template get_info<sycl::info::platform::name>(), backend::get_name(backend::get_type(device)));
+				}
+			} else {
+				CELERITY_DEBUG("Using {} backend for selected platform '{}'.", backend::get_name(backend::get_effective_type(device)),
+				    device.get_platform().template get_info<sycl::info::platform::name>());
+			}
+		}
+
 		return device;
 	}
 
diff --git a/src/backend/CMakeLists.txt b/src/backend/CMakeLists.txt
new file mode 100644
index 000000000..c1c275ff9
--- /dev/null
+++ b/src/backend/CMakeLists.txt
@@ -0,0 +1,21 @@
+find_package(CUDAToolkit QUIET)
+# find_package(LevelZero QUIET) # TODO: Need find module?
+# find_package(ROCM QUIET) # TODO: Need find module?
+
+# TODO: Make conditional on CELERITY_SYCL_IMPL?
+option(CELERITY_ENABLE_CUDA_BACKEND "Enable optimized code paths for CUDA backends" ${CUDAToolkit_FOUND})
+if(CELERITY_ENABLE_CUDA_BACKEND AND NOT CUDAToolkit_FOUND)
+  # Run find_package again to emit error message
+  find_package(CUDAToolkit REQUIRED)
+endif()
+
+add_library(celerity_backends STATIC backend.cc generic_backend.cc)
+target_include_directories(celerity_backends PRIVATE ${CMAKE_SOURCE_DIR}/include)
+add_sycl_to_target(TARGET celerity_backends SOURCES)
+
+if(CELERITY_ENABLE_CUDA_BACKEND)
+  target_sources(celerity_backends PRIVATE cuda_backend.cc)
+  target_link_libraries(celerity_backends PUBLIC CUDA::cudart)
+  target_compile_definitions(celerity_backends PUBLIC "CELERITY_DETAIL_BACKEND_CUDA_ENABLED=1")
+  message(STATUS "CUDA backend enabled")
+endif()
diff --git a/src/backend/backend.cc b/src/backend/backend.cc
new file mode 100644
index 000000000..a4a6a377b
--- /dev/null
+++ b/src/backend/backend.cc
@@ -0,0 +1,25 @@
+#include "backend/backend.h"
+
+namespace celerity::detail::backend {
+
+type get_type(const sycl::device& device) {
+#if defined(__HIPSYCL__) && defined(SYCL_EXT_HIPSYCL_BACKEND_CUDA)
+	if(device.get_backend() == sycl::backend::cuda) { return type::cuda; }
+#endif
+#if defined(__SYCL_COMPILER_VERSION) // DPC++ (TODO: This may break when using OpenSYCL w/ DPC++ as compiler)
+	if(device.get_backend() == sycl::backend::ext_oneapi_cuda) { return type::cuda; }
+#endif
+	return type::unknown;
+}
+
+type get_effective_type(const sycl::device& device) {
+	[[maybe_unused]] const auto b = get_type(device);
+
+#if defined(CELERITY_DETAIL_BACKEND_CUDA_ENABLED)
+	if(b == type::cuda) return b;
+#endif
+
+	return type::generic;
+}
+
+} // namespace celerity::detail::backend
diff --git a/src/backend/cuda_backend.cc b/src/backend/cuda_backend.cc
new file mode 100644
index 000000000..a09959196
--- /dev/null
+++ b/src/backend/cuda_backend.cc
@@ -0,0 +1,49 @@
+#include "backend/cuda_backend.h"
+
+#include <cuda_runtime.h>
+
+#include "ranges.h"
+
+namespace celerity::detail::backend_detail {
+
+void memcpy_strided_device_cuda(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<1>& source_range,
+    const sycl::id<1>& source_offset, const sycl::range<1>& target_range, const sycl::id<1>& target_offset, const sycl::range<1>& copy_range) {
+	(void)queue;
+	const size_t line_size = elem_size * copy_range[0];
+	const auto ret = cudaMemcpy(static_cast<char*>(target_base_ptr) + elem_size * get_linear_index(target_range, target_offset),
+	    static_cast<const char*>(source_base_ptr) + elem_size * get_linear_index(source_range, source_offset), line_size, cudaMemcpyDefault);
+	if(ret != cudaSuccess) throw std::runtime_error("cudaMemcpy failed");
+	// Classic CUDA footgun: Memcpy is not always synchronous (e.g. for D2D)
+	cudaStreamSynchronize(0);
+}
+
+void memcpy_strided_device_cuda(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<2>& source_range,
+    const sycl::id<2>& source_offset, const sycl::range<2>& target_range, const sycl::id<2>& target_offset, const sycl::range<2>& copy_range) {
+	(void)queue;
+	const auto source_base_offset = get_linear_index(source_range, source_offset);
+	const auto target_base_offset = get_linear_index(target_range, target_offset);
+	const auto ret = cudaMemcpy2D(static_cast<char*>(target_base_ptr) + elem_size * target_base_offset, target_range[1] * elem_size,
+	    static_cast<const char*>(source_base_ptr) + elem_size * source_base_offset, source_range[1] * elem_size, copy_range[1] * elem_size, copy_range[0],
+	    cudaMemcpyDefault);
+	if(ret != cudaSuccess) throw std::runtime_error("cudaMemcpy2D failed");
+	// Classic CUDA footgun: Memcpy is not always synchronous (e.g. for D2D)
+	cudaStreamSynchronize(0);
+}
+
+void memcpy_strided_device_cuda(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<3>& source_range,
+    const sycl::id<3>& source_offset, const sycl::range<3>& target_range, const sycl::id<3>& target_offset, const sycl::range<3>& copy_range) {
+	cudaMemcpy3DParms parms = {};
+	parms.srcPos = make_cudaPos(source_offset[2] * elem_size, source_offset[1], source_offset[0]);
+	parms.srcPtr = make_cudaPitchedPtr(
+	    const_cast<void*>(source_base_ptr), source_range[2] * elem_size, source_range[2], source_range[1]); // NOLINT cppcoreguidelines-pro-type-const-cast
+	parms.dstPos = make_cudaPos(target_offset[2] * elem_size, target_offset[1], target_offset[0]);
+	parms.dstPtr = make_cudaPitchedPtr(target_base_ptr, target_range[2] * elem_size, target_range[2], target_range[1]);
+	parms.extent = {copy_range[2] * elem_size, copy_range[1], copy_range[0]};
+	parms.kind = cudaMemcpyDefault;
+	const auto ret = cudaMemcpy3D(&parms);
+	if(ret != cudaSuccess) throw std::runtime_error("cudaMemcpy3D failed");
+	// Classic CUDA footgun: Memcpy is not always synchronous (e.g. for D2D)
+	cudaStreamSynchronize(0);
+}
+
+} // namespace celerity::detail::backend_detail
\ No newline at end of file
diff --git a/src/backend/generic_backend.cc b/src/backend/generic_backend.cc
new file mode 100644
index 000000000..2956c816d
--- /dev/null
+++ b/src/backend/generic_backend.cc
@@ -0,0 +1,48 @@
+#include "backend/generic_backend.h"
+
+#include "ranges.h"
+
+namespace celerity::detail::backend_detail {
+
+void memcpy_strided_device_generic(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<1>& source_range,
+    const sycl::id<1>& source_offset, const sycl::range<1>& target_range, const sycl::id<1>& target_offset, const sycl::range<1>& copy_range) {
+	const size_t line_size = elem_size * copy_range[0];
+	auto evt = queue.memcpy(static_cast<char*>(target_base_ptr) + elem_size * get_linear_index(target_range, target_offset),
+	    static_cast<const char*>(source_base_ptr) + elem_size * get_linear_index(source_range, source_offset), line_size);
+	evt.wait();
+}
+
+// TODO Optimize for contiguous copies?
+void memcpy_strided_device_generic(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<2>& source_range,
+    const sycl::id<2>& source_offset, const sycl::range<2>& target_range, const sycl::id<2>& target_offset, const sycl::range<2>& copy_range) {
+	const auto source_base_offset = get_linear_index(source_range, source_offset);
+	const auto target_base_offset = get_linear_index(target_range, target_offset);
+	const size_t line_size = elem_size * copy_range[1];
+	std::vector<sycl::event> wait_list;
+	wait_list.reserve(copy_range[0]);
+	for(size_t i = 0; i < copy_range[0]; ++i) {
+		auto e = queue.memcpy(static_cast<char*>(target_base_ptr) + elem_size * (target_base_offset + i * target_range[1]),
+		    static_cast<const char*>(source_base_ptr) + elem_size * (source_base_offset + i * source_range[1]), line_size);
+		wait_list.push_back(e);
+	}
+	sycl::event::wait(wait_list);
+}
+
+// TODO Optimize for contiguous copies?
+void memcpy_strided_device_generic(sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size, const sycl::range<3>& source_range,
+    const sycl::id<3>& source_offset, const sycl::range<3>& target_range, const sycl::id<3>& target_offset, const sycl::range<3>& copy_range) {
+	// We simply decompose this into a bunch of 2D copies. Subtract offset on the copy plane, as it will be added again during the 2D copy.
+	const auto source_base_offset =
+	    get_linear_index(source_range, source_offset) - get_linear_index({source_range[1], source_range[2]}, {source_offset[1], source_offset[2]});
+	const auto target_base_offset =
+	    get_linear_index(target_range, target_offset) - get_linear_index({target_range[1], target_range[2]}, {target_offset[1], target_offset[2]});
+
+	for(size_t i = 0; i < copy_range[0]; ++i) {
+		const auto* const source_ptr = static_cast<const char*>(source_base_ptr) + elem_size * (source_base_offset + i * (source_range[1] * source_range[2]));
+		auto* const target_ptr = static_cast<char*>(target_base_ptr) + elem_size * (target_base_offset + i * (target_range[1] * target_range[2]));
+		memcpy_strided_device_generic(queue, source_ptr, target_ptr, elem_size, {source_range[1], source_range[2]}, {source_offset[1], source_offset[2]},
+		    {target_range[1], target_range[2]}, {target_offset[1], target_offset[2]}, {copy_range[1], copy_range[2]});
+	}
+}
+
+} // namespace celerity::detail::backend_detail
\ No newline at end of file
diff --git a/src/buffer_storage.cc b/src/buffer_storage.cc
index 7d2c25f45..1c43b3a25 100644
--- a/src/buffer_storage.cc
+++ b/src/buffer_storage.cc
@@ -42,53 +42,6 @@ namespace detail {
 		}
 	}
 
-	void memcpy_strided_device(cl::sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size,
-	    const cl::sycl::range<1>& source_range, const cl::sycl::id<1>& source_offset, const cl::sycl::range<1>& target_range,
-	    const cl::sycl::id<1>& target_offset, const cl::sycl::range<1>& copy_range) {
-		const size_t line_size = elem_size * copy_range[0];
-		queue
-		    .memcpy(static_cast<char*>(target_base_ptr) + elem_size * get_linear_index(target_range, target_offset),
-		        static_cast<const char*>(source_base_ptr) + elem_size * get_linear_index(source_range, source_offset), line_size)
-		    .wait();
-	}
-
-	// TODO Optimize for contiguous copies?
-	void memcpy_strided_device(cl::sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size,
-	    const cl::sycl::range<2>& source_range, const cl::sycl::id<2>& source_offset, const cl::sycl::range<2>& target_range,
-	    const cl::sycl::id<2>& target_offset, const cl::sycl::range<2>& copy_range) {
-		const size_t line_size = elem_size * copy_range[1];
-		const auto source_base_offset = get_linear_index(source_range, source_offset);
-		const auto target_base_offset = get_linear_index(target_range, target_offset);
-		std::vector<cl::sycl::event> wait_list;
-		wait_list.reserve(copy_range[0]);
-		for(size_t i = 0; i < copy_range[0]; ++i) {
-			auto e = queue.memcpy(static_cast<char*>(target_base_ptr) + elem_size * (target_base_offset + i * target_range[1]),
-			    static_cast<const char*>(source_base_ptr) + elem_size * (source_base_offset + i * source_range[1]), line_size);
-			wait_list.push_back(e);
-		}
-		for(auto& e : wait_list) {
-			e.wait();
-		}
-	}
-
-	// TODO Optimize for contiguous copies?
-	void memcpy_strided_device(cl::sycl::queue& queue, const void* source_base_ptr, void* target_base_ptr, size_t elem_size,
-	    const cl::sycl::range<3>& source_range, const cl::sycl::id<3>& source_offset, const cl::sycl::range<3>& target_range,
-	    const cl::sycl::id<3>& target_offset, const cl::sycl::range<3>& copy_range) {
-		// We simply decompose this into a bunch of 2D copies. Subtract offset on the copy plane, as it will be added again during the 2D copy.
-		const auto source_base_offset =
-		    get_linear_index(source_range, source_offset) - get_linear_index({source_range[1], source_range[2]}, {source_offset[1], source_offset[2]});
-		const auto target_base_offset =
-		    get_linear_index(target_range, target_offset) - get_linear_index({target_range[1], target_range[2]}, {target_offset[1], target_offset[2]});
-		for(size_t i = 0; i < copy_range[0]; ++i) {
-			const auto* const source_ptr =
-			    static_cast<const char*>(source_base_ptr) + elem_size * (source_base_offset + i * (source_range[1] * source_range[2]));
-			auto* const target_ptr = static_cast<char*>(target_base_ptr) + elem_size * (target_base_offset + i * (target_range[1] * target_range[2]));
-			memcpy_strided_device(queue, source_ptr, target_ptr, elem_size, {source_range[1], source_range[2]}, {source_offset[1], source_offset[2]},
-			    {target_range[1], target_range[2]}, {target_offset[1], target_offset[2]}, {copy_range[1], copy_range[2]});
-		}
-	}
-
 	void linearize_subrange(const void* source_base_ptr, void* target_ptr, size_t elem_size, const range<3>& source_range, const subrange<3>& copy_sr) {
 		assert((id_cast<3>(copy_sr.offset) < id_cast<3>(source_range)) == cl::sycl::id<3>(1, 1, 1));
 		assert((id_cast<3>(copy_sr.offset + copy_sr.range) <= id_cast<3>(source_range)) == cl::sycl::id<3>(1, 1, 1));
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 6af6fb132..29fe973c1 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -25,6 +25,7 @@ file(GLOB_RECURSE TEST_INCLUDES *.h)
 
 set(TEST_TARGETS
   accessor_tests
+  backend_tests
   buffer_manager_tests
   graph_generation_tests
   graph_gen_granularity_tests
@@ -76,3 +77,6 @@ target_link_libraries(benchmarks PRIVATE test_main)
 set_test_target_parameters(benchmarks benchmarks.cc)
 
 add_subdirectory(system)
+if(CELERITY_DETAIL_INTEGRATION_TESTING)
+  add_subdirectory(integration)
+endif()
\ No newline at end of file
diff --git a/test/backend_tests.cc b/test/backend_tests.cc
new file mode 100644
index 000000000..3f31503f6
--- /dev/null
+++ b/test/backend_tests.cc
@@ -0,0 +1,157 @@
+#include <catch2/catch_template_test_macros.hpp>
+#include <catch2/catch_test_macros.hpp>
+#include <catch2/generators/catch_generators.hpp>
+#include <catch2/generators/catch_generators_range.hpp>
+
+#include <celerity.h>
+
+#include "test_utils.h"
+
+using namespace celerity;
+using namespace celerity::detail;
+
+// NOTE: We currently test backends on a "best effort" basis, that is we test whatever configuration(s) are available.
+// To exhaustively test all backend configurations we'd need full control over the hardware and software (SYCL implementation)
+// as well as a script to generate the necessary build configurations.
+
+template <int Dims>
+void write_linear_ids(size_t* ptr, const range<Dims>& buffer_range) {
+	const auto br3 = range_cast<3>(buffer_range);
+	for(size_t k = 0; k < br3[0]; ++k) {
+		for(size_t j = 0; j < br3[1]; ++j) {
+			for(size_t i = 0; i < br3[2]; ++i) {
+				const auto linear_id = get_linear_index(buffer_range, id_cast<Dims>(id<3>(k, j, i)));
+				ptr[linear_id] = linear_id;
+			}
+		}
+	}
+}
+
+template <int Dims>
+void verify_copied_linear_ids(const size_t* host_buf, const range<Dims>& source_range, const id<Dims>& source_offset, const range<Dims>& target_range,
+    const id<Dims>& target_offset, const range<Dims>& copy_range) {
+	const auto cr3 = range_cast<3>(copy_range);
+	for(size_t k = 0; k < cr3[0]; ++k) {
+		for(size_t j = 0; j < cr3[1]; ++j) {
+			for(size_t i = 0; i < cr3[2]; ++i) {
+				const auto src_idx = source_offset + id_cast<Dims>(id<3>(k, j, i));
+				const auto tgt_idx = target_offset + id_cast<Dims>(id<3>(k, j, i));
+				REQUIRE_LOOP(host_buf[get_linear_index(target_range, tgt_idx)] == get_linear_index(source_range, src_idx));
+			}
+		}
+	}
+}
+
+template <int Dims>
+struct copy_parameters {
+	range<Dims> source_range = range_cast<Dims>(range<3>(5, 7, 11));
+	range<Dims> target_range = range_cast<Dims>(range<3>(13, 17, 19));
+	range<Dims> copy_range = range_cast<Dims>(range<3>(2, 4, 8));
+	id<Dims> source_offset = id_cast<Dims>(id<3>(2, 2, 2));
+	id<Dims> target_offset = id_cast<Dims>(id<3>(3, 5, 7));
+};
+
+template <int Dims>
+std::vector<size_t> copy_to_host(sycl::queue& q, const size_t* device_ptr, const range<Dims>& buffer_range) {
+	std::vector<size_t> result(buffer_range.size());
+	q.memcpy(result.data(), device_ptr, buffer_range.size() * sizeof(size_t));
+	q.wait_and_throw();
+	return result;
+}
+
+enum class copy_test_type { intra_device, inter_device, host_to_device, device_to_host };
+
+struct host_or_device {
+	std::optional<sycl::queue> queue;
+	size_t* ptr = nullptr;
+
+	template <typename T>
+	void malloc(const size_t count) {
+		if(queue.has_value()) {
+			ptr = sycl::malloc_device<T>(count, *queue);
+		} else {
+			ptr = static_cast<T*>(std::malloc(count * sizeof(T))); // NOLINT cppcoreguidelines-no-malloc
+		}
+	}
+
+	template <typename Fn>
+	void run(const Fn fn) {
+		if(queue.has_value()) {
+			queue->single_task([=]() { fn(); });
+			queue->wait_and_throw();
+		} else {
+			fn();
+		}
+	}
+
+	~host_or_device() {
+		if(ptr != nullptr) {
+			if(queue.has_value()) {
+				sycl::free(ptr, *queue);
+			} else {
+				std::free(ptr); // NOLINT cppcoreguidelines-no-malloc
+			}
+		}
+	}
+};
+
+std::pair<host_or_device, host_or_device> select_source_and_target(const copy_test_type test_type, const sycl::platform& platform) {
+	const auto devices = platform.get_devices();
+	switch(test_type) {
+	case copy_test_type::intra_device: {
+		return std::pair{host_or_device{sycl::queue{devices[0]}}, host_or_device{sycl::queue{devices[0]}}};
+	}
+	case copy_test_type::inter_device: {
+		if(devices.size() < 2) {
+			throw std::runtime_error(fmt::format("Platform {} has less than two devices", platform.get_info<sycl::info::platform::name>()));
+		}
+		return std::pair{host_or_device{sycl::queue{devices[0]}}, host_or_device{sycl::queue{devices[1]}}};
+	}
+	case copy_test_type::host_to_device: {
+		return std::pair{host_or_device{std::nullopt}, host_or_device{sycl::queue{devices[0]}}};
+	}
+	case copy_test_type::device_to_host: {
+		return std::pair{host_or_device{sycl::queue{devices[0]}}, host_or_device{std::nullopt}};
+	}
+	}
+}
+
+TEMPLATE_TEST_CASE_SIG("memcpy_strided_device allows to copy between the same device, different devices on the same platform and between host and device",
+    "[backend]", ((int Dims), Dims), 1, 2, 3) {
+	const copy_parameters<Dims> cp;
+	const size_t platform_id = GENERATE(Catch::Generators::range(size_t(0), sycl::platform::get_platforms().size()));
+	const auto test_type = GENERATE(copy_test_type::intra_device, copy_test_type::inter_device, copy_test_type::host_to_device, copy_test_type::device_to_host);
+	const auto platform = sycl::platform::get_platforms()[platform_id];
+	auto [src, tgt] = ([&] {
+		try {
+			return select_source_and_target(test_type, platform);
+		} catch(std::runtime_error& e) {
+			SKIP(e.what());
+			abort(); // Will never be reached, just to shut up compiler about not all control paths returning a value
+		}
+	})();
+
+	src.template malloc<size_t>(cp.source_range.size());
+	tgt.template malloc<size_t>(cp.target_range.size());
+
+	src.run([ptr = src.ptr, cp = cp]() { write_linear_ids(ptr, cp.source_range); });
+
+	const auto get_a_queue = [](host_or_device& a, host_or_device& b) -> sycl::queue& {
+		if(a.queue.has_value()) return *a.queue;
+		return *b.queue;
+	};
+
+	// Note: This may also be the generic backend
+	SECTION("using automatically selected backend") {
+		backend::memcpy_strided_device(
+		    get_a_queue(src, tgt), src.ptr, tgt.ptr, sizeof(size_t), cp.source_range, cp.source_offset, cp.target_range, cp.target_offset, cp.copy_range);
+	}
+
+	SECTION("using generic backend") {
+		backend_detail::backend_operations<backend::type::generic>::memcpy_strided_device(
+		    get_a_queue(src, tgt), src.ptr, tgt.ptr, sizeof(size_t), cp.source_range, cp.source_offset, cp.target_range, cp.target_offset, cp.copy_range);
+	}
+
+	const auto host_buf = copy_to_host(get_a_queue(tgt, src), tgt.ptr, cp.target_range);
+	verify_copied_linear_ids(host_buf.data(), cp.source_range, cp.source_offset, cp.target_range, cp.target_offset, cp.copy_range);
+}
\ No newline at end of file
diff --git a/test/device_selection_tests.cc b/test/device_selection_tests.cc
index bd45c2943..07ad9dd23 100644
--- a/test/device_selection_tests.cc
+++ b/test/device_selection_tests.cc
@@ -509,3 +509,57 @@ TEST_CASE_METHOD(celerity::test_utils::mpi_fixture, "pick_device does not consid
 	CHECK_THROWS_WITH(
 	    pick_device(cfg, device_selector, std::vector<mock_platform>{mp_0, mp_1, mp_2}), "Device selection with device selector failed: No device available");
 }
+
+TEST_CASE_METHOD(celerity::test_utils::mpi_fixture, "pick_device prints information about device backend") {
+	using namespace celerity;
+	using namespace celerity::detail;
+
+	config cfg(nullptr, nullptr);
+
+	const auto devices = sycl::device::get_devices();
+	std::optional<sycl::device> generic_device;
+	std::optional<sycl::device> cuda_device;
+
+	for(const auto& d : devices) {
+		if(backend::get_type(d) == backend::type::cuda) {
+			cuda_device = d;
+		} else {
+			generic_device = d;
+		}
+	}
+
+	test_utils::log_capture lc(spdlog::level::debug);
+
+	SECTION("warns when using generic backend") {
+		if(!generic_device.has_value()) {
+			SKIP("No generic device available");
+		} else {
+			pick_device(cfg, *generic_device, std::vector<sycl::platform>{});
+			CHECK_THAT(lc.get_log(),
+			    Catch::Matchers::ContainsSubstring(
+			        fmt::format("No backend specialization available for selected platform '{}', falling back to generic. Performance may be degraded.",
+			            generic_device->get_platform().get_info<sycl::info::platform::name>())));
+		}
+	}
+
+	SECTION("informs user when using specialized backend") {
+		if(!cuda_device.has_value() || !backend_detail::is_enabled_v<backend::type::cuda>) {
+			SKIP("No CUDA device available or CUDA backend not enabled");
+		} else {
+			pick_device(cfg, *cuda_device, std::vector<sycl::platform>{});
+			CHECK_THAT(lc.get_log(), Catch::Matchers::ContainsSubstring(fmt::format("Using CUDA backend for selected platform '{}'.",
+			                             cuda_device->get_platform().get_info<sycl::info::platform::name>())));
+		}
+	}
+
+	SECTION("warns when specialized backend is not enabled") {
+		if(!cuda_device.has_value() || backend_detail::is_enabled_v<backend::type::cuda>) {
+			SKIP("No CUDA device available or CUDA backend is enabled");
+		} else {
+			pick_device(cfg, *cuda_device, std::vector<sycl::platform>{});
+			CHECK_THAT(lc.get_log(), Catch::Matchers::ContainsSubstring(
+			                             fmt::format("Selected platform '{}' is compatible with specialized CUDA backend, but it has not been compiled.",
+			                                 cuda_device->get_platform().get_info<sycl::info::platform::name>())));
+		}
+	}
+}
diff --git a/test/integration/CMakeLists.txt b/test/integration/CMakeLists.txt
new file mode 100644
index 000000000..684447331
--- /dev/null
+++ b/test/integration/CMakeLists.txt
@@ -0,0 +1,18 @@
+if(NOT CELERITY_DETAIL_INTEGRATION_TESTING)
+  message(FATAL "Integration tests should not be compiled directly; use Python wrapper script instead")
+endif()
+
+function(add_integration_test TARGET SOURCE)
+  add_executable(${TARGET} ${SOURCE})
+  set_property(TARGET ${TARGET} PROPERTY CXX_STANDARD 17)
+
+  if(MSVC)
+    target_compile_options(${TARGET} PRIVATE /D_CRT_SECURE_NO_WARNINGS /MP /W3 /bigobj)
+  elseif(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang|AppleClang")
+    target_compile_options(${TARGET} PRIVATE -Wall -Wextra -Wextra -Wno-unused-parameter)
+  endif()
+
+  add_celerity_to_target(TARGET ${TARGET} SOURCES ${SOURCE})
+endfunction()
+
+add_integration_test(backend backend.cc)
\ No newline at end of file
diff --git a/test/integration/backend.cc b/test/integration/backend.cc
new file mode 100644
index 000000000..353f95715
--- /dev/null
+++ b/test/integration/backend.cc
@@ -0,0 +1,90 @@
+#include <celerity.h>
+
+#define ASSERT(...)                                                                                                                                            \
+	if(!(__VA_ARGS__)) {                                                                                                                                       \
+		fprintf(stderr, "ASSERT(%s) on line %d failed\n", #__VA_ARGS__, __LINE__);                                                                             \
+		std::_Exit(EXIT_FAILURE);                                                                                                                              \
+	}
+
+template <int Dims>
+void test_copy(celerity::distr_queue& q) {
+	celerity::buffer<size_t, Dims> buf(celerity::detail::range_cast<Dims>(celerity::range<3>{5, 7, 9}));
+
+	// Initialize on device
+	q.submit([=](celerity::handler& cgh) {
+		celerity::accessor acc{buf, cgh, celerity::access::one_to_one<>{}, celerity::write_only, celerity::no_init};
+		cgh.parallel_for(buf.get_range(), [=](celerity::item<Dims> itm) { acc[itm] = itm.get_linear_id(); });
+	});
+
+	// Check and modify partially on host
+	const auto sr = celerity::detail::subrange_cast<Dims>(celerity::subrange<3>{{1, 2, 3}, {3, 4, 5}});
+	const auto sr3 = celerity::detail::subrange_cast<3>(sr);
+	q.submit([=](celerity::handler& cgh) {
+		celerity::accessor acc{buf, cgh, celerity::access::fixed<Dims>{sr}, celerity::read_write_host_task};
+		cgh.host_task(celerity::on_master_node, [=]() {
+			for(size_t k = 0; k < sr3.range[0]; ++k) {
+				for(size_t j = 0; j < sr3.range[1]; ++j) {
+					for(size_t i = 0; i < sr3.range[2]; ++i) {
+						const celerity::id<3> idx{sr3.offset[0] + k, sr3.offset[1] + j, sr3.offset[2] + i};
+						const auto linear_id = celerity::detail::get_linear_index(buf.get_range(), celerity::detail::id_cast<Dims>(idx));
+						ASSERT(acc[celerity::detail::id_cast<Dims>(idx)] == linear_id);
+						acc[celerity::detail::id_cast<Dims>(idx)] *= 2;
+					}
+				}
+			}
+		});
+	});
+
+	// Modify everything on device
+	q.submit([=](celerity::handler& cgh) {
+		celerity::accessor acc{buf, cgh, celerity::access::one_to_one<>{}, celerity::read_write};
+		cgh.parallel_for(buf.get_range(), [=](celerity::item<Dims> itm) { acc[itm] += 1; });
+	});
+
+	// Check everything on host
+	q.submit([=](celerity::handler& cgh) {
+		celerity::accessor acc{buf, cgh, celerity::access::all{}, celerity::read_only_host_task};
+		cgh.host_task(celerity::on_master_node, [=]() {
+			const auto r3 = celerity::detail::range_cast<3>(buf.get_range());
+			for(size_t k = 0; k < r3[0]; ++k) {
+				for(size_t j = 0; j < r3[1]; ++j) {
+					for(size_t i = 0; i < r3[2]; ++i) {
+						const celerity::id<3> idx{k, j, i};
+						const auto is_in_sr =
+						    ((idx >= sr3.offset == celerity::id<3>(true, true, true)) && (idx < sr3.offset + sr3.range == celerity::id<3>(true, true, true)));
+						const auto linear_id = celerity::detail::get_linear_index(buf.get_range(), celerity::detail::id_cast<Dims>(idx));
+						if(is_in_sr) {
+							ASSERT(acc[celerity::detail::id_cast<Dims>(idx)] == 2 * linear_id + 1);
+						} else {
+							ASSERT(acc[celerity::detail::id_cast<Dims>(idx)] == linear_id + 1);
+						}
+					}
+				}
+			}
+		});
+	});
+}
+
+int main(int argc, char* argv[]) {
+	if(argc > 2) {
+		fmt::print(stderr, "Usage: %s [device index]\n", argv[0]);
+		return EXIT_FAILURE;
+	}
+
+	const auto all_devices = sycl::device::get_devices();
+	if(argc == 1) {
+		for(const auto& d : all_devices) {
+			fmt::print("{} {}\n", d.get_platform().get_info<sycl::info::platform::name>(), d.get_info<sycl::info::device::name>());
+		}
+		return 0;
+	}
+
+	const auto device_idx = std::atoi(argv[1]);
+	celerity::distr_queue q{all_devices[device_idx]};
+
+	test_copy<1>(q);
+	test_copy<2>(q);
+	test_copy<3>(q);
+
+	return EXIT_SUCCESS;
+}
\ No newline at end of file
diff --git a/test/integration/run-integration-tests.py b/test/integration/run-integration-tests.py
new file mode 100755
index 000000000..fc917b70c
--- /dev/null
+++ b/test/integration/run-integration-tests.py
@@ -0,0 +1,90 @@
+#!/usr/bin/env python3
+
+import os
+import subprocess
+import sys
+from typing import Literal
+
+
+def run_backend_discovery_and_copy_test(build_dir: str, backend: Literal['CUDA'] | None, enabled: bool):
+    script_dir = os.path.dirname(os.path.abspath(__file__))
+    celerity_dir = os.path.realpath(os.path.join(script_dir, "../.."))
+
+    # Configure and build
+    # TODO: Make sure only one backend is enabled at a time (once we have more than one)
+    backend_config = [f"-DCELERITY_ENABLE_CUDA_BACKEND={'ON' if backend == 'CUDA' and enabled else 'OFF'}"]
+    cmake_args = ["-DCELERITY_DETAIL_INTEGRATION_TESTING=ON"] + backend_config
+    subprocess.run(["cmake", celerity_dir] + cmake_args, cwd=build_dir, check=True)
+    subprocess.run(["cmake", "--build", ".", "--target", "backend"], cwd=build_dir, check=True)
+
+    exe = os.path.join(build_dir, "test/integration/backend")
+
+    # Run once without arguments to get list of available devices
+    output = subprocess.check_output(exe).decode()
+    print("\nAvailable devices:\n", output)
+    devices = [line.strip() for line in output.split("\n") if line.strip()]
+
+    # Try to find a device for the backend as well as an unrelated device
+    backend_device_idx = None
+    other_device_idx = None
+    for i, device in enumerate(devices):
+        if (backend is not None and backend.lower() in device.lower()):
+            if backend_device_idx is None:
+                backend_device_idx = i
+        elif other_device_idx is None:
+            other_device_idx = i
+
+    if backend is not None and backend_device_idx is None:
+        raise RuntimeError(
+            f"No matching device found for backend {backend}")
+
+    env = os.environ.copy()
+    env["CELERITY_LOG_LEVEL"] = "debug"
+
+    # Check that optimized backend is available (or not) for matching device
+    if backend is not None:
+        output = subprocess.run([exe, str(backend_device_idx)], stdout=subprocess.PIPE,
+                                env=env, check=True).stdout.decode()
+        print(output)
+        if enabled:
+            if f"Using {backend} backend for selected platform" not in output:
+                raise RuntimeError(
+                    f"Optimized {backend} backend was not selected for device '{devices[backend_device_idx]}'")
+        else:
+            if f"is compatible with specialized {backend} backend, but it has not been compiled" not in output:
+                raise RuntimeError(f"Did not receive warning about {backend} not being available")
+
+    # Check that no optimized backend is available for other device (if one exists)
+    if other_device_idx is not None:
+        output = subprocess.run([exe, str(other_device_idx)], stdout=subprocess.PIPE,
+                                env=env, check=True).stdout.decode()
+        print(output)
+        if "No backend specialization available" not in output:
+            raise RuntimeError("Did not receive warning about no backend specialization being available")
+
+
+# We currently require an existing build directory to configure and compile
+# integration tests. This is a bit hacky but it avoids having to restructure
+# the entire CI setup...
+if __name__ == '__main__':
+    if len(sys.argv) != 3:
+        print(f"Usage: {sys.argv[0]} <celerity build directory> <platform>")
+        sys.exit(1)
+    build_dir = sys.argv[1]
+    platform = sys.argv[2]
+
+    # Using backend names like "CUDA" or "Level-Zero" might make more sense here,
+    # but the platform is again what we have available in our current CI setup.
+    if platform not in ("intel", "nvidia"):
+        print(f"Invalid platform '{platform}'; must be 'intel' or 'nvidia'")
+        sys.exit(1)
+
+    # Level-Zero is NYI.
+    backend = "CUDA" if platform == "nvidia" else None
+
+    # Run once with backend module enabled
+    run_backend_discovery_and_copy_test(build_dir, backend, True)
+
+    # And once with module disabled
+    if backend is not None:
+        run_backend_discovery_and_copy_test(build_dir, backend, False)