From ef09c4918df687c8179e5e93fb3eb43cf93926bd Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 9 Nov 2017 22:10:55 +0530 Subject: [PATCH 01/19] Texture driver APIs support --- include/hip/hcc_detail/driver_types.h | 46 + include/hip/hcc_detail/hip_runtime_api.h | 24 +- include/hip/hcc_detail/texture_functions.h | 3801 +++++++++++++++++++- include/hip/hcc_detail/texture_types.h | 2 + include/hip/hip_runtime_api.h | 4 +- src/hip_memory.cpp | 138 +- src/hip_module.cpp | 24 +- src/hip_texture.cpp | 92 +- 8 files changed, 4114 insertions(+), 17 deletions(-) diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h index ce5e9789be..5b31e3cd16 100644 --- a/include/hip/hcc_detail/driver_types.h +++ b/include/hip/hcc_detail/driver_types.h @@ -23,6 +23,7 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H #define HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H +typedef void* hipDeviceptr_t; enum hipChannelFormatKind { hipChannelFormatKindSigned = 0, @@ -40,6 +41,29 @@ struct hipChannelFormatDesc enum hipChannelFormatKind f; }; +#define HIP_TRSF_NORMALIZED_COORDINATES 0x02 +#define HIP_TRSF_READ_AS_INTEGER 0x01 +#define HIP_TRSA_OVERRIDE_FORMAT 0x01 + +enum hipArray_Format +{ + HIP_AD_FORMAT_UNSIGNED_INT8 = 0x01, + HIP_AD_FORMAT_UNSIGNED_INT16 = 0x02, + HIP_AD_FORMAT_UNSIGNED_INT32 = 0x03, + HIP_AD_FORMAT_SIGNED_INT8 = 0x08, + HIP_AD_FORMAT_SIGNED_INT16 = 0x09, + HIP_AD_FORMAT_SIGNED_INT32 = 0x0a, + HIP_AD_FORMAT_HALF = 0x10, + HIP_AD_FORMAT_FLOAT = 0x20 +}; + +struct HIP_ARRAY_DESCRIPTOR { + enum hipArray_Format format; + unsigned int numChannels; + size_t width; + size_t height; +}; + struct hipArray { void* data; //FIXME: generalize this struct hipChannelFormatDesc desc; @@ -47,8 +71,30 @@ struct hipArray { unsigned int width; unsigned int height; unsigned int depth; + struct HIP_ARRAY_DESCRIPTOR drvDesc; + bool isDrv; }; +typedef struct hip_Memcpy2D { + size_t height; + size_t widthInBytes; + hipArray* dstArray; + hipDeviceptr_t dstDevice; + void * dstHost; + hipMemoryType dstMemoryType; + size_t dstPitch; + size_t dstXInBytes; + size_t dstY; + hipArray* srcArray; + hipDeviceptr_t srcDevice; + const void * srcHost; + hipMemoryType srcMemoryType; + size_t srcPitch; + size_t srcXInBytes; + size_t srcY; +}hip_Memcpy2D; + + typedef struct hipArray* hipArray_t; typedef const struct hipArray* hipArray_const_t; diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 03be587b0d..9270e8677b 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -84,8 +84,6 @@ typedef struct ihipModule_t *hipModule_t; typedef struct ihipModuleSymbol_t *hipFunction_t; -typedef void* hipDeviceptr_t; - typedef struct ihipEvent_t *hipEvent_t; enum hipLimit_t @@ -621,7 +619,7 @@ hipError_t hipStreamQuery(hipStream_t stream); * * This command is host-synchronous : the host will block until the specified stream is empty. * - * This command follows standard null-stream semantics. Specifically, specifying the null stream will cause the + * This command follows standard null-stream semantics. Specifically, specifying the null stream will cause the * command to wait for other streams on the same device to complete all pending operations. * * This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active or blocking. @@ -644,9 +642,9 @@ hipError_t hipStreamSynchronize(hipStream_t stream); * This function inserts a wait operation into the specified stream. * All future work submitted to @p stream will wait until @p event reports completion before beginning execution. * - * This function only waits for commands in the current stream to complete. Notably,, this function does - * not impliciy wait for commands in the default stream to complete, even if the specified stream is - * created with hipStreamNonBlocking = 0. + * This function only waits for commands in the current stream to complete. Notably,, this function does + * not impliciy wait for commands in the default stream to complete, even if the specified stream is + * created with hipStreamNonBlocking = 0. * * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamDestroy */ @@ -756,7 +754,7 @@ hipError_t hipEventCreate(hipEvent_t* event); * If hipEventRecord() has been previously called on this event, then this call will overwrite any existing state in event. * * If this function is called on a an event that is currently being recorded, results are undefined - either - * outstanding recording may save state into the event, and the order is not guaranteed. + * outstanding recording may save state into the event, and the order is not guaranteed. * * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime * @@ -1318,6 +1316,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags); #endif +hipError_t hipArrayCreate ( hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray ); /** * @brief Frees an array on the device. * @@ -1359,6 +1358,7 @@ hipError_t hipMalloc3DArray(hipArray_t *array, * @see hipMemcpy, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync */ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind); +hipError_t hipMemcpy_2D(const hip_Memcpy2D* pCopy); /** * @brief Copies data between host and device. @@ -1968,6 +1968,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, con */ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name); +hipError_t hipModuleGetTexRef(hipDeviceptr_t *dptr, hipModule_t hmod, const char* name); /** * @brief builds module from code object which resides in host memory. Image is pointer to that location. * @@ -2359,6 +2360,15 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject); hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTextureObject_t textureObject); hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, hipTextureObject_t textureObject); hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject); +hipError_t hipTexRefSetArray ( textureReference* tex, hipArray_const_t array, unsigned int flags ); + +hipError_t hipTexRefSetAddressMode ( textureReference* tex, int dim, hipTextureAddressMode am ); + +hipError_t hipTexRefSetFilterMode ( textureReference* tex, hipTextureFilterMode fm ); + +hipError_t hipTexRefSetFlags ( textureReference* tex, unsigned int flags ); + +hipError_t hipTexRefSetFormat (textureReference* tex, hipArray_Format fmt, int NumPackedComponents ); // doxygen end Texture /** diff --git a/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index d08b429fca..3675c0639d 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/include/hip/hcc_detail/texture_functions.h @@ -39,7 +39,10 @@ union TData { unsigned int ADDRESS_SPACE_2 *i = (unsigned int ADDRESS_SPACE_2*)textureObject; \ unsigned int ADDRESS_SPACE_2 *s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ TData texel; - +#define TEXTURE_REF_PARAMETERS_INIT \ + unsigned int ADDRESS_SPACE_2 *i = (unsigned int ADDRESS_SPACE_2*)texRef.textureObject; \ + unsigned int ADDRESS_SPACE_2 *s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ + TData texel; #define TEXTURE_SET_FLOAT \ *retVal = texel.f.x; @@ -2970,6 +2973,232 @@ __TEXTURE_FUNCTIONS_DECL__ T tex2DLayeredLod(hipTextureObject_t textureObject, f //////////////////////////////////////////////////////////// // Texture Reference APIs //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex1Dfetch(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex1Dfetch(texture texRef, hipTextureObject_t textureObject, int x) { @@ -3194,6 +3423,223 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1Dfetch(texture texR TEXTURE_RETURN_FLOAT_XYZW; } +//////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex1D(texture texRef, int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT_XYZW; +} + //////////////////////////////////////////////////////////// template __TEXTURE_FUNCTIONS_DECL__ char tex1D(texture texRef, hipTextureObject_t textureObject, int x) @@ -3395,6 +3841,14 @@ __TEXTURE_FUNCTIONS_DECL__ float tex1D(texture texRef, hip TEXTURE_RETURN_FLOAT; } +template +__TEXTURE_FUNCTIONS_DECL__ float tex1D(texture texRef, /*hipTextureObject_t textureObject,*/ int x) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1D(i, s, x); + TEXTURE_RETURN_FLOAT; +} + template __TEXTURE_FUNCTIONS_DECL__ float1 tex1D(texture texRef, hipTextureObject_t textureObject, int x) { @@ -3421,6 +3875,231 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1D(texture texRef, h //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_CHAR_X; +} +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex1DLod(texture texRef, float x, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1D(i, s, x, level); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex1DLod(texture texRef, hipTextureObject_t textureObject, float x, float level) { @@ -3646,6 +4325,232 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1DLod(texture texRef //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex1DGrad(texture texRef, float x, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1D(i, s, x, dx, dy); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex1DGrad(texture texRef, hipTextureObject_t textureObject, float x, float dx, float dy) { @@ -3870,6 +4775,202 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1DGrad(texture texRe TEXTURE_RETURN_FLOAT_XYZW; } +//////////////////////////////////////////////////////////// + +template +__TEXTURE_FUNCTIONS_DECL__ char tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + + + //////////////////////////////////////////////////////////// template @@ -4072,6 +5173,22 @@ __TEXTURE_FUNCTIONS_DECL__ float tex2D(texture texRef, hip TEXTURE_RETURN_FLOAT; } +template +__TEXTURE_FUNCTIONS_DECL__ float tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ float1 tex2D(texture texRef, hipTextureObject_t textureObject, float x, float y) { @@ -4080,6 +5197,14 @@ __TEXTURE_FUNCTIONS_DECL__ float1 tex2D(texture texRef, h TEXTURE_RETURN_FLOAT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ float2 tex2D(texture texRef, hipTextureObject_t textureObject, float x, float y) { @@ -4088,6 +5213,14 @@ __TEXTURE_FUNCTIONS_DECL__ float2 tex2D(texture texRef, h TEXTURE_RETURN_FLOAT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex2D(texture texRef, float x, float y) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2D(i, s, hc::short_vector::float2(x, y).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ float4 tex2D(texture texRef, hipTextureObject_t textureObject, float x, float y) { @@ -4098,6 +5231,232 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex2D(texture texRef, h //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex2DLod(texture texRef, float x, float y, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2D(i, s, hc::short_vector::float2(x, y).get_vector(), level); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex2DLod(texture texRef, hipTextureObject_t textureObject, float x, float y, float level) { @@ -4324,6 +5683,316 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex2DLod(texture texRef //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex2DGrad(texture texRef, float x, float y, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2D(i, s, + hc::short_vector::float2(x, y).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex2DGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, float2 dx, float2 dy) { @@ -4634,6 +6303,232 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex2DGrad(texture texRe //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex3D(texture texRef, float x, float y, float z) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex3D(texture texRef, hipTextureObject_t textureObject, float x, float y, float z) { @@ -4860,6 +6755,168 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex3D(texture texRef, h //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex3DLod(texture texRef, float x, float y, float z, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_3D(i, s, hc::short_vector::float4(x, y, z, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex3DLod(texture texRef, hipTextureObject_t textureObject, float x, float y, float z, float level) { @@ -5020,6 +7077,316 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex3DLod(texture texRef TEXTURE_RETURN_FLOAT_XYZW; } +//////////////////////////////////////////////////////////// + +template +__TEXTURE_FUNCTIONS_DECL__ char tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex3DGrad(texture texRef, float x, float y, float z, float4 dx, float4 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f =__ockl_image_sample_grad_3D(i, s, + hc::short_vector::float4(x, y, z, 0.0f).get_vector(), + hc::short_vector::float4(dx.x, dx.y, dx.z, dx.w).get_vector(), + hc::short_vector::float4(dy.x, dy.y, dy.z, dy.w).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + //////////////////////////////////////////////////////////// template __TEXTURE_FUNCTIONS_DECL__ char tex3DGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, float z, float4 dx, float4 dy) @@ -5331,6 +7698,232 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex3DGrad(texture texRe //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayered(texture texRef, float x, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_1Da(i, s, hc::short_vector::float2(x, layer).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex1DLayered(texture texRef, hipTextureObject_t textureObject, float x, int layer) { @@ -5557,6 +8150,232 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayered(texture te //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_CHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_CHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_CHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_CHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UCHAR; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UCHAR_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UCHAR_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UCHAR_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_SHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_SHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_SHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_SHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_USHORT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_USHORT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_USHORT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_USHORT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_INT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_INT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_INT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_INT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UINT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UINT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UINT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_UINT_XYZW; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_FLOAT; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_FLOAT_X; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_FLOAT_XY; +} + +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayeredLod(texture texRef, float x, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), level); + TEXTURE_RETURN_FLOAT_XYZW; +} + +//////////////////////////////////////////////////////////// + template __TEXTURE_FUNCTIONS_DECL__ char tex1DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, int layer, float level) { @@ -5783,6 +8602,14 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayeredLod(texture //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_CHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ char tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5791,6 +8618,14 @@ __TEXTURE_FUNCTIONS_DECL__ char tex1DLayeredGrad(texture te TEXTURE_RETURN_CHAR; } +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_CHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ char1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5799,6 +8634,14 @@ __TEXTURE_FUNCTIONS_DECL__ char1 tex1DLayeredGrad(texture TEXTURE_RETURN_CHAR_X; } +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_CHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ char2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5807,6 +8650,14 @@ __TEXTURE_FUNCTIONS_DECL__ char2 tex1DLayeredGrad(texture TEXTURE_RETURN_CHAR_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_CHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ char4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5815,6 +8666,14 @@ __TEXTURE_FUNCTIONS_DECL__ char4 tex1DLayeredGrad(texture TEXTURE_RETURN_CHAR_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UCHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5823,6 +8682,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned char tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UCHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5831,6 +8698,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar1 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UCHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5839,6 +8714,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar2 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UCHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5847,6 +8730,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar4 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ short tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_SHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ short tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5855,6 +8746,14 @@ __TEXTURE_FUNCTIONS_DECL__ short tex1DLayeredGrad(texture TEXTURE_RETURN_SHORT; } +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_SHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ short1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5863,6 +8762,14 @@ __TEXTURE_FUNCTIONS_DECL__ short1 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ short2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_SHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ short2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5871,6 +8778,14 @@ __TEXTURE_FUNCTIONS_DECL__ short2 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ short4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_SHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ short4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5879,6 +8794,14 @@ __TEXTURE_FUNCTIONS_DECL__ short4 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_USHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5887,6 +8810,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned short tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_USHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5895,6 +8826,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort1 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_USHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5903,6 +8842,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort2 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_USHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5911,6 +8858,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort4 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ int tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_INT; +} + template __TEXTURE_FUNCTIONS_DECL__ int tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5919,6 +8874,14 @@ __TEXTURE_FUNCTIONS_DECL__ int tex1DLayeredGrad(texture texR TEXTURE_RETURN_INT; } +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_INT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ int1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5927,6 +8890,14 @@ __TEXTURE_FUNCTIONS_DECL__ int1 tex1DLayeredGrad(texture te TEXTURE_RETURN_INT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_INT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ int2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5935,6 +8906,14 @@ __TEXTURE_FUNCTIONS_DECL__ int2 tex1DLayeredGrad(texture te TEXTURE_RETURN_INT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_INT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ int4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5943,6 +8922,14 @@ __TEXTURE_FUNCTIONS_DECL__ int4 tex1DLayeredGrad(texture te TEXTURE_RETURN_INT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UINT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5951,6 +8938,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned int tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uint1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UINT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uint1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5959,6 +8954,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint1 tex1DLayeredGrad(texture TEXTURE_RETURN_UINT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UINT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uint2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5967,6 +8970,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint2 tex1DLayeredGrad(texture TEXTURE_RETURN_UINT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_UINT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uint4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5975,6 +8986,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint4 tex1DLayeredGrad(texture TEXTURE_RETURN_UINT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ float tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_FLOAT; +} + template __TEXTURE_FUNCTIONS_DECL__ float tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5983,6 +9002,14 @@ __TEXTURE_FUNCTIONS_DECL__ float tex1DLayeredGrad(texture TEXTURE_RETURN_FLOAT; } +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_FLOAT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ float1 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5991,6 +9018,14 @@ __TEXTURE_FUNCTIONS_DECL__ float1 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ float2 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_FLOAT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ float2 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -5999,6 +9034,14 @@ __TEXTURE_FUNCTIONS_DECL__ float2 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayeredGrad(texture texRef, float x, int layer, float dx, float dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_1Da(i, s, hc::short_vector::float2(x, layer).get_vector(), dx, dy); + TEXTURE_RETURN_FLOAT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, int layer, float dx, float dy) { @@ -6009,6 +9052,14 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex1DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ char tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ char tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6017,6 +9068,14 @@ __TEXTURE_FUNCTIONS_DECL__ char tex2DLayered(texture texRef TEXTURE_RETURN_CHAR; } +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6025,6 +9084,14 @@ __TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayered(texture texR TEXTURE_RETURN_CHAR_X; } +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6033,6 +9100,14 @@ __TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayered(texture texR TEXTURE_RETURN_CHAR_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6041,6 +9116,14 @@ __TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayered(texture texR TEXTURE_RETURN_CHAR_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6049,6 +9132,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayered(texture +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6057,6 +9148,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayered(texture te TEXTURE_RETURN_UCHAR_X; } +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6065,6 +9164,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayered(texture te TEXTURE_RETURN_UCHAR_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6073,6 +9180,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayered(texture te TEXTURE_RETURN_UCHAR_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ short tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ short tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6081,6 +9196,14 @@ __TEXTURE_FUNCTIONS_DECL__ short tex2DLayered(texture texR TEXTURE_RETURN_SHORT; } +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6089,6 +9212,14 @@ __TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayered(texture te TEXTURE_RETURN_SHORT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6097,6 +9228,14 @@ __TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayered(texture te TEXTURE_RETURN_SHORT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6105,6 +9244,14 @@ __TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayered(texture te TEXTURE_RETURN_SHORT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6113,6 +9260,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayered(texture +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6121,6 +9276,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayered(texture TEXTURE_RETURN_USHORT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6129,6 +9292,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayered(texture TEXTURE_RETURN_USHORT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6137,6 +9308,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayered(texture TEXTURE_RETURN_USHORT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ int tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_INT; +} + template __TEXTURE_FUNCTIONS_DECL__ int tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6145,6 +9324,14 @@ __TEXTURE_FUNCTIONS_DECL__ int tex2DLayered(texture texRef, TEXTURE_RETURN_INT; } +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_INT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6153,6 +9340,14 @@ __TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayered(texture texRef TEXTURE_RETURN_INT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_INT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6161,6 +9356,14 @@ __TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayered(texture texRef TEXTURE_RETURN_INT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6169,6 +9372,14 @@ __TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayered(texture texRef TEXTURE_RETURN_INT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UINT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6177,6 +9388,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayered(texture +__TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UINT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6185,6 +9404,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayered(texture texR TEXTURE_RETURN_UINT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6193,6 +9420,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayered(texture texR TEXTURE_RETURN_UINT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6201,6 +9436,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayered(texture texR TEXTURE_RETURN_UINT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ float tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT; +} + template __TEXTURE_FUNCTIONS_DECL__ float tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6209,6 +9452,14 @@ __TEXTURE_FUNCTIONS_DECL__ float tex2DLayered(texture texR TEXTURE_RETURN_FLOAT; } +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6217,6 +9468,14 @@ __TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayered(texture te TEXTURE_RETURN_FLOAT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6225,6 +9484,14 @@ __TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayered(texture te TEXTURE_RETURN_FLOAT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayered(texture texRef, float x, float y, int layer) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayered(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer) { @@ -6235,6 +9502,14 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayered(texture te //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ char tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6243,6 +9518,14 @@ __TEXTURE_FUNCTIONS_DECL__ char tex2DLayeredLod(texture tex TEXTURE_RETURN_CHAR; } +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6251,6 +9534,14 @@ __TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayeredLod(texture t TEXTURE_RETURN_CHAR_X; } +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6259,6 +9550,14 @@ __TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayeredLod(texture t TEXTURE_RETURN_CHAR_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_CHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6267,6 +9566,14 @@ __TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayeredLod(texture t TEXTURE_RETURN_CHAR_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6275,6 +9582,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayeredLod(texture +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6283,6 +9598,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayeredLod(texture TEXTURE_RETURN_UCHAR_X; } +template +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6291,6 +9614,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayeredLod(texture TEXTURE_RETURN_UCHAR_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UCHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6299,6 +9630,14 @@ __TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayeredLod(texture TEXTURE_RETURN_UCHAR_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ short tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_SHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ short tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6307,6 +9646,14 @@ __TEXTURE_FUNCTIONS_DECL__ short tex2DLayeredLod(texture t TEXTURE_RETURN_SHORT; } +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_SHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6315,6 +9662,14 @@ __TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayeredLod(texture TEXTURE_RETURN_SHORT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_SHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6323,6 +9678,14 @@ __TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayeredLod(texture TEXTURE_RETURN_SHORT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_SHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6331,6 +9694,14 @@ __TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayeredLod(texture TEXTURE_RETURN_SHORT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_USHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6339,6 +9710,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayeredLod(texture +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_USHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6347,6 +9726,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayeredLod(texture +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_USHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6355,6 +9742,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayeredLod(texture +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_USHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6363,6 +9758,14 @@ __TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayeredLod(texture +__TEXTURE_FUNCTIONS_DECL__ int tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT; +} + template __TEXTURE_FUNCTIONS_DECL__ int tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6371,6 +9774,14 @@ __TEXTURE_FUNCTIONS_DECL__ int tex2DLayeredLod(texture texRe TEXTURE_RETURN_INT; } +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6379,6 +9790,14 @@ __TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayeredLod(texture tex TEXTURE_RETURN_INT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6387,6 +9806,14 @@ __TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayeredLod(texture tex TEXTURE_RETURN_INT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_INT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6395,6 +9822,14 @@ __TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayeredLod(texture tex TEXTURE_RETURN_INT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6403,6 +9838,14 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayeredLod(texture +__TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6411,6 +9854,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayeredLod(texture t TEXTURE_RETURN_UINT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6419,6 +9870,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayeredLod(texture t TEXTURE_RETURN_UINT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_UINT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6427,6 +9886,14 @@ __TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayeredLod(texture t TEXTURE_RETURN_UINT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ float tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT; +} + template __TEXTURE_FUNCTIONS_DECL__ float tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6435,6 +9902,14 @@ __TEXTURE_FUNCTIONS_DECL__ float tex2DLayeredLod(texture t TEXTURE_RETURN_FLOAT; } +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6443,6 +9918,14 @@ __TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayeredLod(texture TEXTURE_RETURN_FLOAT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6451,6 +9934,14 @@ __TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayeredLod(texture TEXTURE_RETURN_FLOAT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayeredLod(texture texRef, float x, float y, int layer, float level) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_lod_2Da(i, s, hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), level); + TEXTURE_RETURN_FLOAT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayeredLod(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float level) { @@ -6461,6 +9952,17 @@ __TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayeredLod(texture //////////////////////////////////////////////////////////// +template +__TEXTURE_FUNCTIONS_DECL__ char tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ char tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6472,6 +9974,17 @@ __TEXTURE_FUNCTIONS_DECL__ char tex2DLayeredGrad(texture te TEXTURE_RETURN_CHAR; } +template +__TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6483,6 +9996,17 @@ __TEXTURE_FUNCTIONS_DECL__ char1 tex2DLayeredGrad(texture TEXTURE_RETURN_CHAR_X; } +template +__TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6494,6 +10018,17 @@ __TEXTURE_FUNCTIONS_DECL__ char2 tex2DLayeredGrad(texture TEXTURE_RETURN_CHAR_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_CHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6505,6 +10040,17 @@ __TEXTURE_FUNCTIONS_DECL__ char4 tex2DLayeredGrad(texture TEXTURE_RETURN_CHAR_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6516,6 +10062,17 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned char tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6527,6 +10084,17 @@ __TEXTURE_FUNCTIONS_DECL__ uchar1 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6538,6 +10106,17 @@ __TEXTURE_FUNCTIONS_DECL__ uchar2 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UCHAR_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6549,6 +10128,17 @@ __TEXTURE_FUNCTIONS_DECL__ uchar4 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ short tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ short tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6560,6 +10150,17 @@ __TEXTURE_FUNCTIONS_DECL__ short tex2DLayeredGrad(texture TEXTURE_RETURN_SHORT; } +template +__TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6571,6 +10172,17 @@ __TEXTURE_FUNCTIONS_DECL__ short1 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6582,6 +10194,17 @@ __TEXTURE_FUNCTIONS_DECL__ short2 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_SHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6593,6 +10216,17 @@ __TEXTURE_FUNCTIONS_DECL__ short4 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6604,6 +10238,17 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned short tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6615,6 +10260,17 @@ __TEXTURE_FUNCTIONS_DECL__ ushort1 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6626,6 +10282,17 @@ __TEXTURE_FUNCTIONS_DECL__ ushort2 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_USHORT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6637,6 +10304,17 @@ __TEXTURE_FUNCTIONS_DECL__ ushort4 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ int tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT; +} + template __TEXTURE_FUNCTIONS_DECL__ int tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6648,6 +10326,17 @@ __TEXTURE_FUNCTIONS_DECL__ int tex2DLayeredGrad(texture texR TEXTURE_RETURN_INT; } +template +__TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6659,6 +10348,17 @@ __TEXTURE_FUNCTIONS_DECL__ int1 tex2DLayeredGrad(texture te TEXTURE_RETURN_INT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6670,6 +10370,17 @@ __TEXTURE_FUNCTIONS_DECL__ int2 tex2DLayeredGrad(texture te TEXTURE_RETURN_INT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_INT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6681,6 +10392,17 @@ __TEXTURE_FUNCTIONS_DECL__ int4 tex2DLayeredGrad(texture te TEXTURE_RETURN_INT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT; +} + template __TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6692,6 +10414,17 @@ __TEXTURE_FUNCTIONS_DECL__ unsigned int tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6703,6 +10436,17 @@ __TEXTURE_FUNCTIONS_DECL__ uint1 tex2DLayeredGrad(texture TEXTURE_RETURN_UINT_X; } +template +__TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6714,6 +10458,17 @@ __TEXTURE_FUNCTIONS_DECL__ uint2 tex2DLayeredGrad(texture TEXTURE_RETURN_UINT_XY; } +template +__TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_UINT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6725,6 +10480,17 @@ __TEXTURE_FUNCTIONS_DECL__ uint4 tex2DLayeredGrad(texture TEXTURE_RETURN_UINT_XYZW; } +template +__TEXTURE_FUNCTIONS_DECL__ float tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT; +} + template __TEXTURE_FUNCTIONS_DECL__ float tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6736,6 +10502,17 @@ __TEXTURE_FUNCTIONS_DECL__ float tex2DLayeredGrad(texture TEXTURE_RETURN_FLOAT; } +template +__TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT_X; +} + template __TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6747,6 +10524,17 @@ __TEXTURE_FUNCTIONS_DECL__ float1 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT_XY; +} + template __TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { @@ -6758,6 +10546,17 @@ __TEXTURE_FUNCTIONS_DECL__ float2 tex2DLayeredGrad(texture +__TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayeredGrad(texture texRef, float x, float y, int layer, float2 dx, float2 dy) +{ + TEXTURE_REF_PARAMETERS_INIT; + texel.f = __ockl_image_sample_grad_2Da(i, s, + hc::short_vector::float4(x, y, layer, 0.0f).get_vector(), + hc::short_vector::float2(dx.x, dx.y).get_vector(), + hc::short_vector::float2(dy.x, dy.y).get_vector()); + TEXTURE_RETURN_FLOAT_XYZW; +} + template __TEXTURE_FUNCTIONS_DECL__ float4 tex2DLayeredGrad(texture texRef, hipTextureObject_t textureObject, float x, float y, int layer, float2 dx, float2 dy) { diff --git a/include/hip/hcc_detail/texture_types.h b/include/hip/hcc_detail/texture_types.h index 731ed12308..0a99abe451 100644 --- a/include/hip/hcc_detail/texture_types.h +++ b/include/hip/hcc_detail/texture_types.h @@ -93,6 +93,8 @@ struct textureReference float maxMipmapLevelClamp; hipTextureObject_t textureObject; + int numChannels; + enum hipArray_Format format; }; /** diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index a2ebcd27a5..a605406753 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -117,7 +117,9 @@ typedef struct hipDeviceProp_t { */ enum hipMemoryType { hipMemoryTypeHost, ///< Memory is physically located on host - hipMemoryTypeDevice ///< Memory is physically located on device. (see deviceId for specific device) + hipMemoryTypeDevice, ///< Memory is physically located on device. (see deviceId for specific device) + hipMemoryTypeArray, + hipMemoryTypeUnified }; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 8a5225d499..02f38497df 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -410,6 +410,113 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, hsa_ext_image_channel_order_t& channelOrder, hsa_ext_image_channel_type_t& channelType); +hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray ) +{ + HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray); + HIP_SET_DEVICE(); + hipError_t hip_status = hipSuccess; + if(pAllocateArray->width >0) { + auto ctx = ihipGetTlsDefaultCtx(); + + *array = (hipArray*)malloc(sizeof(hipArray)); + array[0]->drvDesc = *pAllocateArray; + array[0]->width = pAllocateArray->width; + array[0]->height = pAllocateArray->height; + array[0]->isDrv = true; + void ** ptr = &array[0]->data; + if (ctx) { + const unsigned am_flags = 0; + size_t size = pAllocateArray->width; + if(pAllocateArray->height > 0) { + size = size * pAllocateArray->height; + } + hsa_ext_image_channel_type_t channelType; + size_t allocSize = 0; + switch(pAllocateArray->format) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + allocSize = size * sizeof(unsigned char); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; + break; + case HIP_AD_FORMAT_UNSIGNED_INT16: + allocSize = size * sizeof(unsigned short); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; + break; + case HIP_AD_FORMAT_UNSIGNED_INT32: + allocSize = size * sizeof(unsigned int); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; + break; + case HIP_AD_FORMAT_SIGNED_INT8: + allocSize = size * sizeof(char); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; + break; + case HIP_AD_FORMAT_SIGNED_INT16: + allocSize = size * sizeof(short); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; + break; + case HIP_AD_FORMAT_SIGNED_INT32: + allocSize = size * sizeof(int); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32; + break; + case HIP_AD_FORMAT_HALF: + allocSize = size * sizeof(short); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT; + break; + case HIP_AD_FORMAT_FLOAT: + allocSize = size * sizeof(float); + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT; + break; + default: + hip_status = hipErrorUnknown; + break; + } + hc::accelerator acc = ctx->getDevice()->_acc; + hsa_agent_t* agent =static_cast(acc.get_hsa_agent()); + + size_t allocGranularity = 0; + hsa_amd_memory_pool_t *allocRegion = static_cast(acc.get_hsa_am_region()); + hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); + + hsa_ext_image_descriptor_t imageDescriptor; + + imageDescriptor.width = pAllocateArray->width; + imageDescriptor.height = pAllocateArray->height; + imageDescriptor.depth = 0; + imageDescriptor.array_size = 0; + + imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; + + hsa_ext_image_channel_order_t channelOrder; + + //getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType); + if (pAllocateArray->numChannels == 4) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; + } else if (pAllocateArray->numChannels == 2) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; + } else if (pAllocateArray->numChannels == 1) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; + } + imageDescriptor.format.channel_order = channelOrder; + imageDescriptor.format.channel_type = channelType; + + hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + hsa_ext_image_data_info_t imageInfo; + hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); + size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; + + *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment); + if (size && (*ptr == NULL)) { + hip_status = hipErrorMemoryAllocation; + } + + } else { + hip_status = hipErrorMemoryAllocation; + } + } else { + hip_status = hipErrorInvalidValue; + } + + return ihipLogStatus(hip_status); +} hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { @@ -425,6 +532,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, array[0]->height = height; array[0]->depth = 1; array[0]->desc = *desc; + array[0]->isDrv = false; void ** ptr = &array[0]->data; @@ -993,13 +1101,11 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h } // TODO - review and optimize -hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, - size_t width, size_t height, hipMemcpyKind kind) { - - HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind); - +hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, + size_t width, size_t height, hipMemcpyKind kind) +{ if(width > dpitch || width > spitch) - return ihipLogStatus(hipErrorUnknown); + return hipErrorUnknown; hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1016,6 +1122,26 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, e = ex._code; } + return e; +} + +hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, + size_t width, size_t height, hipMemcpyKind kind) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind); + hipError_t e = hipSuccess; + e = ihipMemcpy2D(dst,dpitch, src, spitch, width, height, kind); + return ihipLogStatus(e); +} + +hipError_t hipMemcpy_2D(const hip_Memcpy2D* pCopy) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy); + hipError_t e = hipSuccess; + if(pCopy == nullptr) { + e = hipErrorInvalidValue; + } + e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch, pCopy->widthInBytes, pCopy->height, hipMemcpyDefault); return ihipLogStatus(e); } diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 38411f2347..baad204582 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -40,6 +40,8 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" +//TODO Add support for multiple modules +static std::unordered_map coGlobals; //TODO Use Pool APIs from HCC to get memory regions. #include @@ -262,7 +264,9 @@ namespace symbol_section_accessor{self_reader, process_symtab}, x); assert(tmp.first); - + if (coGlobals.count(x) == 0) { + coGlobals.emplace(x, tmp.first); + } void* p = nullptr; hsa_amd_memory_lock( reinterpret_cast(tmp.first), tmp.second, &agent, 1, &p); @@ -833,3 +837,21 @@ hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image, unsigned { return hipModuleLoadData(module, image); } + +hipError_t hipModuleGetTexRef(hipDeviceptr_t *dptr, hipModule_t hmod, const char* name) +{ + HIP_INIT_API(dptr, hmod, name); + hipError_t ret = hipSuccess; + if(dptr == NULL){ + return ihipLogStatus(hipErrorInvalidValue); + } + if(name == NULL || hmod == NULL){ + return ihipLogStatus(hipErrorNotInitialized); + } + else{ + const auto it = coGlobals.find(name); + if (it == coGlobals.end()) return ihipLogStatus(hipErrorInvalidValue); + *dptr = reinterpret_cast(it->second); + return ihipLogStatus(ret); + } +} diff --git a/src/hip_texture.cpp b/src/hip_texture.cpp index 030772f65e..de6100e937 100644 --- a/src/hip_texture.cpp +++ b/src/hip_texture.cpp @@ -32,6 +32,48 @@ void saveTextureInfo(const hipTexture* pTexture, } } +void getDrvChannelOrderAndType(const enum hipArray_Format Format, + unsigned int NumChannels, + hsa_ext_image_channel_order_t& channelOrder, + hsa_ext_image_channel_type_t& channelType) +{ + switch(Format) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; + break; + case HIP_AD_FORMAT_UNSIGNED_INT16: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; + break; + case HIP_AD_FORMAT_UNSIGNED_INT32: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; + break; + case HIP_AD_FORMAT_SIGNED_INT8: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8; + break; + case HIP_AD_FORMAT_SIGNED_INT16: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16; + break; + case HIP_AD_FORMAT_SIGNED_INT32: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32; + break; + case HIP_AD_FORMAT_HALF: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT; + break; + case HIP_AD_FORMAT_FLOAT: + channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT; + break; + default: + break; + } + + if (NumChannels == 4) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; + } else if (NumChannels == 2) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; + } else if (NumChannels == 1) { + channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; + } +} void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureReadMode readMode, hsa_ext_image_channel_order_t& channelOrder, @@ -558,7 +600,11 @@ hipError_t ihipBindTextureToArrayImpl(int dim, hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; - getChannelOrderAndType(desc, readMode, channelOrder, channelType); + if(array->isDrv) { + getDrvChannelOrderAndType(array->drvDesc.format, array->drvDesc.numChannels, channelOrder, channelType); + } else { + getChannelOrderAndType(desc, readMode, channelOrder, channelType); + } imageDescriptor.format.channel_order = channelOrder; imageDescriptor.format.channel_type = channelType; @@ -666,3 +712,47 @@ hipError_t hipGetTextureReference(const textureReference** tex, const void* symb } return ihipLogStatus(hip_status); } + +hipError_t hipTexRefSetFormat (textureReference* tex, hipArray_Format fmt, int NumPackedComponents ) +{ + HIP_INIT_API(tex, fmt, NumPackedComponents); + hipError_t hip_status = hipSuccess; + tex->format = fmt; + tex->numChannels = NumPackedComponents; + return ihipLogStatus(hip_status); +} + +hipError_t hipTexRefSetFlags ( textureReference* tex, unsigned int flags ) +{ + HIP_INIT_API(tex, flags); + hipError_t hip_status = hipSuccess; + tex->normalized = flags; + return ihipLogStatus(hip_status); +} + +hipError_t hipTexRefSetFilterMode ( textureReference* tex, hipTextureFilterMode fm ) +{ + HIP_INIT_API(tex, fm); + hipError_t hip_status = hipSuccess; + tex->filterMode = fm; + return ihipLogStatus(hip_status); +} + +hipError_t hipTexRefSetAddressMode ( textureReference* tex, int dim, hipTextureAddressMode am ) +{ + HIP_INIT_API(tex, dim, am); + hipError_t hip_status = hipSuccess; + tex->addressMode[dim] = am; + return ihipLogStatus(hip_status); +} + +hipError_t hipTexRefSetArray ( textureReference* tex, hipArray_const_t array, unsigned int flags ) +{ + HIP_INIT_API(tex, array, flags); + hipError_t hip_status = hipSuccess; + + return ihipBindTextureToArrayImpl(hipTextureType2D, hipReadModeElementType, + array, array->desc,/* channelOrder, channelType,*/ + tex->addressMode[0], tex->filterMode, tex->normalized, + tex->textureObject); +} \ No newline at end of file From 68cc7766c6dd45914d3eaff11c8489b4530fd0eb Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 9 Nov 2017 22:35:29 +0530 Subject: [PATCH 02/19] Added texture 2D driver API usage example --- samples/2_Cookbook/11_texture_driver/Makefile | 17 ++ .../11_texture_driver/tex2dKernel.cpp | 33 ++++ .../11_texture_driver/texture2dDrv.cpp | 153 ++++++++++++++++++ 3 files changed, 203 insertions(+) create mode 100644 samples/2_Cookbook/11_texture_driver/Makefile create mode 100644 samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp create mode 100644 samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp diff --git a/samples/2_Cookbook/11_texture_driver/Makefile b/samples/2_Cookbook/11_texture_driver/Makefile new file mode 100644 index 0000000000..b68c5c31c7 --- /dev/null +++ b/samples/2_Cookbook/11_texture_driver/Makefile @@ -0,0 +1,17 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif +HIPCC=$(HIP_PATH)/bin/hipcc +HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) + +all: tex2dKernel.code texture2dDrv.out + +texture2dDrv.out: texture2dDrv.cpp + $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + +tex2dKernel.code: tex2dKernel.cpp + $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ + +clean: + rm -f *.code *.out diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp new file mode 100644 index 0000000000..17ed911808 --- /dev/null +++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -0,0 +1,33 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +extern texture tex; + +__global__ void tex2dKernel(hipLaunchParm lp, float* outputData, + int width, + int height) +{ + int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y; + outputData[y*width + x] = tex2D(tex, x, y); +} diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp new file mode 100644 index 0000000000..87e23d6e94 --- /dev/null +++ b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -0,0 +1,153 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include +#include + +#define fileName "tex2dKernel.code" + +texture tex; +bool testResult = true; + +#define HIP_CHECK(cmd) \ +{\ + hipError_t status = cmd;\ + if(status != hipSuccess) {std::cout<<"error: #"<