25 #if defined(__cplusplus)
27 #include <hip/hip_vector_types.h>
28 #include <hip/hip_texture_types.h>
29 #include <hip/amd_detail/ockl_image.h>
31 #if !defined(__HIPCC_RTC__)
32 #include <type_traits>
33 #endif // !defined(__HIPCC_RTC__)
35 #define TEXTURE_PARAMETERS_INIT \
36 unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
37 unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
40 struct __hip_is_tex_channel_type
42 static constexpr
bool value =
43 std::is_same<T, char>::value ||
44 std::is_same<T, unsigned char>::value ||
45 std::is_same<T, short>::value ||
46 std::is_same<T, unsigned short>::value ||
47 std::is_same<T, int>::value ||
48 std::is_same<T, unsigned int>::value ||
49 std::is_same<T, float>::value;
55 struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
57 static constexpr
bool value =
58 __hip_is_tex_channel_type<T>::value &&
65 struct __hip_is_tex_normalized_channel_type
67 static constexpr
bool value =
68 std::is_same<T, char>::value ||
69 std::is_same<T, unsigned char>::value ||
70 std::is_same<T, short>::value ||
71 std::is_same<T, unsigned short>::value;
77 struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
79 static constexpr
bool value =
80 __hip_is_tex_normalized_channel_type<T>::value &&
88 hipTextureReadMode readMode,
89 typename Enable =
void>
92 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
97 hipTextureReadMode readMode>
98 using __hip_tex_ret_t =
typename __hip_tex_ret<T, readMode, bool>::type;
100 template <
typename T>
101 struct __hip_tex_ret<
103 hipReadModeElementType,
104 typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
112 struct __hip_tex_ret<
113 HIP_vector_type<T, rank>,
114 hipReadModeElementType,
115 typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
117 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
121 struct __hip_tex_ret<
123 hipReadModeNormalizedFloat,
124 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
132 struct __hip_tex_ret<
133 HIP_vector_type<T, rank>,
134 hipReadModeNormalizedFloat,
135 typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
137 using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
140 template <
typename T, hipTextureReadMode readMode>
141 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t,
int x)
143 TEXTURE_PARAMETERS_INIT;
144 auto tmp = __ockl_image_load_1Db(i, x);
145 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
148 template <
typename T, hipTextureReadMode readMode>
149 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t,
float x)
151 TEXTURE_PARAMETERS_INIT;
152 auto tmp = __ockl_image_sample_1D(i, s, x);
153 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
156 template <
typename T, hipTextureReadMode readMode>
157 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t,
float x,
float y)
159 TEXTURE_PARAMETERS_INIT;
160 auto tmp = __ockl_image_sample_2D(i, s,
float2(x, y).data);
161 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
164 template <
typename T, hipTextureReadMode readMode>
165 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer)
167 TEXTURE_PARAMETERS_INIT;
168 auto tmp = __ockl_image_sample_1Da(i, s,
float2(x, layer).data);
169 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
172 template <
typename T, hipTextureReadMode readMode>
173 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer)
175 TEXTURE_PARAMETERS_INIT;
176 auto tmp = __ockl_image_sample_2Da(i, s,
float4(x, y, layer, 0.0f).data);
177 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
180 template <
typename T, hipTextureReadMode readMode>
181 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z)
183 TEXTURE_PARAMETERS_INIT;
184 auto tmp = __ockl_image_sample_3D(i, s,
float4(x, y, z, 0.0f).data);
185 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
188 template <
typename T, hipTextureReadMode readMode>
189 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z)
191 TEXTURE_PARAMETERS_INIT;
192 auto tmp = __ockl_image_sample_CM(i, s,
float4(x, y, z, 0.0f).data);
193 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
196 template <
typename T, hipTextureReadMode readMode>
197 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t,
float x,
float level)
199 TEXTURE_PARAMETERS_INIT;
200 auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
201 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
204 template <
typename T, hipTextureReadMode readMode>
205 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float level)
207 TEXTURE_PARAMETERS_INIT;
208 auto tmp = __ockl_image_sample_lod_2D(i, s,
float2(x, y).data, level);
209 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
212 template <
typename T, hipTextureReadMode readMode>
213 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float level)
215 TEXTURE_PARAMETERS_INIT;
216 auto tmp = __ockl_image_sample_lod_1Da(i, s,
float2(x, layer).data, level);
217 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
220 template <
typename T, hipTextureReadMode readMode>
221 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t,
float x,
float y,
int layer,
float level)
223 TEXTURE_PARAMETERS_INIT;
224 auto tmp = __ockl_image_sample_lod_2Da(i, s,
float4(x, y, layer, 0.0f).data, level);
225 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
228 template <
typename T, hipTextureReadMode readMode>
229 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t,
float x,
float y,
float z,
float level)
231 TEXTURE_PARAMETERS_INIT;
232 auto tmp = __ockl_image_sample_lod_3D(i, s,
float4(x, y, z, 0.0f).data, level);
233 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
236 template <
typename T, hipTextureReadMode readMode>
237 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t,
float x,
float y,
float z,
float level)
239 TEXTURE_PARAMETERS_INIT;
240 auto tmp = __ockl_image_sample_lod_CM(i, s,
float4(x, y, z, 0.0f).data, level);
241 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
244 template <
typename T, hipTextureReadMode readMode>
245 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t,
float x,
float y,
float z,
int layer)
247 TEXTURE_PARAMETERS_INIT;
248 auto tmp = __ockl_image_sample_CMa(i, s,
float4(x, y, z, layer).data);
249 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
252 template <
typename T, hipTextureReadMode readMode>
253 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)
255 TEXTURE_PARAMETERS_INIT;
256 auto tmp = __ockl_image_sample_lod_CMa(i, s,
float4(x, y, z, layer).data, level);
257 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
260 template <
typename T, hipTextureReadMode readMode>
261 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)
263 TEXTURE_PARAMETERS_INIT;
270 template <
typename T, hipTextureReadMode readMode>
271 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)
273 TEXTURE_PARAMETERS_INIT;
280 template <
typename T, hipTextureReadMode readMode>
281 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t,
float x,
float dPdx,
float dPdy)
283 TEXTURE_PARAMETERS_INIT;
284 auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
285 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
288 template <
typename T, hipTextureReadMode readMode>
289 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
float2 dPdx,
float2 dPdy)
291 TEXTURE_PARAMETERS_INIT;
292 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);
293 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
296 template <
typename T, hipTextureReadMode readMode>
297 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t,
float x,
int layer,
float dPdx,
float dPdy)
299 TEXTURE_PARAMETERS_INIT;
300 auto tmp = __ockl_image_sample_grad_1Da(i, s,
float2(x, layer).data, dPdx, dPdy);
301 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
304 template <
typename T, hipTextureReadMode readMode>
305 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)
307 TEXTURE_PARAMETERS_INIT;
308 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);
309 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
312 template <
typename T, hipTextureReadMode readMode>
313 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)
315 TEXTURE_PARAMETERS_INIT;
316 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);
317 return *
reinterpret_cast<__hip_tex_ret_t<T, readMode>*
>(&tmp);
322 hipTextureReadMode readMode,
323 typename Enable =
void>
324 struct __hip_tex2dgather_ret
326 static_assert(std::is_same<Enable, void>::value,
"Invalid channel type!");
331 hipTextureReadMode readMode>
332 using __hip_tex2dgather_ret_t =
typename __hip_tex2dgather_ret<T, readMode, bool>::type;
334 template <
typename T>
335 struct __hip_tex2dgather_ret<
337 hipReadModeElementType,
338 typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
340 using type = HIP_vector_type<T, 4>;
346 struct __hip_tex2dgather_ret<
347 HIP_vector_type<T, rank>,
348 hipReadModeElementType,
349 typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
351 using type = HIP_vector_type<T, 4>;
354 template <
typename T>
355 struct __hip_tex2dgather_ret<
357 hipReadModeNormalizedFloat,
358 typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
363 template <
typename T, hipTextureReadMode readMode>
364 static __forceinline__ __device__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t,
float x,
float y,
int comp=0)
366 TEXTURE_PARAMETERS_INIT;
369 auto tmp = __ockl_image_gather4g_2D(i, s,
float2(x, y).data);
370 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);
373 auto tmp = __ockl_image_gather4b_2D(i, s,
float2(x, y).data);
374 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);
377 auto tmp = __ockl_image_gather4a_2D(i, s,
float2(x, y).data);
378 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);
381 auto tmp = __ockl_image_gather4r_2D(i, s,
float2(x, y).data);
382 return *
reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*
>(&tmp);