Skip to content

Commit 24ce45c

Browse files
DBDuncanDBDuncan
authored
[SYCL][Bindless][Exp] 3D images accept 3 component vecs instead of 4 (#12581)
Update read/write image functions to only accept coords with three arguments instead of the current four, for 3D images. Before this patch, when reading or writing to 3D bindless images, 4D coordinates had to be used, where the last coord is always ignored. This patch aligns bindless images to the rest of SYCL and other solutions.
1 parent 491e6e4 commit 24ce45c

File tree

13 files changed

+146
-152
lines changed

13 files changed

+146
-152
lines changed

‎libclc/ptx-nvidiacl/libspirv/images/image.cl

Lines changed: 98 additions & 96 deletions
Large diffs are not rendered by default.

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

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1030,15 +1030,11 @@ of the <<recognized_standard_types>>.
10301030
Sampled images cannot be written to using `write_image`.
10311031

10321032
For reading and writing of unsampled images, coordinates are specified by `int`,
1033-
`sycl::vec<int, 2>`, and `sycl::vec<int, 4>` for 1D, 2D, and 3D images,
1033+
`sycl::vec<int, 2>`, and `sycl::vec<int, 3>` for 1D, 2D, and 3D images,
10341034
respectively.
10351035

10361036
Sampled image reads take `float`, `sycl::vec<float, 2>`, and
1037-
`sycl::vec<float, 4>` coordinate types for 1D, 2D, and 3D images, respectively.
1038-
1039-
Note that in the case of 3D reads or writes, coordinates for 3D images take a
1040-
vector of size 4, not 3, as the fourth element in the coordinate vector is
1041-
ignored.
1037+
`sycl::vec<float, 3>` coordinate types for 1D, 2D, and 3D images, respectively.
10421038

10431039
Note also that all images must be used in either read-only or write-only fashion
10441040
within a single kernel invocation; read/write images are not supported.
@@ -1061,7 +1057,7 @@ standard types.
10611057

10621058
* All POD types (`char`, `short`, `int`, `float`, etc.) excluding `double`
10631059
* `sycl::half`
1064-
* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`, `2`, or `4`
1060+
* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`, `2`, or `3`
10651061

10661062
Any other types are classified as user-defined types.
10671063

@@ -1080,7 +1076,7 @@ struct my_short2 {
10801076
```
10811077

10821078
When providing the above types as `DataT` parameters to an image read function,
1083-
the corresponding `HintT` parameters to use would be `sycl::vec<float, 4>` and
1079+
the corresponding `HintT` parameters to use would be `sycl::vec<float, 4>` and
10841080
`sycl::vec<short, 2>`, respectively.
10851081

10861082
== Mipmapped images

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

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -783,8 +783,8 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
783783
const CoordT &coords [[maybe_unused]]) {
784784
detail::assert_unsampled_coords<CoordT>();
785785
constexpr size_t coordSize = detail::coord_size<CoordT>();
786-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
787-
"Expected input coordinate to be have 1, 2, or 4 components "
786+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
787+
"Expected input coordinate to be have 1, 2, or 3 components "
788788
"for 1D, 2D and 3D images, respectively.");
789789

790790
#ifdef __SYCL_DEVICE_ONLY__
@@ -829,8 +829,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
829829
const CoordT &coords [[maybe_unused]]) {
830830
detail::assert_sampled_coords<CoordT>();
831831
constexpr size_t coordSize = detail::coord_size<CoordT>();
832-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
833-
"Expected input coordinate to be have 1, 2, or 4 components "
832+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
833+
"Expected input coordinate to be have 1, 2, or 3 components "
834834
"for 1D, 2D and 3D images, respectively.");
835835

836836
#ifdef __SYCL_DEVICE_ONLY__
@@ -871,8 +871,8 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
871871
const float level [[maybe_unused]]) {
872872
detail::assert_sampled_coords<CoordT>();
873873
constexpr size_t coordSize = detail::coord_size<CoordT>();
874-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
875-
"Expected input coordinate to be have 1, 2, or 4 components "
874+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
875+
"Expected input coordinate to be have 1, 2, or 3 components "
876876
"for 1D, 2D and 3D images, respectively.");
877877

878878
#ifdef __SYCL_DEVICE_ONLY__
@@ -915,8 +915,8 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
915915
const CoordT &dY [[maybe_unused]]) {
916916
detail::assert_sampled_coords<CoordT>();
917917
constexpr size_t coordSize = detail::coord_size<CoordT>();
918-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
919-
"Expected input coordinates and gradients to have 1, 2, or 4 "
918+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
919+
"Expected input coordinates and gradients to have 1, 2, or 3 "
920920
"components for 1D, 2D, and 3D images, respectively.");
921921

922922
#ifdef __SYCL_DEVICE_ONLY__
@@ -961,8 +961,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
961961
const float level [[maybe_unused]]) {
962962
detail::assert_sampled_coords<CoordT>();
963963
constexpr size_t coordSize = detail::coord_size<CoordT>();
964-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
965-
"Expected input coordinate to be have 1, 2, or 4 components "
964+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
965+
"Expected input coordinate to be have 1, 2, or 3 components "
966966
"for 1D, 2D and 3D images, respectively.");
967967

968968
#ifdef __SYCL_DEVICE_ONLY__
@@ -1008,8 +1008,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
10081008
const CoordT &dY [[maybe_unused]]) {
10091009
detail::assert_sampled_coords<CoordT>();
10101010
constexpr size_t coordSize = detail::coord_size<CoordT>();
1011-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
1012-
"Expected input coordinates and gradients to have 1, 2, or 4 "
1011+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
1012+
"Expected input coordinates and gradients to have 1, 2, or 3 "
10131013
"components for 1D, 2D, and 3D images, respectively.");
10141014

10151015
#ifdef __SYCL_DEVICE_ONLY__
@@ -1045,8 +1045,8 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
10451045
const DataT &color [[maybe_unused]]) {
10461046
detail::assert_unsampled_coords<CoordT>();
10471047
constexpr size_t coordSize = detail::coord_size<CoordT>();
1048-
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
1049-
"Expected input coordinate to be have 1, 2, or 4 components "
1048+
static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
1049+
"Expected input coordinate to be have 1, 2, or 3 components "
10501050
"for 1D, 2D and 3D images, respectively.");
10511051

10521052
#ifdef __SYCL_DEVICE_ONLY__

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

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
#include <sycl/sycl.hpp>
99

1010
// Uncomment to print additional test information
11-
#define VERBOSE_PRINT
11+
// #define VERBOSE_PRINT
1212

1313
template <typename DType, sycl::image_channel_type CType> class kernel;
1414

@@ -102,9 +102,8 @@ template <typename DType, sycl::image_channel_type CType> bool runTest() {
102102
// Extension: read mipmap with anisotropic filtering with zero
103103
// viewing gradients
104104
VecType px1 = sycl::ext::oneapi::experimental::read_mipmap<VecType>(
105-
mipHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0),
106-
sycl::float4(0.0f, 0.0f, 0.0f, 0.0f),
107-
sycl::float4(0.0f, 0.0f, 0.0f, 0.0f));
105+
mipHandle, sycl::float3(fdim0, fdim1, fdim2),
106+
sycl::float3(0.0f, 0.0f, 0.0f), sycl::float3(0.0f, 0.0f, 0.0f));
108107

109108
outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1[0];
110109
});

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -76,10 +76,10 @@ int main() {
7676
// Extension: read image data from handle
7777
sycl::float4 px1 =
7878
sycl::ext::oneapi::experimental::read_image<sycl::float4>(
79-
imgHandle1, sycl::int4(dim0, dim1, dim2, 0));
79+
imgHandle1, sycl::int3(dim0, dim1, dim2));
8080
sycl::float4 px2 =
8181
sycl::ext::oneapi::experimental::read_image<sycl::float4>(
82-
imgHandle2, sycl::int4(dim0, dim1, dim2, 0));
82+
imgHandle2, sycl::int3(dim0, dim1, dim2));
8383

8484
sum = px1[0] + px2[0];
8585
outAcc[sycl::id<3>{dim2, dim1, dim0}] = sum;

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

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -71,9 +71,8 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
7171
syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel);
7272
} else if constexpr (NDims == 3) {
7373
OutputType pixel = syclexp::read_image<OutputType>(
74-
imgIn, sycl::float4(dim0, dim1, dim2, 0));
75-
syclexp::write_image(imgOut, sycl::int4(dim0, dim1, dim2, 0),
76-
pixel);
74+
imgIn, sycl::float3(dim0, dim1, dim2));
75+
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel);
7776
}
7877
});
7978
});

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -76,15 +76,15 @@ int main() {
7676
// Extension: read image data from handle
7777
sycl::float4 px1 =
7878
sycl::ext::oneapi::experimental::read_image<sycl::float4>(
79-
imgIn1, sycl::int4(dim0, dim1, dim2, 0));
79+
imgIn1, sycl::int3(dim0, dim1, dim2));
8080
sycl::float4 px2 =
8181
sycl::ext::oneapi::experimental::read_image<sycl::float4>(
82-
imgIn2, sycl::int4(dim0, dim1, dim2, 0));
82+
imgIn2, sycl::int3(dim0, dim1, dim2));
8383

8484
sum = px1[0] + px2[0];
8585
// Extension: write to image with handle
8686
sycl::ext::oneapi::experimental::write_image<sycl::float4>(
87-
imgOut, sycl::int4(dim0, dim1, dim2, 0), sycl::float4(sum));
87+
imgOut, sycl::int3(dim0, dim1, dim2), sycl::float4(sum));
8888
});
8989
});
9090

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -111,14 +111,14 @@ int main() {
111111
float sum = 0;
112112
// Extension: read image data from handle
113113
float px1 = sycl::ext::oneapi::experimental::read_image<float>(
114-
imgHandle1, sycl::int4(dim0, dim1, dim2, 0));
114+
imgHandle1, sycl::int3(dim0, dim1, dim2));
115115
float px2 = sycl::ext::oneapi::experimental::read_image<float>(
116-
imgHandle2, sycl::int4(dim0, dim1, dim2, 0));
116+
imgHandle2, sycl::int3(dim0, dim1, dim2));
117117

118118
sum = px1 + px2;
119119
// Extension: write to image with handle
120120
sycl::ext::oneapi::experimental::write_image<float>(
121-
imgHandle3, sycl::int4(dim0, dim1, dim2, 0), sum);
121+
imgHandle3, sycl::int3(dim0, dim1, dim2), sum);
122122
});
123123
});
124124

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

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -99,24 +99,24 @@ struct util {
9999
if constexpr (NChannels >= 1) {
100100
VecType px1 =
101101
sycl::ext::oneapi::experimental::read_image<VecType>(
102-
input_0, sycl::int4(dim0, dim1, dim2, 0));
102+
input_0, sycl::int3(dim0, dim1, dim2));
103103
VecType px2 =
104104
sycl::ext::oneapi::experimental::read_image<VecType>(
105-
input_1, sycl::int4(dim0, dim1, dim2, 0));
105+
input_1, sycl::int3(dim0, dim1, dim2));
106106

107107
auto sum =
108108
VecType(util::add_kernel<DType, NChannels>(px1, px2));
109109
sycl::ext::oneapi::experimental::write_image<VecType>(
110-
output, sycl::int4(dim0, dim1, dim2, 0), VecType(sum));
110+
output, sycl::int3(dim0, dim1, dim2), VecType(sum));
111111
} else {
112112
DType px1 = sycl::ext::oneapi::experimental::read_image<DType>(
113-
input_0, sycl::int4(dim0, dim1, dim2, 0));
113+
input_0, sycl::int3(dim0, dim1, dim2));
114114
DType px2 = sycl::ext::oneapi::experimental::read_image<DType>(
115-
input_1, sycl::int4(dim0, dim1, dim2, 0));
115+
input_1, sycl::int3(dim0, dim1, dim2));
116116

117117
auto sum = DType(util::add_kernel<DType, NChannels>(px1, px2));
118118
sycl::ext::oneapi::experimental::write_image<DType>(
119-
output, sycl::int4(dim0, dim1, dim2, 0), DType(sum));
119+
output, sycl::int3(dim0, dim1, dim2), DType(sum));
120120
}
121121
});
122122
});

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -80,7 +80,7 @@ int main() {
8080
// Extension: read image data from handle
8181
sycl::float4 px1 =
8282
sycl::ext::oneapi::experimental::read_image<sycl::float4>(
83-
imgHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0));
83+
imgHandle, sycl::float3(fdim0, fdim1, fdim2));
8484

8585
outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1[0];
8686
});

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -90,7 +90,7 @@ int main() {
9090

9191
// Extension: read image data from handle
9292
float px1 = syclexp::read_image<float>(
93-
imgHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0));
93+
imgHandle, sycl::float3(fdim0, fdim1, fdim2));
9494

9595
outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1;
9696
});

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -151,7 +151,7 @@ bool run_sycl(sycl::range<NDims> globalSize, sycl::range<NDims> localSize,
151151
VecType pixel;
152152
pixel = syclexp::read_image<
153153
std::conditional_t<NChannels == 1, DType, VecType>>(
154-
handles.imgInput, sycl::float4(fdim0, fdim1, fdim2, 0));
154+
handles.imgInput, sycl::float3(fdim0, fdim1, fdim2));
155155

156156
pixel *= static_cast<DType>(10.1f);
157157
outAcc[sycl::id{dim2, dim1, dim0}] = pixel;

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

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -209,25 +209,23 @@ void run_ndim_test(sycl::range<NDims> global_size,
209209

210210
if constexpr (NChannels > 1) {
211211
VecType px1 = syclexp::read_image<VecType>(
212-
handles.input_1, sycl::int4(dim0, dim1, dim2, 0));
212+
handles.input_1, sycl::int3(dim0, dim1, dim2));
213213
VecType px2 = syclexp::read_image<VecType>(
214-
handles.input_2, sycl::int4(dim0, dim1, dim2, 0));
214+
handles.input_2, sycl::int3(dim0, dim1, dim2));
215215

216216
auto sum =
217217
VecType(util::add_kernel<VecType, NChannels>(px1, px2));
218-
syclexp::write_image<VecType>(handles.output,
219-
sycl::int4(dim0, dim1, dim2, 0),
220-
VecType(sum));
218+
syclexp::write_image<VecType>(
219+
handles.output, sycl::int3(dim0, dim1, dim2), VecType(sum));
221220
} else {
222221
DType px1 = syclexp::read_image<DType>(
223-
handles.input_1, sycl::int4(dim0, dim1, dim2, 0));
222+
handles.input_1, sycl::int3(dim0, dim1, dim2));
224223
DType px2 = syclexp::read_image<DType>(
225-
handles.input_2, sycl::int4(dim0, dim1, dim2, 0));
224+
handles.input_2, sycl::int3(dim0, dim1, dim2));
226225

227226
auto sum = DType(util::add_kernel<DType, NChannels>(px1, px2));
228-
syclexp::write_image<DType>(handles.output,
229-
sycl::int4(dim0, dim1, dim2, 0),
230-
DType(sum));
227+
syclexp::write_image<DType>(
228+
handles.output, sycl::int3(dim0, dim1, dim2), DType(sum));
231229
}
232230
}
233231
});

0 commit comments

Comments
 (0)