diff --git a/Cargo.lock b/Cargo.lock
index 18ca49e..f34410e 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -19,6 +19,20 @@ name = "bytemuck"
 version = "1.25.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "d6aedf8ae72766347502cf3cb4f41cf5e9cc37d28bee90f1fdaaae15f9cf9424"
+dependencies = [
+ "bytemuck_derive",
+]
+
+[[package]]
+name = "bytemuck_derive"
+version = "1.12.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fc0e56a716f1e132ff6bf4bdac1c944a3fcdc1cae65f70a4a2a1ac3b401d2d1f"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 3.0.3",
+]
 
 [[package]]
 name = "cc"
@@ -30,6 +44,12 @@ dependencies = [
  "shlex",
 ]
 
+[[package]]
+name = "cfg-if"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801"
+
 [[package]]
 name = "clap"
 version = "4.6.1"
@@ -67,6 +87,12 @@ dependencies = [
  "unicode-width",
 ]
 
+[[package]]
+name = "crunchy"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5"
+
 [[package]]
 name = "cxx"
 version = "1.0.194"
@@ -235,6 +261,18 @@ dependencies = [
  "slab",
 ]
 
+[[package]]
+name = "half"
+version = "2.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b"
+dependencies = [
+ "bytemuck",
+ "cfg-if",
+ "crunchy",
+ "zerocopy",
+]
+
 [[package]]
 name = "hashbrown"
 version = "0.17.1"
@@ -393,6 +431,7 @@ dependencies = [
  "allocator-api2",
  "bytemuck",
  "cxx",
+ "half",
  "pin-project",
  "sycl-rs-derive",
  "sycl-rs-sys",
@@ -575,3 +614,23 @@ checksum = "23b97319f7b8343df12cc98938e5c3eb436064524c8d2b4e30a1d3a36eecdf81"
 dependencies = [
  "memchr",
 ]
+
+[[package]]
+name = "zerocopy"
+version = "0.8.56"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "556764e583adb45a9f8d413c2a147fa7e8d821e48e12b14fd560b607998b75eb"
+dependencies = [
+ "zerocopy-derive",
+]
+
+[[package]]
+name = "zerocopy-derive"
+version = "0.8.56"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f2ab42fc20575779bd240faa45f94a74256f755c0fa9e89f0ede20d91d0cdfc1"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn 2.0.118",
+]
diff --git a/sycl/sycl-rs/Cargo.toml b/sycl/sycl-rs/Cargo.toml
index dcd382b..f584f66 100644
--- a/sycl/sycl-rs/Cargo.toml
+++ b/sycl/sycl-rs/Cargo.toml
@@ -29,4 +29,5 @@ pin-project = "1.1.13"
 thiserror = "2.0.18"
 
 [dev-dependencies]
+half = { version = "2.7.1", features = ["bytemuck"] }
 tokio = { version = "1.52.3", features = ["macros", "rt", "rt-multi-thread"] }
diff --git a/sycl/sycl-rs/examples/kernel_launch.rs b/sycl/sycl-rs/examples/kernel_launch.rs
index 1ee43ab..57cc6c8 100644
--- a/sycl/sycl-rs/examples/kernel_launch.rs
+++ b/sycl/sycl-rs/examples/kernel_launch.rs
@@ -6,6 +6,7 @@
 // SPDX-License-Identifier: MIT OR Apache-2.0
 //
 
+use half::f16;
 use sycl_rs::prelude::*;
 
 static IOTA_SRC: &str = r#"
@@ -15,16 +16,16 @@ namespace syclexp = sycl::ext::oneapi::experimental;
 
 extern "C"
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
-void iota(float start, float *ptr) {
+void iota(sycl::half start, sycl::half *ptr) {
     size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
-    ptr[id] = start + static_cast<float>(id);
+    ptr[id] = start + static_cast<sycl::half>(id);
 }
 "#;
 
 #[tokio::main]
 async fn main() -> sycl_rs::Result<()> {
     let mut queue = Queue::new();
-    let mut device_array = queue.alloc_device::<f32>(1024)?.await?;
+    let mut device_array = queue.alloc_device::<f16>(1024)?.await?;
 
     let kernel = queue
         .get_context()
@@ -36,12 +37,12 @@ async fn main() -> sycl_rs::Result<()> {
         queue.launch(
             NdRange::new([1024], [16]),
             &kernel,
-            (3.14_f32, &mut device_array),
+            (f16::from_f32(3.14), &mut device_array),
         )
     }?
     .await?;
 
-    let mut host_array = queue.alloc_host::<f32>(1024)?.await?;
+    let mut host_array = queue.alloc_host::<f16>(1024)?.await?;
 
     queue.copy(&device_array, &mut host_array)?.await?;