25 #if defined(__cplusplus)
27 #include <hip/hip_vector_types.h>
28 #include <hip/texture_types.h>
29 #include <hip/hcc_detail/ockl_image.h>
31 #include <type_traits>
33 #define TEXTURE_PARAMETERS_INIT \
34 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
35 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
38 struct __hip_is_tex_channel_type
40 static constexpr
bool value =
41 std::is_same<T, char>::value ||
42 std::is_same<T, unsigned char>::value ||
43 std::is_same<T, short>::value ||
44 std::is_same<T, unsigned short>::value ||
45 std::is_same<T, int>::value ||
46 std::is_same<T, unsigned int>::value ||
47 std::is_same<T, float>::value;
53 struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
55 static constexpr
bool value =
56 __hip_is_tex_channel_type<T>::value &&
63 struct __hip_is_tex_normalized_channel_type
65 static constexpr
bool value =
66 std::is_same<T, char>::value ||
67 std::is_same<T, unsigned char>::value ||
68 std::is_same<T, short>::value ||
69 std::is_same<T, unsigned short>::value;
75 struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
77 static constexpr
bool value =
78 __hip_is_tex_normalized_channel_type<T>::value &&
86 hipTextureReadMode readMode,
87 typename Enable =
void>
90 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
95 hipTextureReadMode readMode>
96 using __hip_tex_ret_t =
typename __hip_tex_ret<T, readMode, bool>::type;
101 hipReadModeElementType,
102 typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
110 struct __hip_tex_ret<
111 HIP_vector_type<T, rank>,
112 hipReadModeElementType,
113 typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
115 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
119 struct __hip_tex_ret<
121 hipReadModeNormalizedFloat,
122 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
130 struct __hip_tex_ret<
131 HIP_vector_type<T, rank>,
132 hipReadModeNormalizedFloat,
133 typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
135 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
138 template <
typename T, hipTextureReadMode readMode>
139 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t,
int x)
141 TEXTURE_PARAMETERS_INIT;
142 auto tmp = __ockl_image_load_1Db(i, x);
143 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
146 template <
typename T, hipTextureReadMode readMode>
147 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t,
float x)
149 TEXTURE_PARAMETERS_INIT;
150 auto tmp = __ockl_image_sample_1D(i, s, x);
151 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
154 template <
typename T, hipTextureReadMode readMode>
155 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t,
float x,
float y)
157 TEXTURE_PARAMETERS_INIT;
158 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
159 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
162 template <
typename T, hipTextureReadMode readMode>
163 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer)
165 TEXTURE_PARAMETERS_INIT;
166 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
167 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
170 template <
typename T, hipTextureReadMode readMode>
171 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer)
173 TEXTURE_PARAMETERS_INIT;
174 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
175 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
178 template <
typename T, hipTextureReadMode readMode>
179 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z)
181 TEXTURE_PARAMETERS_INIT;
182 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
183 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
186 template <
typename T, hipTextureReadMode readMode>
187 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z)
189 TEXTURE_PARAMETERS_INIT;
190 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
191 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
194 template <
typename T, hipTextureReadMode readMode>
195 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t,
float x,
float level)
197 TEXTURE_PARAMETERS_INIT;
198 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
199 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
202 template <
typename T, hipTextureReadMode readMode>
203 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float level)
205 TEXTURE_PARAMETERS_INIT;
206 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
207 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
210 template <
typename T, hipTextureReadMode readMode>
211 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float level)
213 TEXTURE_PARAMETERS_INIT;
214 auto tmp = __ockl_image_sample_lod_1Da(i, s,
float2(x, layer).data, level);
215 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
218 template <
typename T, hipTextureReadMode readMode>
219 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer,
float level)
221 TEXTURE_PARAMETERS_INIT;
222 auto tmp = __ockl_image_sample_lod_2Da(i, s,
float4(x, y, layer, 0.0f).data, level);
223 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
226 template <
typename T, hipTextureReadMode readMode>
227 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z,
float level)
229 TEXTURE_PARAMETERS_INIT;
230 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
231 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
234 template <
typename T, hipTextureReadMode readMode>
235 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z,
float level)
237 TEXTURE_PARAMETERS_INIT;
238 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
239 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
242 template <
typename T, hipTextureReadMode readMode>
243 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer)
245 TEXTURE_PARAMETERS_INIT;
246 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
247 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
250 template <
typename T, hipTextureReadMode readMode>
251 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredLod(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer,
float level)
253 TEXTURE_PARAMETERS_INIT;
254 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
255 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
258 template <
typename T, hipTextureReadMode readMode>
259 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
261 TEXTURE_PARAMETERS_INIT;
268 template <
typename T, hipTextureReadMode readMode>
269 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer,
float4 dPdx,
float4 dPdy)
271 TEXTURE_PARAMETERS_INIT;
278 template <
typename T, hipTextureReadMode readMode>
279 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t,
float x,
float dPdx,
float dPdy)
281 TEXTURE_PARAMETERS_INIT;
282 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
283 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
286 template <
typename T, hipTextureReadMode readMode>
287 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float2 dPdx,
float2 dPdy)
289 TEXTURE_PARAMETERS_INIT;
290 auto tmp = __ockl_image_sample_grad_2D(i, s,
float2(x, y).data,
float2(dPdx.x, dPdx.y).data,
float2(dPdy.x, dPdy.y).data);
291 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
294 template <
typename T, hipTextureReadMode readMode>
295 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float dPdx,
float dPdy)
297 TEXTURE_PARAMETERS_INIT;
298 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
299 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
302 template <
typename T, hipTextureReadMode readMode>
303 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer,
float2 dPdx,
float2 dPdy)
305 TEXTURE_PARAMETERS_INIT;
306 auto tmp = __ockl_image_sample_grad_2Da(i, s,
float4(x, y, layer, 0.0f).data,
float2(dPdx.x, dPdx.y).data,
float2(dPdy.x, dPdy.y).data);
307 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
310 template <
typename T, hipTextureReadMode readMode>
311 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z,
float4 dPdx,
float4 dPdy)
313 TEXTURE_PARAMETERS_INIT;
314 auto tmp = __ockl_image_sample_grad_3D(i, s,
float4(x, y, z, 0.0f).data,
float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data,
float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
315 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
320 hipTextureReadMode readMode,
321 typename Enable =
void>
322 struct __hip_tex2dgather_ret
324 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
329 hipTextureReadMode readMode>
330 using __hip_tex2dgather_ret_t =
typename __hip_tex2dgather_ret<T, readMode, bool>::type;
332 template <
typename T>
333 struct __hip_tex2dgather_ret<
335 hipReadModeElementType,
336 typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
338 using type = HIP_vector_type<T, 4>;
344 struct __hip_tex2dgather_ret<
345 HIP_vector_type<T, rank>,
346 hipReadModeElementType,
347 typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
349 using type = HIP_vector_type<T, 4>;
352 template <
typename T>
353 struct __hip_tex2dgather_ret<
355 hipReadModeNormalizedFloat,
356 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
361 template <
typename T, hipTextureReadMode readMode>
362 static __forceinline__ __device__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
int comp=0)
364 TEXTURE_PARAMETERS_INIT;
367 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
368 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);
371 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
372 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);
375 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
376 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);
379 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
380 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);