[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));
  }
  }
  });