DPC++ Runtime
Runtime libraries for oneAPI DPC++
image_ocl_types.hpp
Go to the documentation of this file.
1 //===-- Image_ocl_types.hpp - Image OpenCL types --------- ------*- C++ -*-===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 // This file is to define some utility functions and declare the structs with
9 // type as appropriate opencl image types based on Dims, AccessMode and
10 // AccessTarget. The macros essentially expand to -
11 //
12 // template <> struct
13 // opencl_image_type<1, access::mode::read, access::target::image> {
14 // using type = __ocl_image1d_ro_t;
15 // };
16 //
17 // template <>
18 // struct opencl_image_type<1, access::mode::write, access::target::image> {
19 // using type = __ocl_image1d_array_wo_t;
20 // };
21 //
22 // As an example, this can be
23 // used as below:
24 // detail::opencl_image_type<2, access::mode::read, access::target::image>::type
25 // MyImage;
26 //
27 #pragma once
28 
29 #ifdef __SYCL_DEVICE_ONLY__
30 
31 #include <sycl/access/access.hpp>
33 
34 #include <CL/__spirv/spirv_ops.hpp>
35 
36 namespace sycl {
38 namespace detail {
39 
40 // Type trait to get the associated sampled image type for a given image type.
41 template <typename ImageType> struct sampled_opencl_image_type;
42 
43 } // namespace detail
44 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
45 } // namespace sycl
46 
47 #define __SYCL_INVOKE_SPIRV_CALL_ARG1(call) \
48  template <typename R, typename T1> inline R __invoke_##call(T1 ParT1) { \
49  using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
50  T1 Arg1 = ParT1; \
51  Ret RetVar = __spirv_##call<Ret, T1>(Arg1); \
52  return sycl::detail::convertDataToType<Ret, R>(RetVar); \
53  }
54 
55 // The macro defines the function __invoke_ImageXXXX,
56 // The functions contains the spirv call to __spirv_ImageXXXX.
57 __SYCL_INVOKE_SPIRV_CALL_ARG1(ImageQuerySize)
58 __SYCL_INVOKE_SPIRV_CALL_ARG1(ImageQueryFormat)
59 __SYCL_INVOKE_SPIRV_CALL_ARG1(ImageQueryOrder)
60 
61 template <typename ImageT, typename CoordT, typename ValT>
62 static void __invoke__ImageWrite(ImageT Img, CoordT Coords, ValT Val) {
63 
64  // Convert from sycl types to builtin types to get correct function mangling.
65  using TmpValT = sycl::detail::ConvertToOpenCLType_t<ValT>;
66  using TmpCoordT = sycl::detail::ConvertToOpenCLType_t<CoordT>;
67 
68  TmpCoordT TmpCoord =
69  sycl::detail::convertDataToType<CoordT, TmpCoordT>(Coords);
70  TmpValT TmpVal = sycl::detail::convertDataToType<ValT, TmpValT>(Val);
71  __spirv_ImageWrite<ImageT, TmpCoordT, TmpValT>(Img, TmpCoord, TmpVal);
72 }
73 
74 template <typename RetType, typename ImageT, typename CoordT>
75 static RetType __invoke__ImageRead(ImageT Img, CoordT Coords) {
76 
77  // Convert from sycl types to builtin types to get correct function mangling.
78  using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
79  using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;
80 
81  TempArgT Arg = sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
82  TempRetT Ret = __spirv_ImageRead<TempRetT, ImageT, TempArgT>(Img, Arg);
83  return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
84 }
85 
86 template <typename RetType, typename ImageT, typename CoordT>
87 static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords,
88  const __ocl_sampler_t &Smpl) {
89 
90  // Convert from sycl types to builtin types to get correct function mangling.
91  using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
92  using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;
93  using SampledT =
94  typename sycl::detail::sampled_opencl_image_type<ImageT>::type;
95 
96  TempArgT TmpCoords =
97  sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
98  // According to validation rules(SPIR-V specification, section 2.16.1) result
99  // of __spirv_SampledImage is allowed to be an operand of image lookup
100  // and query instructions explicitly specified to take an operand whose
101  // type is OpTypeSampledImage.
102  //
103  // According to SPIR-V specification section 3.32.10 at least one operand
104  // setting the level of detail must be present. The last two arguments of
105  // __spirv_ImageSampleExplicitLod represent image operand type and value.
106  // From the SPIR-V specification section 3.14:
107  enum ImageOperands { Lod = 0x2 };
108 
109  // Lod value is zero as mipmap is not supported.
110  TempRetT Ret = __spirv_ImageSampleExplicitLod<SampledT, TempRetT, TempArgT>(
111  __spirv_SampledImage<ImageT, SampledT>(Img, Smpl), TmpCoords,
112  ImageOperands::Lod, 0.0f);
113  return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
114 }
115 
116 namespace sycl {
118 namespace detail {
119 
120 // Function to return the number of channels for Image Channel Order returned by
121 // SPIR-V call to OpImageQueryOrder.
122 // The returned int value represents an enum from Image Channel Order. The enums
123 // for Image Channel Order are mapped differently in sycl and SPIR-V spec.
124 inline int getSPIRVNumChannels(int ImageChannelOrder) {
125  switch (ImageChannelOrder) {
126  case 0: // R
127  case 1: // A
128  case 10: // Rx
129  case 8: // Intensity
130  case 9: // Luminance
131  return 1;
132  case 2: // RG
133  case 3: // RA
134  case 11: // RGx
135  return 2;
136  case 4: // RGB
137  return 3;
138  case 5: // RGBA
139  case 6: // BGRA
140  case 7: // ARGB
141  case 12: // RGBx
142  case 19: // ABGR
143  case 17: // sRGBA
144  return 4;
145  case 13: // Depth
146  case 14: // DepthStencil
147  case 18: // sBGRA
148  // TODO: Enable the below assert after assert is supported for device
149  // compiler. assert(!"Unhandled image channel order in sycl.");
150  default:
151  return 0;
152  }
153 }
154 
155 // Function to compute the Element Size for a given Image Channel Type and Image
156 // Channel Order, returned by SPIR-V calls to OpImageQueryFormat and
157 // OpImageQueryOrder respectively.
158 // The returned int value from OpImageQueryFormat represents an enum from Image
159 // Channel Data Type. The enums for Image Channel Data Type are mapped
160 // differently in sycl and SPIR-V spec.
161 inline int getSPIRVElementSize(int ImageChannelType, int ImageChannelOrder) {
162  int NumChannels = getSPIRVNumChannels(ImageChannelOrder);
163  switch (ImageChannelType) {
164  case 0: // SnormInt8
165  case 2: // UnormInt8
166  case 7: // SignedInt8
167  case 10: // UnsignedInt8
168  return NumChannels;
169  case 1: // SnormInt16
170  case 3: // UnormInt16
171  case 8: // SignedInt16
172  case 11: // UnsignedInt16
173  case 13: // HalfFloat
174  return 2 * NumChannels;
175  case 4: // UnormShort565
176  case 5: // UnormShort555
177  return 2;
178  case 6: // UnormInt101010
179  return 4;
180  case 9: // SignedInt32
181  case 12: // UnsignedInt32
182  case 14: // Float
183  return 4 * NumChannels;
184  case 15: // UnormInt24
185  case 16: // UnormInt101010_2
186  default:
187  // TODO: Enable the below assert after assert is supported for device
188  // compiler. assert(!"Unhandled image channel type in sycl.");
189  return 0;
190  }
191 }
192 
193 template <int Dimensions, access::mode AccessMode, access::target AccessTarget>
194 struct opencl_image_type;
195 
196 // Creation of dummy ocl types for host_image targets.
197 // These dummy ocl types are needed by the compiler parser for the compilation
198 // of host application code with __SYCL_DEVICE_ONLY__ macro set.
199 template <int Dimensions, access::mode AccessMode>
200 struct opencl_image_type<Dimensions, AccessMode, access::target::host_image> {
201  using type =
202  opencl_image_type<Dimensions, AccessMode, access::target::host_image> *;
203 };
204 template <typename T> struct sampled_opencl_image_type<T *> {
205  using type = void *;
206 };
207 
208 #define __SYCL_IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, Ifarray_) \
209  template <> \
210  struct opencl_image_type<Dim, access::mode::AccessMode, \
211  access::target::Target> { \
212  using type = __ocl_image##Dim##d_##Ifarray_##AMSuffix##_t; \
213  };
214 #define __SYCL_SAMPLED_AND_IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, \
215  Ifarray_) \
216  __SYCL_IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, Ifarray_) \
217  template <> \
218  struct sampled_opencl_image_type<typename opencl_image_type< \
219  Dim, access::mode::AccessMode, access::target::Target>::type> { \
220  using type = __ocl_sampled_image##Dim##d_##Ifarray_##AMSuffix##_t; \
221  };
222 
223 #define __SYCL_IMAGETY_READ_3_DIM_IMAGE \
224  __SYCL_SAMPLED_AND_IMAGETY_DEFINE(1, read, ro, image, ) \
225  __SYCL_SAMPLED_AND_IMAGETY_DEFINE(2, read, ro, image, ) \
226  __SYCL_SAMPLED_AND_IMAGETY_DEFINE(3, read, ro, image, )
227 
228 #define __SYCL_IMAGETY_WRITE_3_DIM_IMAGE \
229  __SYCL_IMAGETY_DEFINE(1, write, wo, image, ) \
230  __SYCL_IMAGETY_DEFINE(2, write, wo, image, ) \
231  __SYCL_IMAGETY_DEFINE(3, write, wo, image, )
232 
233 #define __SYCL_IMAGETY_DISCARD_WRITE_3_DIM_IMAGE \
234  __SYCL_IMAGETY_DEFINE(1, discard_write, wo, image, ) \
235  __SYCL_IMAGETY_DEFINE(2, discard_write, wo, image, ) \
236  __SYCL_IMAGETY_DEFINE(3, discard_write, wo, image, )
237 
238 #define __SYCL_IMAGETY_READ_2_DIM_IARRAY \
239  __SYCL_SAMPLED_AND_IMAGETY_DEFINE(1, read, ro, image_array, array_) \
240  __SYCL_SAMPLED_AND_IMAGETY_DEFINE(2, read, ro, image_array, array_)
241 
242 #define __SYCL_IMAGETY_WRITE_2_DIM_IARRAY \
243  __SYCL_IMAGETY_DEFINE(1, write, wo, image_array, array_) \
244  __SYCL_IMAGETY_DEFINE(2, write, wo, image_array, array_)
245 
246 #define __SYCL_IMAGETY_DISCARD_WRITE_2_DIM_IARRAY \
247  __SYCL_IMAGETY_DEFINE(1, discard_write, wo, image_array, array_) \
248  __SYCL_IMAGETY_DEFINE(2, discard_write, wo, image_array, array_)
249 
250 __SYCL_IMAGETY_READ_3_DIM_IMAGE
251 __SYCL_IMAGETY_WRITE_3_DIM_IMAGE
252 __SYCL_IMAGETY_DISCARD_WRITE_3_DIM_IMAGE
253 
254 __SYCL_IMAGETY_READ_2_DIM_IARRAY
255 __SYCL_IMAGETY_WRITE_2_DIM_IARRAY
256 __SYCL_IMAGETY_DISCARD_WRITE_2_DIM_IARRAY
257 
258 } // namespace detail
259 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
260 } // namespace sycl
261 
262 #undef __SYCL_SAMPLED_AND_IMAGETY_DEFINE
263 #undef __SYCL_INVOKE_SPIRV_CALL_ARG1
264 #undef __SYCL_IMAGETY_DEFINE
265 #undef __SYCL_IMAGETY_DISCARD_WRITE_3_DIM_IMAGE
266 #undef __SYCL_IMAGETY_READ_3_DIM_IMAGE
267 #undef __SYCL_IMAGETY_WRITE_3_DIM_IMAGE
268 #undef __SYCL_IMAGETY_DISCARD_WRITE_2_DIM_IARRAY
269 #undef __SYCL_IMAGETY_READ_2_DIM_IARRAY
270 #undef __SYCL_IMAGETY_WRITE_2_DIM_IARRAY
271 #endif // #ifdef __SYCL_DEVICE_ONLY__
#define __SYCL_INLINE_VER_NAMESPACE(X)
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS Dimensions
Definition: accessor.hpp:2782
class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor class __SYCL_EBO __SYCL_SPECIAL_CLASS AccessMode
Definition: accessor.hpp:2782
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
void * __ocl_sampler_t