[SYCL][Bindless][Exp] Add Support For Unsampled Image Arrays #12464

isaacault · 2024-01-23T12:50:30Z

Creation / destruction of unsampled image arrays
Fetching / writing of unsampled image arrays
sycl::ext::oneapi::experimental::image_type::array enum value added
sycl::ext::oneapi::experimental::image_descriptor::array_size member added
sycl::ext::oneapi::experimental::image_descriptor::verify() member function added

Correlated UR PR: [Bindless][Exp] Add Support For Image Arrays #1274

- Creation / destruction of unsampled image arrays - Reading / writing of unsampled image arrays - sycl::ext::oneapi::experimental::image_type::array enum value added - sycl::ext::oneapi::experimental::image_descriptor::array_size member added - sycl::ext::oneapi::experimental::image_descriptor::verify() member function added

sycl/test-e2e/bindless_images/bindless_helpers.hpp

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

sycl/test-e2e/bindless_images/bindless_helpers.hpp

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

sycl/include/sycl/ext/oneapi/bindless_images.hpp

- Use bindless_helpers funcs for vulkan interop testing - Rename fill_rand func to conform to snake case - Use new assertion on unsampled coords in read_image_array - Declare and use syclexp namespace in image array test

jchlanda

I've only gone through libclc changes.

libclc/ptx-nvidiacl/libspirv/images/image.cl

* Remove cast reshuffling * Wrap image array read/write intrinsics in macros

isaacault · 2024-02-06T14:15:28Z

@jchlanda My apologies, I accidentally rebased. I believe I've corrected my mistake, and you can see my changes addressing your feedback in 6b2ca14.

aelovikov-intel · 2024-02-06T16:54:56Z

sycl/include/sycl/ext/oneapi/bindless_images.hpp

+
+#ifdef __SYCL_DEVICE_ONLY__
+#if defined(__NVPTX__)
+  __invoke__ImageArrayWrite((uint64_t)imageHandle.raw_handle, coords,


Can we use more C++'ish cast here? I'd expect it to be sycl::bit_cast or static_cast.

aelovikov-intel · 2024-02-06T16:57:07Z

sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp

+            "Standard images cannot have num_levels greater than 1! Use "
+            "image_type::mipmap for mipmap images.");
+      }
+    } else if (this->type == image_type::array) {


Suggested change

} else if (this->type == image_type::array) {

return;

}

