From 88e359238088d9c0bb390c1ac82194dbe021d35d Mon Sep 17 00:00:00 2001
From: "Chen, Sheng S" <sheng.s.chen@intel.com>
Date: Tue, 19 Nov 2024 13:53:47 +0800
Subject: [PATCH] [SYCLCompat] Add wrapper for byte addresss image API
 migration.

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>
---
 sycl/include/syclcompat/bindless_image.hpp | 53 ++++++++++++++++++++++
 sycl/include/syclcompat/syclcompat.hpp     |  1 +
 2 files changed, 54 insertions(+)
 create mode 100644 sycl/include/syclcompat/bindless_image.hpp

diff --git a/sycl/include/syclcompat/bindless_image.hpp b/sycl/include/syclcompat/bindless_image.hpp
new file mode 100644
index 0000000000000..cbdf40f2b65b0
--- /dev/null
+++ b/sycl/include/syclcompat/bindless_image.hpp
@@ -0,0 +1,53 @@
+/***************************************************************************
+ *
+ *  Copyright (C) Codeplay Software Ltd.
+ *
+ *  Part of the LLVM Project, under the Apache License v2.0 with LLVM
+ *  Exceptions. See https://llvm.org/LICENSE.txt for license information.
+ *  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *  Unless required by applicable law or agreed to in writing, software
+ *  distributed under the License is distributed on an "AS IS" BASIS,
+ *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ *  See the License for the specific language governing permissions and
+ *  limitations under the License.
+ *
+ *  SYCL compatibility extension
+ *
+ *  bindless_image.hpp
+ *
+ *  Description:
+ *    bindless image functionality for the SYCL compatibility extension
+ **************************************************************************/
+
+// The original source was under the license below:
+//==---- bindless_image.hpp -------------------------------*- C++
+//-*----------------==//
+//
+// Copyright (C) Intel Corporation
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// See https://llvm.org/LICENSE.txt for license information.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+#include <sycl/ext/oneapi/bindless_images.hpp>
+
+namespace syclcompat {
+namespace experimental {
+
+template <typename DataT, typename HintT = DataT, typename CoordT>
+DataT sample_image(
+    const sycl::ext::oneapi::experimental::sampled_image_handle &imageHandle,
+    CoordT &&coords) {
+  if constexpr (std::is_scalar_v<CoordT>) {
+    return sycl::ext::oneapi::experimental::sample_image<DataT, HintT, CoordT>(
+        imageHandle, coords / sizeof(DataT));
+  } else {
+    coords[0] = coords[0] / sizeof(DataT);
+    return sycl::ext::oneapi::experimental::sample_image<DataT, HintT, CoordT>(
+        imageHandle, coords);
+  }
+}
+} // namespace experimental
+} // namespace syclcompat
diff --git a/sycl/include/syclcompat/syclcompat.hpp b/sycl/include/syclcompat/syclcompat.hpp
index 8c5f693794948..7f6df3537e047 100644
--- a/sycl/include/syclcompat/syclcompat.hpp
+++ b/sycl/include/syclcompat/syclcompat.hpp
@@ -32,3 +32,4 @@
 #include <syclcompat/math.hpp>
 #include <syclcompat/memory.hpp>
 #include <syclcompat/util.hpp>
+#include <syclcompat/bindless_image.hpp>