[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]>
intel · Dec 20, 2024 · 437f972 · 437f972
1 parent 38b5829
commit 437f972
Show file tree

Hide file tree

Showing 50 changed files with 1,215 additions and 55 deletions.
diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES b/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES
@@ -15,6 +15,7 @@ atomic/atomic_max.cl
 atomic/atomic_sub.cl
 atomic/atomic_store.cl
 synchronization/barrier.cl
+images/image.cl
 math/acos.cl
 math/acosh.cl
 math/asin.cl

diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/images/image.cl b/libclc/libspirv/lib/amdgcn-amdhsa/images/image.cl
diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake
@@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
+  set(UNIFIED_RUNTIME_REPO "https://github.com/GeorgeWeb/unified-runtime.git")
   include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")

diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake
@@ -4,4 +4,4 @@
 # Date:   Thu Dec 19 11:26:01 2024 +0000
 #     Merge pull request #2277 from igchor/cooperative_fix
 #     [Spec] fix urKernelSuggestMaxCooperativeGroupCountExp
-set(UNIFIED_RUNTIME_TAG ea0f3a1f5f15f9af7bf40bd13669afeb9ada569c)
+set(UNIFIED_RUNTIME_TAG georgi/bindless-hip)
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc
@@ -2039,7 +2039,7 @@ void release_external_semaphore(external_semaphore semaphoreHandle,
 ```cpp
 #include <sycl/sycl.hpp>
 
-include::../../../test-e2e/bindless_images/examples/example_1_1D_read_write.cpp[lines=9..-1]
+include::../../../test-e2e/bindless_images/examples/example_1_1D_read_write.cpp[lines=12..-1]
 ```
 
 === 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
 ```cpp
 #include <sycl/sycl.hpp>
 
-include::../../../test-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp[lines=9..-1]
+include::../../../test-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp[lines=10..-1]
 ```
 
 === 1D image array read/write
 ```cpp
 #include <sycl/sycl.hpp>
 
-include::../../../test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp[lines=9..-1]
+include::../../../test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp[lines=14..-1]
 ```
 
 === Sampling a cubemap
 
 ```c++
 #include <sycl/sycl.hpp>
 
-include::../../../test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp[lines=9..-1]
+include::../../../test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp[lines=10..-1]
 ```
 
 === Using imported memory and semaphore objects
 
 ```c++
 #include <sycl/sycl.hpp>
 
-include::../../../test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp[lines=8..-1]
+include::../../../test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp[lines=14..-1]
 ```
 
 == Implementation notes

diff --git a/sycl/test-e2e/bindless_images/3_channel_format.cpp b/sycl/test-e2e/bindless_images/3_channel_format.cpp
@@ -1,4 +1,8 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip || level_zero
+// UNSUPPORTED-INTENDED: Unimplemented in the HIP adapter yet.
+// Also, the feature is not fully implemented in the Level Zero stack.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/array/fetch_sampled_array.cpp b/sycl/test-e2e/bindless_images/array/fetch_sampled_array.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_image_array
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/array/read_sampled_array.cpp b/sycl/test-e2e/bindless_images/array/read_sampled_array.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_image_array
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
@@ -137,7 +137,17 @@ static bool runTest(sycl::range<NDims> dims, sycl::range<NDims> localSize,
                     unsigned int seed = 0) {
   using VecType = sycl::vec<DType, NChannels>;
 
-  sycl::device dev;
+  sycl::device dev{};
+  // skip half tests if the device does not support the aspect.
+  if constexpr (std::is_same_v<DType, sycl::half>) {
+    if (!dev.has(sycl::aspect::fp16)) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Test skipped due to lack of device support for fp16\n";
+#endif
+      return false;
+    }
+  }
+
   sycl::queue q(dev);
   auto ctxt = q.get_context();
 

diff --git a/sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp b/sycl/test-e2e/bindless_images/array/read_write_1d_subregion.cpp
@@ -1,7 +1,7 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_image_array
 
-// RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
+// RUN: %if !any-device-is-hip %{ %{build} -o %t.out %}
+// RUN: %if !any-device-is-hip %{ %{run} %t.out %}
 
 #include <iostream>
 #include <sycl/detail/core.hpp>

diff --git a/sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp b/sycl/test-e2e/bindless_images/array/read_write_2d_subregion.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_image_array
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp b/sycl/test-e2e/bindless_images/array/read_write_unsampled_array.cpp
@@ -1,5 +1,4 @@
-// REQUIRES: linux
-// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_image_array
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out

diff --git a/sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp b/sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp
@@ -1,6 +1,4 @@
-// REQUIRES: cuda,aspect-ext_oneapi_cubemap
 // REQUIRES: aspect-ext_oneapi_cubemap_seamless_filtering
-// REQUIRES: build-and-run-mode
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/cubemap/cubemap_unsampled.cpp b/sycl/test-e2e/bindless_images/cubemap/cubemap_unsampled.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda,aspect-ext_oneapi_cubemap
+// REQUIRES: aspect-ext_oneapi_cubemap
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out

diff --git a/sycl/test-e2e/bindless_images/device_to_device_copy.cpp b/sycl/test-e2e/bindless_images/device_to_device_copy.cpp
@@ -1,4 +1,7 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in this test.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/device_to_device_copy_1D_subregion.cpp b/sycl/test-e2e/bindless_images/device_to_device_copy_1D_subregion.cpp
@@ -1,4 +1,7 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in this test.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
@@ -15,9 +18,9 @@ namespace syclexp = sycl::ext::oneapi::experimental;
 
 void copy_image_mem_handle_to_image_mem_handle(
     const syclexp::image_descriptor &dataInDesc,
-    const syclexp::image_descriptor &outDesc,
-    const std::vector<float> &dataIn1, const std::vector<float> &dataIn2,
-    sycl::device dev, sycl::queue q, std::vector<float> &out) {
+    const syclexp::image_descriptor &outDesc, const std::vector<float> &dataIn1,
+    const std::vector<float> &dataIn2, sycl::device dev, sycl::queue q,
+    std::vector<float> &out) {
 
   // Check that output image is double size of input images
   assert(outDesc.width == dataInDesc.width * 2);

diff --git a/sycl/test-e2e/bindless_images/device_to_device_copy_2D_subregion.cpp b/sycl/test-e2e/bindless_images/device_to_device_copy_2D_subregion.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/device_to_device_copy_3D_subregion.cpp b/sycl/test-e2e/bindless_images/device_to_device_copy_3D_subregion.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled_semaphore.cpp b/sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled_semaphore.cpp
@@ -1,6 +1,7 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
 // REQUIRES: windows
-// XFAIL: *
+
+// XFAIL: run-mode
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/15851
 
 // RUN: %{build} -l d3d12 -l dxgi -l dxguid -o %t.out

diff --git a/sycl/test-e2e/bindless_images/examples/example_1_1D_read_write.cpp b/sycl/test-e2e/bindless_images/examples/example_1_1D_read_write.cpp
@@ -1,4 +1,7 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/examples/example_2_2D_dynamic_read.cpp b/sycl/test-e2e/bindless_images/examples/example_2_2D_dynamic_read.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/...est-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp b/...est-e2e/bindless_images/examples/example_3_1D_mipmap_anisotropic_filtering_and_levels.cpp
@@ -1,4 +1,5 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_mipmap
+// REQUIRES: aspect-ext_oneapi_mipmap_anisotropy
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp b/sycl/test-e2e/bindless_images/examples/example_4_1D_array_read_write.cpp
@@ -1,4 +1,7 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

diff --git a/sycl/test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp b/sycl/test-e2e/bindless_images/examples/example_5_sample_cubemap.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_cubemap
 // REQUIRES: build-and-run-mode
 
 // RUN: %{build} -o %t.out

diff --git a/sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp b/sycl/test-e2e/bindless_images/examples/example_6_import_memory_and_semaphores.cpp
@@ -1,6 +1,12 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_external_memory_import
+// REQUIRES: aspect-ext_oneapi_external_semaphore_import
 
 // RUN: %{build} -o %t.out
+// This test is not being executed via the {run} command due to using invalid
+// external input and output file descriptors for the external resource that is
+// being imported. The purpose of this test is to showcase the interop APIs and
+// in order to properly obtain those descriptors we would need a lot of Vulkan
+// context and texture setup as a prerequisite to the example and complicate it.
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/bindless_images.hpp>

diff --git a/sycl/test-e2e/bindless_images/image_get_info.cpp b/sycl/test-e2e/bindless_images/image_get_info.cpp
@@ -1,4 +1,8 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip || level_zero
+// UNSUPPORTED-INTENDED: Image channels queries not working correctly on HIP.
+// Also, the feature is not fully implemented in the Level Zero stack.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out

diff --git a/sycl/test-e2e/bindless_images/mipmap/mipmap_read_1D.cpp b/sycl/test-e2e/bindless_images/mipmap/mipmap_read_1D.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_mipmap
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out

diff --git a/sycl/test-e2e/bindless_images/mipmap/mipmap_read_2D.cpp b/sycl/test-e2e/bindless_images/mipmap/mipmap_read_2D.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_mipmap
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out

diff --git a/sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp b/sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_mipmap
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out

diff --git a/sycl/test-e2e/bindless_images/read_norm_types.cpp b/sycl/test-e2e/bindless_images/read_norm_types.cpp
@@ -1,5 +1,8 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
 
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Returning non fp[32/16] values from sampling fails.
+
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
 

diff --git a/sycl/test-e2e/bindless_images/read_sampled.cpp b/sycl/test-e2e/bindless_images/read_sampled.cpp
@@ -1,4 +1,8 @@
-// REQUIRES: cuda
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip || level_zero
+// UNSUPPORTED-INTENDED: Returning non-FP values from fetching fails on HIP.
+// Also, the feature is not fully implemented in the Level Zero stack.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out
@@ -127,7 +131,17 @@ static bool runTest(sycl::range<NDims> dims, sycl::range<NDims> localSize,
                     unsigned int seed = 0) {
   using VecType = sycl::vec<DType, NChannels>;
 
-  sycl::device dev;
+  sycl::device dev{};
+  // skip half tests if not supported
+  if constexpr (std::is_same_v<DType, sycl::half>) {
+    if (!dev.has(sycl::aspect::fp16)) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Test skipped due to lack of device support for fp16\n";
+#endif
+      return false;
+    }
+  }
+
   sycl::queue q(dev);
   auto ctxt = q.get_context();
 

diff --git a/sycl/test-e2e/bindless_images/read_write_1D.cpp b/sycl/test-e2e/bindless_images/read_write_1D.cpp
@@ -1,5 +1,8 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
 
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
+
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
 

diff --git a/sycl/test-e2e/bindless_images/read_write_1D_subregion.cpp b/sycl/test-e2e/bindless_images/read_write_1D_subregion.cpp
@@ -1,5 +1,8 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
 
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
+
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
 

diff --git a/sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp b/sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp
@@ -1,5 +1,8 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
 
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test.
+
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
 

diff --git a/sycl/test-e2e/bindless_images/read_write_unsampled.cpp b/sycl/test-e2e/bindless_images/read_write_unsampled.cpp
@@ -1,5 +1,8 @@
 // REQUIRES: aspect-ext_oneapi_bindless_images
 
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Returning non fp[32/16] values from sampling fails.
+
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
 

diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM.cpp
@@ -1,4 +1,3 @@
-// REQUIRES: cuda
 // REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_1d_usm
 
 // RUN: %{build} -o %t.out

diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp
@@ -1,4 +1,3 @@
-// REQUIRES: cuda
 // REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d
 
 // RUN: %{build} -o %t.out

diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM.cpp
@@ -1,4 +1,3 @@
-// REQUIRES: cuda
 // REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d_usm
 
 // RUN: %{build} -o %t.out

diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_3D.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_3D.cpp
@@ -1,4 +1,3 @@
-// REQUIRES: cuda
 // REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_3d
 
 // RUN: %{build} -o %t.out

diff --git a/sycl/test-e2e/bindless_images/sampling_2D.cpp b/sycl/test-e2e/bindless_images/sampling_2D.cpp
@@ -1,4 +1,4 @@
-// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_bindless_images_2d_usm
 
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out

diff --git a/sycl/test-e2e/bindless_images/sampling_2D_USM_shared.cpp b/sycl/test-e2e/bindless_images/sampling_2D_USM_shared.cpp
@@ -1,6 +1,12 @@
-// REQUIRES: cuda
 // REQUIRES: aspect-ext_oneapi_bindless_images_shared_usm
 
+// This test is unstable (sometimes passes) on HIP-AMD platforms.
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: While rarely, urBindlessImagesSampledImageCreateExp for
+// USM image memory type (with linear sampler) sometimes returns an unsupported
+// feature result code (1:1 mapping from the native errc from the HIP runtime).
+// We think this is likely an issue in the ROCm drivers(could be arch-specific).
+
 // RUN: %{build} -o %t.out
 // RUN: %{run-unfiltered-devices} %t.out