[SYCL][Bindless][Exp] 3D images accept 3 component vecs instead of 4

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -1030,15 +1030,11 @@ of the <<recognized_standard_types>>.
 Sampled images cannot be written to using `write_image`.
 
 For reading and writing of unsampled images, coordinates are specified by `int`, 
-`sycl::vec<int, 2>`, and `sycl::vec<int, 4>` for 1D, 2D, and 3D images, 
+`sycl::vec<int, 2>`, and `sycl::vec<int, 3>` for 1D, 2D, and 3D images, 
 respectively.
 
 Sampled image reads take `float`, `sycl::vec<float, 2>`, and 
-`sycl::vec<float, 4>` coordinate types for 1D, 2D, and 3D images, respectively.
-
-Note that in the case of 3D reads or writes, coordinates for 3D images take a 
-vector of size 4, not 3, as the fourth element in the coordinate vector is 
-ignored.
+`sycl::vec<float, 3>` coordinate types for 1D, 2D, and 3D images, respectively.
 
 Note also that all images must be used in either read-only or write-only fashion 
 within a single kernel invocation; read/write images are not supported.
@@ -1061,7 +1057,7 @@ standard types.
 
 * All POD types (`char`, `short`, `int`, `float`, etc.) excluding `double`
 * `sycl::half`
-* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`, `2`, or `4`
+* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`, `2`, or `3`
 
 Any other types are classified as user-defined types.
 
