From 061dddc50d7c5f420c5c873d3cdd7045544e1832 Mon Sep 17 00:00:00 2001
From: Ben Ashbaugh <ben.ashbaugh@intel.com>
Date: Mon, 3 Aug 2026 10:29:17 -0700
Subject: [PATCH] add a simple USM sample

---
 samples/CMakeLists.txt                       |  2 +-
 samples/dpcpp/usm/00_hellousm/CMakeLists.txt | 10 +++++
 samples/dpcpp/usm/00_hellousm/README.md      | 16 ++++++++
 samples/dpcpp/usm/00_hellousm/main.cpp       | 41 ++++++++++++++++++++
 samples/dpcpp/usm/CMakeLists.txt             |  1 +
 5 files changed, 69 insertions(+), 1 deletion(-)
 create mode 100644 samples/dpcpp/usm/00_hellousm/CMakeLists.txt
 create mode 100644 samples/dpcpp/usm/00_hellousm/README.md
 create mode 100644 samples/dpcpp/usm/00_hellousm/main.cpp

diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt
index 83bfe3d..c478533 100644
--- a/samples/CMakeLists.txt
+++ b/samples/CMakeLists.txt
@@ -26,7 +26,7 @@ function(add_sycl_sample)
 
     target_include_directories(${SYCL_SAMPLE_TARGET} PRIVATE ${SYCL_SAMPLE_INCLUDES})
 
-    target_compile_options(${SYCL_SAMPLE_TARGET} PRIVATE -fsycl -fsycl-unnamed-lambda)
+    target_compile_options(${SYCL_SAMPLE_TARGET} PRIVATE -fsycl -fsycl-unnamed-lambda -fsycl-range-rounding=disable)
     target_compile_options(${SYCL_SAMPLE_TARGET} PRIVATE ${SYCL_SAMPLE_ADDITIONAL_COMPILE_OPTIONS})
 
     # Passing -fsycl via target_link_libraries for older versions of CMake:
diff --git a/samples/dpcpp/usm/00_hellousm/CMakeLists.txt b/samples/dpcpp/usm/00_hellousm/CMakeLists.txt
new file mode 100644
index 0000000..e5171c0
--- /dev/null
+++ b/samples/dpcpp/usm/00_hellousm/CMakeLists.txt
@@ -0,0 +1,10 @@
+# Copyright (c) 2026 Ben Ashbaugh
+#
+# SPDX-License-Identifier: MIT
+
+add_sycl_sample(
+    TEST
+    NUMBER 00
+    TARGET hellousm
+    CATEGORY usm
+    SOURCES main.cpp)
diff --git a/samples/dpcpp/usm/00_hellousm/README.md b/samples/dpcpp/usm/00_hellousm/README.md
new file mode 100644
index 0000000..a07a0f5
--- /dev/null
+++ b/samples/dpcpp/usm/00_hellousm/README.md
@@ -0,0 +1,16 @@
+# hellousm
+
+## Sample Purpose
+
+This is an equivalent of the hellosycl sample, rewritten to use USM.
+
+It uses device USM, so it should be able to run on many devices.
+Because it uses device USM, though, it needs to explicitly copy memory between the host and the device.
+
+## Key APIs and Concepts
+
+This sample demonstrates how to allocate, write, and copy USM memory.
+
+## Command Line Options
+
+None.
diff --git a/samples/dpcpp/usm/00_hellousm/main.cpp b/samples/dpcpp/usm/00_hellousm/main.cpp
new file mode 100644
index 0000000..988f9eb
--- /dev/null
+++ b/samples/dpcpp/usm/00_hellousm/main.cpp
@@ -0,0 +1,41 @@
+/*
+// Copyright (c) 2026 Ben Ashbaugh
+//
+// SPDX-License-Identifier: MIT
+*/
+
+#include <sycl/sycl.hpp>
+#include <iostream>
+
+int main()
+{
+    const size_t array_size = 16;
+
+    sycl::queue q{sycl::property::queue::in_order()};
+
+    std::cout << "Hello from SYCL, with USM!\n";
+    std::cout << "Running on default SYCL device " << q.get_device().get_info<sycl::info::device::name>() << std::endl;
+
+    if (q.get_device().has(sycl::aspect::usm_device_allocations)) {
+        std::cout << "This device supports aspect::usm_device_allocations.\n";
+    } else {
+        std::cerr << "This device does not support aspect::usm_device_allocations. Exiting.\n";
+        return -1;
+    }
+
+    int* d_data = sycl::malloc_device<int>(array_size, q);
+    q.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> i) {
+        d_data[i] = i.get(0);
+    });
+
+    int h_data[array_size];
+    q.copy(d_data, h_data, array_size);
+    q.wait();
+
+    for( int i = 0; i < array_size; i++ )
+    {
+        std::cout << "data[" << i << "] = " << h_data[i] << std::endl;
+    }
+
+    return 0;
+}
diff --git a/samples/dpcpp/usm/CMakeLists.txt b/samples/dpcpp/usm/CMakeLists.txt
index b8cbadc..c012ba4 100644
--- a/samples/dpcpp/usm/CMakeLists.txt
+++ b/samples/dpcpp/usm/CMakeLists.txt
@@ -2,6 +2,7 @@
 #
 # SPDX-License-Identifier: MIT
 
+add_subdirectory( 00_hellousm )
 add_subdirectory( 00_usmqueries )
 
 add_subdirectory( 100_dmemhelloworld )