if (this->type == image_type::array) {

same below. That way the reader would know immediately that there are no other checks after the "cascade" of "if"s. Or maybe even change to the switch(this->type).

Good point- changed to switch statement with returns.

aelovikov-intel · 2024-02-06T16:58:19Z

sycl/plugins/unified_runtime/CMakeLists.txt

+  set(UNIFIED_RUNTIME_REPO "https://github.com/isaacault/unified-runtime.git")
+  # commit f6ab4ed9846a472d8cb173c12fb9affae21c0fd0
+  # Author: Isaac Ault <[email protected]>
+  # Date:   Tue Jan 23 12:34:08 2024 +0000
+  #
+  #     [Bindless][Exp] Add Support For Image Arrays
+  set(UNIFIED_RUNTIME_TAG f6ab4ed9846a472d8cb173c12fb9affae21c0fd0)


Can't approve before it's fixed.

https://github.com/orgs/intel/teams/unified-runtime-reviewers will not approve until this is updated, but we do have a policy in our repo to get approvals from other reviewers here before contributions there can proceed. We take full responsibility for blocking this kind of change

aelovikov-intel · 2024-02-06T16:58:52Z

sycl/source/detail/bindless_images.cpp

-                                                        : PI_MEM_TYPE_IMAGE1D);
+
+  if (desc.array_size > 1) {
+    // Image Array


Add periods after each comment here and in other places/fiels.

aelovikov-intel · 2024-02-06T17:00:42Z

sycl/source/detail/bindless_images.cpp

+    // Image Array
+    if (desc.depth > 0) {
+      // Image arrays must be 1D or 2D
+      throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),


Is it missing support or they simply can't be per the specification? If the latter, why can't we rely on verify to throw instead? Same in handler.cpp.

It is the latter. I've removed this here and in handler.cpp and added calls to verify(). Note, there isn't a call added here as the calling functions in this file all call verify() themselves.

aelovikov-intel · 2024-02-06T17:04:27Z

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

+// REQUIRES: linux
+// REQUIRES: cuda
+
+// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out


Suggested change

// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out

// RUN: %{build} -o %t.out

aelovikov-intel · 2024-02-06T17:04:41Z

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

+// REQUIRES: cuda
+
+// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
+// RUN: %t.out


Suggested change

// RUN: %t.out

// RUN: %{run} %t.out

aelovikov-intel · 2024-02-06T17:05:12Z

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

+// Uncomment to print additional test information
+// #define VERBOSE_PRINT


Suggested change

// Uncomment to print additional test information

// #define VERBOSE_PRINT

#define VERBOSE_PRINT 0

and update the checks accordingly. Or (even better) switch to inline constexpr bool and if constexpr.

I'd like to leave this as is. It's how we're implementing verbosity in all of the bindless images tests, and we have an internal task tracking this to be cleaned up.

aelovikov-intel · 2024-02-06T17:09:44Z

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

+  // parallel_for 2D
+  template <int NDims, typename DType, int NChannels, typename KernelName,
+            typename = std::enable_if_t<NDims == 2>>
+  static void run_ndim_test(sycl::queue q, sycl::range<2> globalSize,


I think metaprogramming magic can unify these two versions neatly, but I won't insist.

Keeping as is, for consistency with other bindless images tests.

aelovikov-intel · 2024-02-06T17:11:30Z

sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp

+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  // skip half tests if not supported


I'd add a static_assert against double that would otherwise require aspect:fp64.

We actively don't use doubles in our tests, and this keeps consistency with the rest. I'd like to keep this as is.

jchlanda · 2024-02-07T09:52:49Z

@jchlanda My apologies, I accidentally rebased. I believe I've corrected my mistake, and you can see my changes addressing your feedback in 6b2ca14.

Not a problem, I've left one more question, otherwise the libclc part looks good. Thank you.

* Use image descriptor verify where applicable * Periods after comments

aelovikov-intel

SYCL RT changes are good, cursory reading of e2e tests is fine stylistically, I'd expect the logic/completeness of the tests to be checked by images domain experts from Codeplay.

jchlanda

libclc 👍

* Vulkan test passing * `read_image_array` accepts `HintT` template

* Use convertToOpenCLType helper

kbenzie

oneapi-src/unified-runtime#1274 has been merged, please pull in the latest sycl branch changes and update the UR tag as suggested.

sycl/plugins/unified_runtime/CMakeLists.txt

isaacault · 2024-02-23T18:01:33Z

@intel/llvm-gatekeepers Can this get merged? CI failure unrelated & already reported in #12797 and #12798

kbenzie · 2024-02-26T15:56:43Z

@intel/llvm-gatekeepers please merge, the Windows failure is being tracked in #12798

isaacault temporarily deployed to WindowsCILock January 23, 2024 13:01 — with GitHub Actions Inactive

isaacault temporarily deployed to WindowsCILock January 23, 2024 13:33 — with GitHub Actions Inactive

isaacault mentioned this pull request Jan 23, 2024

[Bindless][Exp] Add Support For Image Arrays oneapi-src/unified-runtime#1274

Merged

isaacault self-assigned this Jan 23, 2024

cppchedy reviewed Jan 23, 2024

View reviewed changes

isaacault added 2 commits January 24, 2024 10:38

Address Feedback:

560a552

- Use bindless_helpers funcs for vulkan interop testing - Rename fill_rand func to conform to snake case - Use new assertion on unsampled coords in read_image_array - Declare and use syclexp namespace in image array test

Merge branch 'sycl' into iault/unsampled_image_arrays

eccbdb1

isaacault temporarily deployed to WindowsCILock January 24, 2024 11:07 — with GitHub Actions Inactive

cppchedy approved these changes Jan 24, 2024

View reviewed changes

isaacault temporarily deployed to WindowsCILock January 24, 2024 11:37 — with GitHub Actions Inactive

Merge branch 'sycl' into iault/unsampled_image_arrays

ee82ed8

isaacault temporarily deployed to WindowsCILock January 25, 2024 10:06 — with GitHub Actions Inactive

isaacault temporarily deployed to WindowsCILock January 25, 2024 10:36 — with GitHub Actions Inactive

isaacault marked this pull request as ready for review January 30, 2024 09:28

isaacault requested review from a team as code owners January 30, 2024 09:28

isaacault requested review from aelovikov-intel and jchlanda January 30, 2024 09:28

Seanst98 approved these changes Jan 30, 2024

View reviewed changes

jchlanda reviewed Jan 30, 2024

View reviewed changes

Address Feedback:

6b2ca14

* Remove cast reshuffling * Wrap image array read/write intrinsics in macros

isaacault force-pushed the iault/unsampled_image_arrays branch from ee82ed8 to 855efbd Compare February 6, 2024 13:57

isaacault had a problem deploying to WindowsCILock February 6, 2024 13:58 — with GitHub Actions Failure

Merge branch 'sycl' into iault/unsampled_image_arrays

e782d6b

isaacault force-pushed the iault/unsampled_image_arrays branch from 855efbd to e782d6b Compare February 6, 2024 14:07

isaacault temporarily deployed to WindowsCILock February 6, 2024 14:07 — with GitHub Actions Inactive

aelovikov-intel reviewed Feb 6, 2024

View reviewed changes

Address Feedback:

7a45785

* Use image descriptor verify where applicable * Periods after comments

aelovikov-intel approved these changes Feb 8, 2024

View reviewed changes

isaacault requested a review from jchlanda February 12, 2024 09:32

jchlanda approved these changes Feb 15, 2024

View reviewed changes

Merge branch 'sycl' into iault/unsampled_image_arrays

5cb9101

isaacault temporarily deployed to WindowsCILock February 15, 2024 11:12 — with GitHub Actions Inactive

isaacault temporarily deployed to WindowsCILock February 15, 2024 11:40 — with GitHub Actions Inactive

isaacault added 2 commits February 20, 2024 17:17

Fix merge

5afc472

* Vulkan test passing * `read_image_array` accepts `HintT` template

Merge branch 'sycl' into iault/unsampled_image_arrays

ddcb510

isaacault temporarily deployed to WindowsCILock February 20, 2024 19:01 — with GitHub Actions Inactive

isaacault temporarily deployed to WindowsCILock February 20, 2024 19:36 — with GitHub Actions Inactive

Fix Merge

efd0fc5

* Use convertToOpenCLType helper

isaacault force-pushed the iault/unsampled_image_arrays branch from 0f5b076 to efd0fc5 Compare February 21, 2024 09:56

isaacault temporarily deployed to WindowsCILock February 21, 2024 10:00 — with GitHub Actions Inactive

isaacault had a problem deploying to WindowsCILock February 21, 2024 10:28 — with GitHub Actions Failure

kbenzie reviewed Feb 23, 2024

View reviewed changes

sycl/plugins/unified_runtime/CMakeLists.txt Outdated Show resolved Hide resolved

Merge branch 'sycl' into iault/unsampled_image_arrays

77113e3

isaacault had a problem deploying to WindowsCILock February 23, 2024 15:09 — with GitHub Actions Error

Fix merge formatting

217545a

isaacault temporarily deployed to WindowsCILock February 23, 2024 15:14 — with GitHub Actions Inactive

kbenzie approved these changes Feb 23, 2024

View reviewed changes

isaacault had a problem deploying to WindowsCILock February 23, 2024 15:42 — with GitHub Actions Failure

Merge branch 'sycl' into iault/unsampled_image_arrays

410f7f7

isaacault temporarily deployed to WindowsCILock February 26, 2024 09:47 — with GitHub Actions Inactive

isaacault had a problem deploying to WindowsCILock February 26, 2024 10:17 — with GitHub Actions Failure

aelovikov-intel merged commit 76ec3f0 into intel:sycl Feb 26, 2024
10 of 11 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL][Bindless][Exp] Add Support For Unsampled Image Arrays #12464

[SYCL][Bindless][Exp] Add Support For Unsampled Image Arrays #12464

isaacault commented Jan 23, 2024 •

edited

Loading

jchlanda left a comment

isaacault commented Feb 6, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 7, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 7, 2024

aelovikov-intel Feb 6, 2024

aarongreig Feb 8, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024 •

edited

Loading

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024

aelovikov-intel Feb 6, 2024

isaacault Feb 8, 2024 •

edited

Loading

jchlanda commented Feb 7, 2024

aelovikov-intel left a comment

jchlanda left a comment

kbenzie left a comment

isaacault commented Feb 23, 2024

kbenzie commented Feb 26, 2024

	// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
	// RUN: %{build} -o %t.out

		// Uncomment to print additional test information
		// #define VERBOSE_PRINT

	// Uncomment to print additional test information
	// #define VERBOSE_PRINT
	#define VERBOSE_PRINT 0

[SYCL][Bindless][Exp] Add Support For Unsampled Image Arrays #12464

[SYCL][Bindless][Exp] Add Support For Unsampled Image Arrays #12464

Conversation

isaacault commented Jan 23, 2024 • edited Loading

jchlanda left a comment

Choose a reason for hiding this comment

isaacault commented Feb 6, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

isaacault Feb 8, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

isaacault Feb 8, 2024 • edited Loading

Choose a reason for hiding this comment

jchlanda commented Feb 7, 2024

aelovikov-intel left a comment

Choose a reason for hiding this comment

jchlanda left a comment

Choose a reason for hiding this comment

kbenzie left a comment

Choose a reason for hiding this comment

isaacault commented Feb 23, 2024

kbenzie commented Feb 26, 2024

isaacault commented Jan 23, 2024 •

edited

Loading

isaacault Feb 8, 2024 •

edited

Loading

isaacault Feb 8, 2024 •

edited

Loading