@@ -1080,7 +1076,7 @@ struct my_short2 {
 ```
 
 When providing the above types as `DataT` parameters to an image read function, 
-the corresponding `HintT` parameters to use would be `sycl::vec<float, 4>` and 
+the corresponding `HintT` parameters to use would be `sycl::vec<float, 4>` and
 `sycl::vec<short, 2>`, respectively.
 
 == Mipmapped images

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

@@ -783,8 +783,8 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]]) {
   detail::assert_unsampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinate to be have 1, 2, or 4 components "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinate to be have 1, 2, or 3 components "
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -829,8 +829,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]]) {
   detail::assert_sampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinate to be have 1, 2, or 4 components "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinate to be have 1, 2, or 3 components "
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -871,8 +871,8 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                   const float level [[maybe_unused]]) {
   detail::assert_sampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinate to be have 1, 2, or 4 components "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinate to be have 1, 2, or 3 components "
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -915,8 +915,8 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                   const CoordT &dY [[maybe_unused]]) {
   detail::assert_sampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinates and gradients to have 1, 2, or 4 "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinates and gradients to have 1, 2, or 3 "
                 "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -961,8 +961,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                  const float level [[maybe_unused]]) {
   detail::assert_sampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinate to be have 1, 2, or 4 components "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinate to be have 1, 2, or 3 components "
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -1008,8 +1008,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &dY [[maybe_unused]]) {
   detail::assert_sampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinates and gradients to have 1, 2, or 4 "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinates and gradients to have 1, 2, or 3 "
                 "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -1045,8 +1045,8 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                  const DataT &color [[maybe_unused]]) {
   detail::assert_unsampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
-  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinate to be have 1, 2, or 4 components "
+  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
+                "Expected input coordinate to be have 1, 2, or 3 components "
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__

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

@@ -8,7 +8,7 @@
 #include <sycl/sycl.hpp>
 
 // Uncomment to print additional test information
-#define VERBOSE_PRINT
+// #define VERBOSE_PRINT
 
 template <typename DType, sycl::image_channel_type CType> class kernel;
 
@@ -102,9 +102,8 @@ template <typename DType, sycl::image_channel_type CType> bool runTest() {
             // Extension: read mipmap with anisotropic filtering with zero
             // viewing gradients
             VecType px1 = sycl::ext::oneapi::experimental::read_mipmap<VecType>(
-                mipHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0),
-                sycl::float4(0.0f, 0.0f, 0.0f, 0.0f),
-                sycl::float4(0.0f, 0.0f, 0.0f, 0.0f));
+                mipHandle, sycl::float3(fdim0, fdim1, fdim2),
+                sycl::float3(0.0f, 0.0f, 0.0f), sycl::float3(0.0f, 0.0f, 0.0f));
 
             outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1[0];
           });

sycl/test-e2e/bindless_images/read_3D.cpp

@@ -76,10 +76,10 @@ int main() {
             // Extension: read image data from handle
             sycl::float4 px1 =
                 sycl::ext::oneapi::experimental::read_image<sycl::float4>(
-                    imgHandle1, sycl::int4(dim0, dim1, dim2, 0));
+                    imgHandle1, sycl::int3(dim0, dim1, dim2));
             sycl::float4 px2 =
                 sycl::ext::oneapi::experimental::read_image<sycl::float4>(
-                    imgHandle2, sycl::int4(dim0, dim1, dim2, 0));
+                    imgHandle2, sycl::int3(dim0, dim1, dim2));
 
             sum = px1[0] + px2[0];
             outAcc[sycl::id<3>{dim2, dim1, dim0}] = sum;

sycl/test-e2e/bindless_images/read_norm_types.cpp

@@ -71,9 +71,8 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
               syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel);
             } else if constexpr (NDims == 3) {
               OutputType pixel = syclexp::read_image<OutputType>(
-                  imgIn, sycl::float4(dim0, dim1, dim2, 0));
-              syclexp::write_image(imgOut, sycl::int4(dim0, dim1, dim2, 0),
-                                   pixel);
+                  imgIn, sycl::float3(dim0, dim1, dim2));
+              syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel);
             }
           });
     });

sycl/test-e2e/bindless_images/read_write_3D.cpp

@@ -76,15 +76,15 @@ int main() {
             // Extension: read image data from handle
             sycl::float4 px1 =
                 sycl::ext::oneapi::experimental::read_image<sycl::float4>(
-                    imgIn1, sycl::int4(dim0, dim1, dim2, 0));
+                    imgIn1, sycl::int3(dim0, dim1, dim2));
             sycl::float4 px2 =
                 sycl::ext::oneapi::experimental::read_image<sycl::float4>(
-                    imgIn2, sycl::int4(dim0, dim1, dim2, 0));
+                    imgIn2, sycl::int3(dim0, dim1, dim2));
 
             sum = px1[0] + px2[0];
             // Extension: write to image with handle
             sycl::ext::oneapi::experimental::write_image<sycl::float4>(
-                imgOut, sycl::int4(dim0, dim1, dim2, 0), sycl::float4(sum));
+                imgOut, sycl::int3(dim0, dim1, dim2), sycl::float4(sum));
           });
     });
 

sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp

@@ -111,14 +111,14 @@ int main() {
             float sum = 0;
             // Extension: read image data from handle
             float px1 = sycl::ext::oneapi::experimental::read_image<float>(
-                imgHandle1, sycl::int4(dim0, dim1, dim2, 0));
+                imgHandle1, sycl::int3(dim0, dim1, dim2));
             float px2 = sycl::ext::oneapi::experimental::read_image<float>(
-                imgHandle2, sycl::int4(dim0, dim1, dim2, 0));
+                imgHandle2, sycl::int3(dim0, dim1, dim2));
 
             sum = px1 + px2;
             // Extension: write to image with handle
             sycl::ext::oneapi::experimental::write_image<float>(
-                imgHandle3, sycl::int4(dim0, dim1, dim2, 0), sum);
+                imgHandle3, sycl::int3(dim0, dim1, dim2), sum);
           });
     });
 

sycl/test-e2e/bindless_images/read_write_unsampled.cpp

@@ -99,24 +99,24 @@ struct util {
               if constexpr (NChannels >= 1) {
                 VecType px1 =
                     sycl::ext::oneapi::experimental::read_image<VecType>(
-                        input_0, sycl::int4(dim0, dim1, dim2, 0));
+                        input_0, sycl::int3(dim0, dim1, dim2));
                 VecType px2 =
                     sycl::ext::oneapi::experimental::read_image<VecType>(
-                        input_1, sycl::int4(dim0, dim1, dim2, 0));
+                        input_1, sycl::int3(dim0, dim1, dim2));
 
                 auto sum =
                     VecType(util::add_kernel<DType, NChannels>(px1, px2));
                 sycl::ext::oneapi::experimental::write_image<VecType>(
-                    output, sycl::int4(dim0, dim1, dim2, 0), VecType(sum));
+                    output, sycl::int3(dim0, dim1, dim2), VecType(sum));
               } else {
                 DType px1 = sycl::ext::oneapi::experimental::read_image<DType>(
-                    input_0, sycl::int4(dim0, dim1, dim2, 0));
+                    input_0, sycl::int3(dim0, dim1, dim2));
                 DType px2 = sycl::ext::oneapi::experimental::read_image<DType>(
-                    input_1, sycl::int4(dim0, dim1, dim2, 0));
+                    input_1, sycl::int3(dim0, dim1, dim2));
 
                 auto sum = DType(util::add_kernel<DType, NChannels>(px1, px2));
                 sycl::ext::oneapi::experimental::write_image<DType>(
-                    output, sycl::int4(dim0, dim1, dim2, 0), DType(sum));
+                    output, sycl::int3(dim0, dim1, dim2), DType(sum));
               }
             });
       });

sycl/test-e2e/bindless_images/sampling_3D.cpp

@@ -80,7 +80,7 @@ int main() {
             // Extension: read image data from handle
             sycl::float4 px1 =
                 sycl::ext::oneapi::experimental::read_image<sycl::float4>(
-                    imgHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0));
+                    imgHandle, sycl::float3(fdim0, fdim1, fdim2));
 
             outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1[0];
           });

sycl/test-e2e/bindless_images/sampling_unique_addr_modes.cpp

@@ -90,7 +90,7 @@ int main() {
 
             // Extension: read image data from handle
             float px1 = syclexp::read_image<float>(
-                imgHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0));
+                imgHandle, sycl::float3(fdim0, fdim1, fdim2));
 
             outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1;
           });

sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp

@@ -151,7 +151,7 @@ bool run_sycl(sycl::range<NDims> globalSize, sycl::range<NDims> localSize,
               VecType pixel;
               pixel = syclexp::read_image<
                   std::conditional_t<NChannels == 1, DType, VecType>>(
-                  handles.imgInput, sycl::float4(fdim0, fdim1, fdim2, 0));
+                  handles.imgInput, sycl::float3(fdim0, fdim1, fdim2));
 
               pixel *= static_cast<DType>(10.1f);
               outAcc[sycl::id{dim2, dim1, dim0}] = pixel;

sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp

@@ -209,25 +209,23 @@ void run_ndim_test(sycl::range<NDims> global_size,
 
               if constexpr (NChannels > 1) {
                 VecType px1 = syclexp::read_image<VecType>(
-                    handles.input_1, sycl::int4(dim0, dim1, dim2, 0));
+                    handles.input_1, sycl::int3(dim0, dim1, dim2));
                 VecType px2 = syclexp::read_image<VecType>(
-                    handles.input_2, sycl::int4(dim0, dim1, dim2, 0));
+                    handles.input_2, sycl::int3(dim0, dim1, dim2));
 
                 auto sum =
                     VecType(util::add_kernel<VecType, NChannels>(px1, px2));
-                syclexp::write_image<VecType>(handles.output,
-                                              sycl::int4(dim0, dim1, dim2, 0),
-                                              VecType(sum));
+                syclexp::write_image<VecType>(
+                    handles.output, sycl::int3(dim0, dim1, dim2), VecType(sum));
               } else {
                 DType px1 = syclexp::read_image<DType>(
-                    handles.input_1, sycl::int4(dim0, dim1, dim2, 0));
+                    handles.input_1, sycl::int3(dim0, dim1, dim2));
                 DType px2 = syclexp::read_image<DType>(
-                    handles.input_2, sycl::int4(dim0, dim1, dim2, 0));
+                    handles.input_2, sycl::int3(dim0, dim1, dim2));
 
                 auto sum = DType(util::add_kernel<DType, NChannels>(px1, px2));
-                syclexp::write_image<DType>(handles.output,
-                                            sycl::int4(dim0, dim1, dim2, 0),
-                                            DType(sum));
+                syclexp::write_image<DType>(
+                    handles.output, sycl::int3(dim0, dim1, dim2), DType(sum));
               }
             }
           });