[SYCL][HIP] Disable legacy images on HIP (#17296)

npmiller · web-flow · commit b0f9bedb4277 · 2025-03-05T11:28:17.000+01:00
Legacy image support for HIP was initially hidden behind an environment
variable as it's not really supported, bindless images should be used
instead.

But after looking a bit closer at the existing support, it's pretty much
entirely broken, almost none of our existing legacy images tests even
compile for AMD.

So this patch fully disables the legacy images for HIP and cleans up all
the related `UNSUPPORTED`, in a lot of cases `target-amd` had to be kept
since the `build-only` mode doesn't know about the device aspects.

It also cleans up existing `UNSUPPORTED` tags for `cuda` in the same
tests.
diff --git a/sycl/test-e2e/Adapters/enqueue-arg-order-image.cpp b/sycl/test-e2e/Adapters/enqueue-arg-order-image.cpp
@@ -1,5 +1,7 @@
 // REQUIRES: aspect-ext_intel_legacy_image
 // UNSUPPORTED: target-amd
+// UNSUPPORTED-INTENDED: Legacy images aren't supported on AMD, but also fail
+// to compile. Bindless images should be used instead.
 
 // spir-v gen for legacy images at O0 not working
 // UNSUPPORTED: O0
diff --git a/sycl/test-e2e/Basic/image/image.cpp b/sycl/test-e2e/Basic/image/image.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Basic/image/image_accessor_range.cpp b/sycl/test-e2e/Basic/image/image_accessor_range.cpp
@@ -2,8 +2,9 @@
 // FIXME: Investigate OS-agnostic failures
 // UNSUPPORTED: true
 
-// UNSUPPORTED: cuda || hip
-// CUDA does not support SYCL 1.2.1 images.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA doesn't fully support SYCL 1.2.1 images. Bindless
+// images should be used instead.
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp b/sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp
@@ -1,6 +1,7 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: cuda || target-amd
-// CUDA cannot support SYCL 1.2.1 images.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA doesn't fully support SYCL 1.2.1 images. Bindless
+// images should be used instead.
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_accessor_readwrite.cpp b/sycl/test-e2e/Basic/image/image_accessor_readwrite.cpp
@@ -1,6 +1,7 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: cuda || target-amd
-// CUDA cannot support SYCL 1.2.1 images.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA doesn't fully support SYCL 1.2.1 images. Bindless
+// images should be used instead.
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_accessor_readwrite_half.cpp b/sycl/test-e2e/Basic/image/image_accessor_readwrite_half.cpp
@@ -1,7 +1,8 @@
 // REQUIRES: aspect-fp16, aspect-ext_intel_legacy_image
 
-// UNSUPPORTED: cuda, target-amd
-// CUDA cannot support SYCL 1.2.1 images.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA doesn't fully support SYCL 1.2.1 images. Bindless
+// images should be used instead.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_array.cpp b/sycl/test-e2e/Basic/image/image_array.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_max_size.cpp b/sycl/test-e2e/Basic/image/image_max_size.cpp
@@ -2,8 +2,9 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// UNSUPPORTED: cuda || target-amd
-// CUDA does not support info::device::image3d_max_width query.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA does not support info::device::image3d_max_width
+// query. Bindless images should be used instead.
 
 // The test checks that 'image' with max allowed sizes is handled correctly.
 
diff --git a/sycl/test-e2e/Basic/image/image_read.cpp b/sycl/test-e2e/Basic/image/image_read.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // Temporarily add explicit '-O2' to avoid GPU hang issue with O0 optimization.
 // RUN: %{build} -O2 -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_read_fp16.cpp b/sycl/test-e2e/Basic/image/image_read_fp16.cpp
@@ -1,7 +1,5 @@
 // REQUIRES: aspect-fp16, aspect-ext_intel_legacy_image
 
-// UNSUPPORTED: target-amd
-
 // Temporarily add explicit '-O2' to avoid GPU hang issue with O0 optimization.
 // RUN: %{build} -O2 -o %t.out
 // RUN: %{run} %t.out
diff --git a/sycl/test-e2e/Basic/image/image_sample.cpp b/sycl/test-e2e/Basic/image/image_sample.cpp
@@ -1,7 +1,6 @@
 // REQUIRES: aspect-ext_intel_legacy_image
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
-// UNSUPPORTED: target-amd
 
 #include <sycl/accessor_image.hpp>
 #include <sycl/detail/core.hpp>
diff --git a/sycl/test-e2e/Basic/image/image_write.cpp b/sycl/test-e2e/Basic/image/image_write.cpp
@@ -2,11 +2,6 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// UNSUPPORTED: target-amd
-// TODO: re-enable on cuda device.
-// See https://github.com/intel/llvm/issues/1542#issuecomment-707877817 for more
-// details.
-
 #include "image_write.h"
 
 int main() {
diff --git a/sycl/test-e2e/Basic/image/image_write_fp16.cpp b/sycl/test-e2e/Basic/image/image_write_fp16.cpp
@@ -1,7 +1,5 @@
 // REQUIRES: aspect-fp16, aspect-ext_intel_legacy_image
 
-// UNSUPPORTED: target-amd
-
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Basic/image/lit.local.cfg b/sycl/test-e2e/Basic/image/lit.local.cfg
@@ -2,4 +2,7 @@
 # the OpenCL environments support of the "Unknown" image format:
 # https://github.com/KhronosGroup/SPIRV-Headers/issues/487
 # After the issue is resolved we will re-enable Image support.
-config.unsupported_features += ['spirv-backend']
+#
+# Legacy images aren't supported on AMD and also don't compile, so mark them
+# unsupported here. Bindless images should be used instead.
+config.unsupported_features += ['spirv-backend', 'target-amd']
diff --git a/sycl/test-e2e/Basic/image/srgba-read.cpp b/sycl/test-e2e/Basic/image/srgba-read.cpp
@@ -3,9 +3,6 @@
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/14387
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s
-//
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-INTENDED: legacy image not supported on AMD
 
 #include <iostream>
 #include <sycl/accessor_image.hpp>
diff --git a/sycl/test-e2e/Regression/image_access.cpp b/sycl/test-e2e/Regression/image_access.cpp
@@ -1,9 +1,6 @@
 // REQUIRES: aspect-ext_intel_legacy_image
 // RUN: %{build} -o %t.out
 // RUN: env SYCL_UR_TRACE=2 %{run} %t.out 2>&1 | FileCheck %s
-//
-// UNSUPPORTED: hip
-// CUDA doesn't fully support OpenCL spec conform images.
 
 //==-------------- image_access.cpp - SYCL image accessors test  -----------==//
 //
diff --git a/sycl/test-e2e/Sampler/basic-rw-float.cpp b/sycl/test-e2e/Sampler/basic-rw-float.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/basic-rw.cpp b/sycl/test-e2e/Sampler/basic-rw.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/lit.local.cfg b/sycl/test-e2e/Sampler/lit.local.cfg
@@ -5,4 +5,7 @@
 # merged with intel/llvm to address an issue of mapping from SPIR-V friendly builtins
 # to Image Read/Write instructions
 # After the 1 issue is resolved and 2 is merged we will re-enable Image support.
-config.unsupported_features += ['spirv-backend']
+#
+# Legacy images aren't supported on AMD and also don't compile, so mark them
+# unsupported here. Bindless images should be used instead.
+config.unsupported_features += ['spirv-backend', 'target-amd']
diff --git a/sycl/test-e2e/Sampler/normalized-clamp-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-clamp-linear-float.cpp
@@ -1,5 +1,9 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
+//
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA doesn't fully support SYCL 1.2.1 images. Bindless
+// images should be used instead.
+//
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/normalized-clamp-nearest.cpp b/sycl/test-e2e/Sampler/normalized-clamp-nearest.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/normalized-clampedge-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-clampedge-linear-float.cpp
@@ -1,12 +1,12 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: cuda
+//
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
-// XFAIL: target-amd
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/14732
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/normalized-clampedge-nearest.cpp b/sycl/test-e2e/Sampler/normalized-clampedge-nearest.cpp
@@ -1,11 +1,6 @@
 // REQUIRES: aspect-ext_intel_legacy_image
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
-//
-// Missing __spirv_ImageWrite, __spirv_SampledImage,
-// __spirv_ImageSampleExplicitLod on AMD
-// XFAIL: target-amd
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/14732
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/normalized-mirror-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-mirror-linear-float.cpp
@@ -1,10 +1,11 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/normalized-mirror-nearest.cpp b/sycl/test-e2e/Sampler/normalized-mirror-nearest.cpp
@@ -1,10 +1,10 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA is not handling repeat or mirror correctly with normalized coordinates.
-// Waiting on a fix.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA is not handling repeat or mirror correctly with
+// normalized coordinates. Bindless images should be used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/normalized-none-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-none-linear-float.cpp
@@ -1,10 +1,11 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/normalized-none-nearest.cpp b/sycl/test-e2e/Sampler/normalized-none-nearest.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/normalized-repeat-linear-float.cpp b/sycl/test-e2e/Sampler/normalized-repeat-linear-float.cpp
@@ -1,10 +1,11 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/normalized-repeat-nearest.cpp b/sycl/test-e2e/Sampler/normalized-repeat-nearest.cpp
@@ -1,10 +1,10 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA is not handling repeat or mirror correctly with normalized coordinates.
-// Waiting on a fix.
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA is not handling repeat or mirror correctly with
+// normalized coordinates. Bindless images should be used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/unnormalized-clamp-linear-float.cpp b/sycl/test-e2e/Sampler/unnormalized-clamp-linear-float.cpp
@@ -1,10 +1,11 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/unnormalized-clamp-nearest.cpp b/sycl/test-e2e/Sampler/unnormalized-clamp-nearest.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/unnormalized-clampedge-linear-float.cpp b/sycl/test-e2e/Sampler/unnormalized-clampedge-linear-float.cpp
@@ -1,10 +1,11 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/unnormalized-clampedge-nearest.cpp b/sycl/test-e2e/Sampler/unnormalized-clampedge-nearest.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Sampler/unnormalized-none-linear-float.cpp b/sycl/test-e2e/Sampler/unnormalized-none-linear-float.cpp
@@ -1,10 +1,11 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd, cuda
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
-// type (such as unorm_int8)
+// UNSUPPORTED: cuda
+// UNSUPPORTED-INTENDED: CUDA works with image_channel_type::fp32, but not with
+// any 8-bit per channel type (such as unorm_int8). Bindless images should be
+// used instead.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/sycl/test-e2e/Sampler/unnormalized-none-nearest.cpp b/sycl/test-e2e/Sampler/unnormalized-none-nearest.cpp
@@ -1,5 +1,4 @@
 // REQUIRES: aspect-ext_intel_legacy_image
-// UNSUPPORTED: target-amd
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
diff --git a/sycl/test-e2e/Tracing/image_printers.cpp b/sycl/test-e2e/Tracing/image_printers.cpp
@@ -31,14 +31,11 @@ int main() {
     sycl::image<2> Img(ImgHostData.data(), ChanOrder, ChanType, ImgSize);
     queue Q;
 
-// legacy Images uses an API that is not supported in hip 4.x
-#if HIP_VERSION_MAJOR >= 5
     Q.submit([&](sycl::handler &CGH) {
       auto ImgAcc = Img.get_access<sycl::float4, SYCLWrite>(CGH);
 
       CGH.single_task<class EmptyTask>([=]() {});
     });
-#endif
   }
   return 0;
 }
diff --git a/sycl/test-e2e/format.py b/sycl/test-e2e/format.py
@@ -283,9 +283,6 @@ def get_extra_env(sycl_devices):
             if "cuda:gpu" in sycl_devices:
                 extra_env.append("UR_CUDA_ENABLE_IMAGE_SUPPORT=1")
 
-            if "hip:gpu" in sycl_devices:
-                extra_env.append("UR_HIP_ENABLE_IMAGE_SUPPORT=1")
-
             return extra_env
 
         extra_env = get_extra_env(devices_for_test)
diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py
@@ -754,8 +754,6 @@ def open_check_file(file_name):
     env["ONEAPI_DEVICE_SELECTOR"] = sycl_device
     if sycl_device.startswith("cuda:"):
         env["UR_CUDA_ENABLE_IMAGE_SUPPORT"] = "1"
-    if sycl_device.startswith("hip:"):
-        env["UR_HIP_ENABLE_IMAGE_SUPPORT"] = "1"
     # When using the ONEAPI_DEVICE_SELECTOR environment variable, sycl-ls
     # prints warnings that might derail a user thinking something is wrong
     # with their test run. It's just us filtering here, so silence them unless
diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp
diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp

Original file line number	Diff line number	Diff line change
`@@ -1,5 +1,4 @@`
`1`	`1`	`// REQUIRES: aspect-ext_intel_legacy_image`
`2`		`-// UNSUPPORTED: target-amd`
`3`	`2`	`// RUN: %{build} -o %t.out`
`4`	`3`	`// RUN: %{run} %t.out`
`5`	`4`