Skip to content

Commit 9384f74

Browse files
GeorgeWebGeorgeWeb
committed
[SYCL][Ext][Bindless] Initial implementation of image spirv builtins on HIP
TODO: Complete the description with more information. Signed-off-by: Georgi Mirazchiyski <[email protected]>
1 parent 38b5829 commit 9384f74

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

50 files changed

+1215
-55
lines changed

‎libclc/libspirv/lib/amdgcn-amdhsa/SOURCES

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@ atomic/atomic_max.cl
1515
atomic/atomic_sub.cl
1616
atomic/atomic_store.cl
1717
synchronization/barrier.cl
18+
images/image.cl
1819
math/acos.cl
1920
math/acosh.cl
2021
math/asin.cl

‎libclc/libspirv/lib/amdgcn-amdhsa/images/image.cl

Lines changed: 1076 additions & 0 deletions
Large diffs are not rendered by default.

‎sycl/cmake/modules/FetchUnifiedRuntime.cmake

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT)
116116
CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
117117
endfunction()
118118

119-
set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
119+
set(UNIFIED_RUNTIME_REPO "https://git.office.codeplay.com/sycl/unified-runtime.git")
120120
include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake)
121121

122122
set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")

‎sycl/cmake/modules/UnifiedRuntimeTag.cmake

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,4 +4,4 @@
44
# Date: Thu Dec 19 11:26:01 2024 +0000
55
# Merge pull request #2277 from igchor/cooperative_fix
66
# [Spec] fix urKernelSuggestMaxCooperativeGroupCountExp
7-
set(UNIFIED_RUNTIME_TAG ea0f3a1f5f15f9af7bf40bd13669afeb9ada569c)
7+
set(UNIFIED_RUNTIME_TAG peter/bindless-hip)

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2039,7 +2039,7 @@ void release_external_semaphore(external_semaphore semaphoreHandle,
20392039
```cpp
20402040
#include <sycl/sycl.hpp>
20412041

2042-
include::../../../test-e2e/bindless_images/examples/example_1_1D_read_write.cpp[lines=9..-1]
2042+
include::../../../test-e2e/bindless_images/examples/example_1_1D_read_write.cpp[lines=12..-1]
20432043
```
20442044

20452045
=== Reading from a dynamically sized array of 2D images
@@ -2055,30 +2055,30 @@ include::../../../test-e2e/bindless_images/examples/example_2_2D_dynamic_read.cp
20552055
```cpp
20562056
#include <sycl/sycl.hpp>
20572057

2058-
include::../../../test-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp[lines=9..-1]
2058+
include::../../../test-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp[lines=10..-1]
20592059
```
20602060

20612061
=== 1D image array read/write
20622062
```cpp
20632063
#include <sycl/sycl.hpp>
20642064

2065-
include::../../../test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp[lines=9..-1]
2065+
include::../../../test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp[lines=14..-1]
20662066
```
20672067

20682068
=== Sampling a cubemap
20692069

20702070
```c++
20712071
#include <sycl/sycl.hpp>
20722072

2073-
include::../../../test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp[lines=9..-1]
2073+
include::../../../test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp[lines=10..-1]
20742074
```
20752075

20762076
=== Using imported memory and semaphore objects
20772077

20782078
```c++
20792079
#include <sycl/sycl.hpp>
20802080

2081-
include::../../../test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp[lines=8..-1]
2081+
include::../../../test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp[lines=14..-1]
20822082
```
20832083

20842084
== Implementation notes

‎sycl/test-e2e/bindless_images/3_channel_format.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,8 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip || level_zero
4+
// UNSUPPORTED-INTENDED: Unimplemented in the HIP adapter yet.
5+
// Also, the feature is not fully implemented in the Level Zero stack.
26

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

‎sycl/test-e2e/bindless_images/array/fetch_sampled_array.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_image_array
22

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

‎sycl/test-e2e/bindless_images/array/read_sampled_array.cpp

Lines changed: 12 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_image_array
22

33
// RUN: %{build} -o %t.out
44
// RUN: %{run} %t.out
@@ -137,7 +137,17 @@ static bool runTest(sycl::range<NDims> dims, sycl::range<NDims> localSize,
137137
unsigned int seed = 0) {
138138
using VecType = sycl::vec<DType, NChannels>;
139139

140-
sycl::device dev;
140+
sycl::device dev{};
141+
// skip half tests if the device does not support the aspect.
142+
if constexpr (std::is_same_v<DType, sycl::half>) {
143+
if (!dev.has(sycl::aspect::fp16)) {
144+
#ifdef VERBOSE_PRINT
145+
std::cout << "Test skipped due to lack of device support for fp16\n";
146+
#endif
147+
return false;
148+
}
149+
}
150+
141151
sycl::queue q(dev);
142152
auto ctxt = q.get_context();
143153

‎sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_image_array
22

3-
// RUN: %{build} -o %t.out
4-
// RUN: %{run} %t.out
3+
// RUN: %if !any-device-is-hip %{ %{build} -o %t.out %}
4+
// RUN: %if !any-device-is-hip %{ %{run} %t.out %}
55

66
#include <iostream>
77
#include <sycl/detail/core.hpp>

‎sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_image_array
22

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

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
1-
// REQUIRES: linux
2-
// REQUIRES: aspect-ext_oneapi_bindless_images
1+
// REQUIRES: aspect-ext_oneapi_image_array
32

43
// RUN: %{build} -o %t.out
54
// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out

‎sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,4 @@
1-
// REQUIRES: cuda,aspect-ext_oneapi_cubemap
21
// REQUIRES: aspect-ext_oneapi_cubemap_seamless_filtering
3-
// REQUIRES: build-and-run-mode
42

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

‎sycl/test-e2e/bindless_images/cubemap/cubemap_unsampled.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda,aspect-ext_oneapi_cubemap
1+
// REQUIRES: aspect-ext_oneapi_cubemap
22

33
// RUN: %{build} -o %t.out
44
// RUN: %{run-unfiltered-devices} %t.out

‎sycl/test-e2e/bindless_images/device_to_device_copy.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,7 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in this test.
25

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

‎sycl/test-e2e/bindless_images/device_to_device_copy_1D_subregion.cpp

Lines changed: 7 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,7 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in this test.
25

36
// RUN: %{build} -o %t.out
47
// RUN: %{run} %t.out
@@ -15,9 +18,9 @@ namespace syclexp = sycl::ext::oneapi::experimental;
1518

1619
void copy_image_mem_handle_to_image_mem_handle(
1720
const syclexp::image_descriptor &dataInDesc,
18-
const syclexp::image_descriptor &outDesc,
19-
const std::vector<float> &dataIn1, const std::vector<float> &dataIn2,
20-
sycl::device dev, sycl::queue q, std::vector<float> &out) {
21+
const syclexp::image_descriptor &outDesc, const std::vector<float> &dataIn1,
22+
const std::vector<float> &dataIn2, sycl::device dev, sycl::queue q,
23+
std::vector<float> &out) {
2124

2225
// Check that output image is double size of input images
2326
assert(outDesc.width == dataInDesc.width * 2);

‎sycl/test-e2e/bindless_images/device_to_device_copy_2D_subregion.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
22

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

‎sycl/test-e2e/bindless_images/device_to_device_copy_3D_subregion.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
22

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

‎sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled_semaphore.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
22
// REQUIRES: windows
3-
// XFAIL: *
3+
4+
// XFAIL: run-mode
45
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/15851
56

67
// RUN: %{build} -l d3d12 -l dxgi -l dxguid -o %t.out

‎sycl/test-e2e/bindless_images/examples/example_1_1D_read_write.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,7 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
25

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

‎sycl/test-e2e/bindless_images/examples/example_2_2D_dynamic_read.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
22

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

‎sycl/test-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_mipmap
2+
// REQUIRES: aspect-ext_oneapi_mipmap_anisotropy
23

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

‎sycl/test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,7 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
25

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

‎sycl/test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_cubemap
22
// REQUIRES: build-and-run-mode
33

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

‎sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,12 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_external_memory_import
2+
// REQUIRES: aspect-ext_oneapi_external_semaphore_import
23

34
// RUN: %{build} -o %t.out
5+
// This test is not being executed via the {run} command due to using invalid
6+
// external input and output file descriptors for the external resource that is
7+
// being imported. The purpose of this test is to showcase the interop APIs and
8+
// in order to properly obtain those descriptors we would need a lot of Vulkan
9+
// context and texture setup as a prerequisite to the example and complicate it.
410

511
#include <sycl/detail/core.hpp>
612
#include <sycl/ext/oneapi/bindless_images.hpp>

‎sycl/test-e2e/bindless_images/image_get_info.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,8 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip || level_zero
4+
// UNSUPPORTED-INTENDED: Image channels queries not working correctly on HIP.
5+
// Also, the feature is not fully implemented in the Level Zero stack.
26

37
// RUN: %{build} -o %t.out
48
// RUN: %{run-unfiltered-devices} %t.out

‎sycl/test-e2e/bindless_images/mipmap/mipmap_read_1D.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_mipmap
22

33
// RUN: %{build} -o %t.out
44
// RUN: %{run-unfiltered-devices} %t.out

‎sycl/test-e2e/bindless_images/mipmap/mipmap_read_2D.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_mipmap
22

33
// RUN: %{build} -o %t.out
44
// RUN: %{run-unfiltered-devices} %t.out

‎sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_mipmap
22

33
// RUN: %{build} -o %t.out
44
// RUN: %{run-unfiltered-devices} %t.out

‎sycl/test-e2e/bindless_images/read_norm_types.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,8 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
22

3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Returning non fp[32/16] values from sampling fails.
5+
36
// RUN: %{build} -o %t.out
47
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
58

‎sycl/test-e2e/bindless_images/read_sampled.cpp

Lines changed: 16 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,8 @@
1-
// REQUIRES: cuda
1+
// REQUIRES: aspect-ext_oneapi_bindless_images
2+
3+
// UNSUPPORTED: hip || level_zero
4+
// UNSUPPORTED-INTENDED: Returning non-FP values from fetching fails on HIP.
5+
// Also, the feature is not fully implemented in the Level Zero stack.
26

37
// RUN: %{build} -o %t.out
48
// RUN: %{run-unfiltered-devices} %t.out
@@ -127,7 +131,17 @@ static bool runTest(sycl::range<NDims> dims, sycl::range<NDims> localSize,
127131
unsigned int seed = 0) {
128132
using VecType = sycl::vec<DType, NChannels>;
129133

130-
sycl::device dev;
134+
sycl::device dev{};
135+
// skip half tests if not supported
136+
if constexpr (std::is_same_v<DType, sycl::half>) {
137+
if (!dev.has(sycl::aspect::fp16)) {
138+
#ifdef VERBOSE_PRINT
139+
std::cout << "Test skipped due to lack of device support for fp16\n";
140+
#endif
141+
return false;
142+
}
143+
}
144+
131145
sycl::queue q(dev);
132146
auto ctxt = q.get_context();
133147

‎sycl/test-e2e/bindless_images/read_write_1D.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,8 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
22

3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
5+
36
// RUN: %{build} -o %t.out
47
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
58

‎sycl/test-e2e/bindless_images/read_write_1D_subregion.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,8 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
22

3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
5+
36
// RUN: %{build} -o %t.out
47
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
58

‎sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,8 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
22

3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
5+
36
// RUN: %{build} -o %t.out
47
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
58

‎sycl/test-e2e/bindless_images/read_write_unsampled.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,8 @@
11
// REQUIRES: aspect-ext_oneapi_bindless_images
22

3+
// UNSUPPORTED: hip
4+
// UNSUPPORTED-INTENDED: Returning non fp[32/16] values from sampling fails.
5+
36
// RUN: %{build} -o %t.out
47
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
58

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,3 @@
1-
// REQUIRES: cuda
21
// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_1d_usm
32

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

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,3 @@
1-
// REQUIRES: cuda
21
// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d
32

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

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,3 @@
1-
// REQUIRES: cuda
21
// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d_usm
32

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

‎sycl/test-e2e/bindless_images/sampled_fetch/fetch_3D.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,3 @@
1-
// REQUIRES: cuda
21
// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_3d
32

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

‎sycl/test-e2e/bindless_images/sampling_2D.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// REQUIRES: aspect-ext_oneapi_bindless_images
1+
// REQUIRES: aspect-ext_oneapi_bindless_images_2d_usm
22

33
// RUN: %{build} -o %t.out
44
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out

‎sycl/test-e2e/bindless_images/sampling_2D_USM_shared.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,12 @@
1-
// REQUIRES: cuda
21
// REQUIRES: aspect-ext_oneapi_bindless_images_shared_usm
32

3+
// This test is unstable (sometimes passes) on HIP-AMD platforms.
4+
// UNSUPPORTED: hip
5+
// UNSUPPORTED-INTENDED: While rarely, urBindlessImagesSampledImageCreateExp for
6+
// USM image memory type (with linear sampler) sometimes returns an unsupported
7+
// feature result code (1:1 mapping from the native errc from the HIP runtime).
8+
// We think this is likely an issue in the ROCm drivers(could be arch-specific).
9+
410
// RUN: %{build} -o %t.out
511
// RUN: %{run-unfiltered-devices} %t.out
612

0 commit comments

Comments
 (0)