sycl: host-only copies go through the CPU - the Arc's copy engine hangs on pinned-host <-> pageable copies

A queue memcpy/memset with neither end in device memory (USM host pinned <-> pageable host) runs on the Arc's
copy engine, which hangs on it on a B70: dmesg "Engine reset: engine_class=bcs", and after repeated runs the
device is declared wedged (needs a reboot). conversation_snapshot_test's fill() into its pinned KV pools hit it
on every run (the urEventWait "hang" was the driver resetting the engine under the copy), and so would the
engine's conversation_kv_save / restore of a --kv-resident session, whose pools are pinned host memory.
transfer() and the test's fill / fill_after now memcpy / memset on the CPU when sycl::get_pointer_type says both
ends are host or pageable; copies with a device end stay on the queue.

sycl/src/core/conversation_snapshot.cpp

@@ -111,9 +111,21 @@ bool valid(const QsaState& st, const Layout& l, int64_t upto, std::string& error
     return true;
 }
 
+// SYCL port: a copy with neither end in device memory (pinned USM host <-> pageable host) goes through memcpy on the
+// CPU. Through the queue it runs on the Arc's copy engine, which hung on it (dmesg "Engine reset: engine_class=bcs";
+// repeated, the B70 was declared wedged): conversation_snapshot_test, and the parked KV of a --kv-resident session.
+inline bool host_only_copy(void* dst, const void* src) {
+    const sycl::context ctx = dpct::get_in_order_queue().get_context();
+    const auto host = [&](const void* p) {
+        const sycl::usm::alloc t = sycl::get_pointer_type(p, ctx);
+        return t == sycl::usm::alloc::host || t == sycl::usm::alloc::unknown;
+    };
+    return host(dst) && host(src);
+}
 bool transfer(void *dst, const void *src, size_t n, std::string &error) try {
     if (!n) return true;
     if (!src || !dst) { error = "conversation snapshot: missing state buffer"; return false; }
+    if (host_only_copy(dst, src)) { std::memcpy(dst, src, n); return true; }
     // Default handles both device allocations and device-mapped host pool aliases.
     const dpct::err0 e =
         DPCT_CHECK_ERROR(dpct::get_in_order_queue().memcpy(dst, src, n).wait());

sycl/src/core/conversation_snapshot_test.cpp

@@ -13,12 +13,24 @@
 #include <iterator>
 #include <string>
 #include <cstdlib>
+#include <cstring>
 #include <limits>
 #include <vector>
 
 using namespace strata::core;
 using namespace strata::kernels;
 
+// SYCL port: a copy with neither end in device memory (pinned USM host <-> pageable host) goes through memcpy on the
+// CPU. Through the queue it runs on the Arc's copy engine, which hung on it (dmesg "Engine reset: engine_class=bcs";
+// repeated, the B70 was declared wedged): conversation_snapshot_test, and the parked KV of a --kv-resident session.
+static bool host_only_copy(void* dst, const void* src) {
+    const sycl::context ctx = dpct::get_in_order_queue().get_context();
+    const auto host = [&](const void* p) {
+        const sycl::usm::alloc t = sycl::get_pointer_type(p, ctx);
+        return t == sycl::usm::alloc::host || t == sycl::usm::alloc::unknown;
+    };
+    return host(dst) && host(src);
+}
 namespace {
 int checks = 0;
 void check(bool ok, const char* label) {
@@ -107,7 +119,8 @@ struct Fixture {
         for (size_t i=0;i<sources.size();++i) {
             std::vector<uint8_t> data(sizes[i]);
             for (size_t j=0;j<data.size();++j) data[j]=(uint8_t)(salt+i*31+j*7+j/257);
-            if (!data.empty()) cuda_check(DPCT_CHECK_ERROR(
+            if (!data.empty() && host_only_copy(sources[i], data.data())) std::memcpy(sources[i], data.data(), data.size());
+            else if (!data.empty()) cuda_check(DPCT_CHECK_ERROR(
                 dpct::get_in_order_queue()
                     .memcpy(sources[i], data.data(), data.size())
                     .wait()));
@@ -117,7 +130,9 @@ struct Fixture {
         for (size_t i=0;i<sources.size();++i) {
             const size_t offset = i == 4 ? size_t(first_dirty/4)*g.idx_key_dim*4
                                          : (sizes[i]/size_t(state.max_cells))*size_t((first_dirty/4)*4);
-            if (sizes[i] > offset)
+            if (sizes[i] > offset && host_only_copy(sources[i], sources[i]))
+                std::memset(static_cast<uint8_t *>(sources[i]) + offset, salt, sizes[i] - offset);
+            else if (sizes[i] > offset)
                 cuda_check(DPCT_CHECK_ERROR(
                     dpct::get_in_order_queue()
                         .memset(static_cast<uint8_t *>(sources[i]) + offset,