AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge in the rocclr based hip runtime (#2032)

Browse Source 
* Merge master-next changes in master (include vdi development in master branch)
This commit is contained in:
Maneesh Gupta
2020-04-23 21:42:06 +05:30
committed by GitHub
parent 386a0e0123
commit a0b5dfd625
136 changed files with 29756 additions and 307 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/channel_descriptor.h
+6
View File
@@ -29,8 +29,14 @@ THE SOFTWARE.
#ifdef __cplusplus
#if __HIP_VDI__
extern "C" {
#endif
HIP_PUBLIC_API
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
#if __HIP_VDI__
}
#endif
static inline hipChannelFormatDesc hipCreateChannelDescHalf() {
int e = (int)sizeof(unsigned short) * 8;
include/hip/hcc_detail/driver_types.h
+152 -20
View File
@@ -135,6 +135,47 @@ typedef enum hipResourceType {
hipResourceTypePitch2D = 0x03
}hipResourceType;
typedef enum HIPresourcetype_enum {
HIP_RESOURCE_TYPE_ARRAY = 0x00, /**< Array resoure */
HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY = 0x01, /**< Mipmapped array resource */
HIP_RESOURCE_TYPE_LINEAR = 0x02, /**< Linear resource */
HIP_RESOURCE_TYPE_PITCH2D = 0x03 /**< Pitch 2D resource */
} HIPresourcetype;
/**
* hip address modes
*/
typedef enum HIPaddress_mode_enum {
HIP_TR_ADDRESS_MODE_WRAP = 0,
HIP_TR_ADDRESS_MODE_CLAMP = 1,
HIP_TR_ADDRESS_MODE_MIRROR = 2,
HIP_TR_ADDRESS_MODE_BORDER = 3
} HIPaddress_mode;
/**
* hip filter modes
*/
typedef enum HIPfilter_mode_enum {
HIP_TR_FILTER_MODE_POINT = 0,
HIP_TR_FILTER_MODE_LINEAR = 1
} HIPfilter_mode;
/**
* Texture descriptor
*/
typedef struct HIP_TEXTURE_DESC_st {
HIPaddress_mode addressMode[3]; /**< Address modes */
HIPfilter_mode filterMode; /**< Filter mode */
unsigned int flags; /**< Flags */
unsigned int maxAnisotropy; /**< Maximum anisotropy ratio */
HIPfilter_mode mipmapFilterMode; /**< Mipmap filter mode */
float mipmapLevelBias; /**< Mipmap level bias */
float minMipmapLevelClamp; /**< Mipmap minimum level clamp */
float maxMipmapLevelClamp; /**< Mipmap maximum level clamp */
float borderColor[4]; /**< Border Color */
int reserved[12];
} HIP_TEXTURE_DESC;
/**
* hip texture resource view formats
*/
@@ -176,6 +217,45 @@ typedef enum hipResourceViewFormat {
hipResViewFormatUnsignedBlockCompressed7 = 0x22
}hipResourceViewFormat;
typedef enum HIPresourceViewFormat_enum
{
HIP_RES_VIEW_FORMAT_NONE = 0x00, /**< No resource view format (use underlying resource format) */
HIP_RES_VIEW_FORMAT_UINT_1X8 = 0x01, /**< 1 channel unsigned 8-bit integers */
HIP_RES_VIEW_FORMAT_UINT_2X8 = 0x02, /**< 2 channel unsigned 8-bit integers */
HIP_RES_VIEW_FORMAT_UINT_4X8 = 0x03, /**< 4 channel unsigned 8-bit integers */
HIP_RES_VIEW_FORMAT_SINT_1X8 = 0x04, /**< 1 channel signed 8-bit integers */
HIP_RES_VIEW_FORMAT_SINT_2X8 = 0x05, /**< 2 channel signed 8-bit integers */
HIP_RES_VIEW_FORMAT_SINT_4X8 = 0x06, /**< 4 channel signed 8-bit integers */
HIP_RES_VIEW_FORMAT_UINT_1X16 = 0x07, /**< 1 channel unsigned 16-bit integers */
HIP_RES_VIEW_FORMAT_UINT_2X16 = 0x08, /**< 2 channel unsigned 16-bit integers */
HIP_RES_VIEW_FORMAT_UINT_4X16 = 0x09, /**< 4 channel unsigned 16-bit integers */
HIP_RES_VIEW_FORMAT_SINT_1X16 = 0x0a, /**< 1 channel signed 16-bit integers */
HIP_RES_VIEW_FORMAT_SINT_2X16 = 0x0b, /**< 2 channel signed 16-bit integers */
HIP_RES_VIEW_FORMAT_SINT_4X16 = 0x0c, /**< 4 channel signed 16-bit integers */
HIP_RES_VIEW_FORMAT_UINT_1X32 = 0x0d, /**< 1 channel unsigned 32-bit integers */
HIP_RES_VIEW_FORMAT_UINT_2X32 = 0x0e, /**< 2 channel unsigned 32-bit integers */
HIP_RES_VIEW_FORMAT_UINT_4X32 = 0x0f, /**< 4 channel unsigned 32-bit integers */
HIP_RES_VIEW_FORMAT_SINT_1X32 = 0x10, /**< 1 channel signed 32-bit integers */
HIP_RES_VIEW_FORMAT_SINT_2X32 = 0x11, /**< 2 channel signed 32-bit integers */
HIP_RES_VIEW_FORMAT_SINT_4X32 = 0x12, /**< 4 channel signed 32-bit integers */
HIP_RES_VIEW_FORMAT_FLOAT_1X16 = 0x13, /**< 1 channel 16-bit floating point */
HIP_RES_VIEW_FORMAT_FLOAT_2X16 = 0x14, /**< 2 channel 16-bit floating point */
HIP_RES_VIEW_FORMAT_FLOAT_4X16 = 0x15, /**< 4 channel 16-bit floating point */
HIP_RES_VIEW_FORMAT_FLOAT_1X32 = 0x16, /**< 1 channel 32-bit floating point */
HIP_RES_VIEW_FORMAT_FLOAT_2X32 = 0x17, /**< 2 channel 32-bit floating point */
HIP_RES_VIEW_FORMAT_FLOAT_4X32 = 0x18, /**< 4 channel 32-bit floating point */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC1 = 0x19, /**< Block compressed 1 */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC2 = 0x1a, /**< Block compressed 2 */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC3 = 0x1b, /**< Block compressed 3 */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC4 = 0x1c, /**< Block compressed 4 unsigned */
HIP_RES_VIEW_FORMAT_SIGNED_BC4 = 0x1d, /**< Block compressed 4 signed */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC5 = 0x1e, /**< Block compressed 5 unsigned */
HIP_RES_VIEW_FORMAT_SIGNED_BC5 = 0x1f, /**< Block compressed 5 signed */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC6H = 0x20, /**< Block compressed 6 unsigned half-float */
HIP_RES_VIEW_FORMAT_SIGNED_BC6H = 0x21, /**< Block compressed 6 signed half-float */
HIP_RES_VIEW_FORMAT_UNSIGNED_BC7 = 0x22 /**< Block compressed 7 */
} HIPresourceViewFormat;
/**
* HIP resource descriptor
*/
@@ -204,6 +284,39 @@ typedef struct hipResourceDesc {
} res;
}hipResourceDesc;
typedef struct HIP_RESOURCE_DESC_st
{
HIPresourcetype resType; /**< Resource type */
union {
struct {
hipArray_t hArray; /**< HIP array */
} array;
struct {
hipMipmappedArray_t hMipmappedArray; /**< HIP mipmapped array */
} mipmap;
struct {
hipDeviceptr_t devPtr; /**< Device pointer */
hipArray_Format format; /**< Array format */
unsigned int numChannels; /**< Channels per array element */
size_t sizeInBytes; /**< Size in bytes */
} linear;
struct {
hipDeviceptr_t devPtr; /**< Device pointer */
hipArray_Format format; /**< Array format */
unsigned int numChannels; /**< Channels per array element */
size_t width; /**< Width of the array in elements */
size_t height; /**< Height of the array in elements */
size_t pitchInBytes; /**< Pitch between two rows in bytes */
} pitch2D;
struct {
int reserved[32];
} reserved;
} res;
unsigned int flags; /**< Flags (must be zero) */
} HIP_RESOURCE_DESC;
/**
* hip resource view descriptor
*/
@@ -218,6 +331,22 @@ struct hipResourceViewDesc {
unsigned int lastLayer;
};
/**
* Resource view descriptor
*/
typedef struct HIP_RESOURCE_VIEW_DESC_st
{
HIPresourceViewFormat format; /**< Resource view format */
size_t width; /**< Width of the resource view */
size_t height; /**< Height of the resource view */
size_t depth; /**< Depth of the resource view */
unsigned int firstMipmapLevel; /**< First defined mipmap level */
unsigned int lastMipmapLevel; /**< Last defined mipmap level */
unsigned int firstLayer; /**< First layer index */
unsigned int lastLayer; /**< Last layer index */
unsigned int reserved[16];
} HIP_RESOURCE_VIEW_DESC;
/**
* Memory copy types
*
@@ -263,26 +392,29 @@ typedef struct hipMemcpy3DParms {
} hipMemcpy3DParms;
typedef struct HIP_MEMCPY3D {
size_t Depth;
size_t Height;
size_t WidthInBytes;
hipDeviceptr_t dstDevice;
size_t dstHeight;
void* dstHost;
size_t dstLOD;
hipMemoryType dstMemoryType;
size_t dstPitch;
size_t dstXInBytes;
size_t dstY;
size_t dstZ;
void* reserved0;
void* reserved1;
hipDeviceptr_t srcDevice;
size_t srcHeight;
const void* srcHost;
size_t srcLOD;
hipMemoryType srcMemoryType;
size_t srcPitch;
unsigned int srcXInBytes;
unsigned int srcY;
unsigned int srcZ;
unsigned int srcLOD;
hipMemoryType srcMemoryType;
const void* srcHost;
hipDeviceptr_t srcDevice;
hipArray_t srcArray;
unsigned int srcPitch;
unsigned int srcHeight;
unsigned int dstXInBytes;
unsigned int dstY;
unsigned int dstZ;
unsigned int dstLOD;
hipMemoryType dstMemoryType;
void* dstHost;
hipDeviceptr_t dstDevice;
hipArray_t dstArray;
unsigned int dstPitch;
unsigned int dstHeight;
unsigned int WidthInBytes;
unsigned int Height;
unsigned int Depth;
} HIP_MEMCPY3D;
static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz,
include/hip/hcc_detail/functional_grid_launch.hpp
-10
View File
@@ -192,16 +192,6 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
stream, &config[0]);
}
inline
__attribute__((visibility("hidden")))
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
hip_impl::hip_init();
auto& ps = hip_impl::get_program_state();
return ihipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags, ps);
}
template <typename F>
inline
__attribute__((visibility("hidden")))
include/hip/hcc_detail/hip_fp16.h
+1 -1
View File
@@ -229,7 +229,7 @@ THE SOFTWARE.
__host__ __device__
operator __half_raw() const { return __half_raw{data}; }
__host__ __device__
operator volatile __half_raw() const volatile
operator __half_raw() const volatile
{
return __half_raw{data};
}
include/hip/hcc_detail/hip_runtime.h
+5 -2
View File
@@ -108,9 +108,12 @@ extern int HIP_TRACE_API;
#include <hip/hcc_detail/host_defines.h>
#include <hip/hcc_detail/device_functions.h>
#include <hip/hcc_detail/surface_functions.h>
#include <hip/hcc_detail/texture_functions.h>
#if __HCC__
#include <hip/hcc_detail/math_functions.h>
#include <hip/hcc_detail/texture_functions.h>
#else
#include <hip/hcc_detail/texture_fetch_functions.h>
#include <hip/hcc_detail/texture_indirect_functions.h>
#endif
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
#if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
@@ -385,7 +388,7 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
#elif defined(__clang__) && defined(__HIP__)
#define HIP_KERNEL_NAME(...) __VA_ARGS__
#define HIP_SYMBOL(X) #X
#define HIP_SYMBOL(X) X
typedef int hipLaunchParm;
include/hip/hcc_detail/hip_runtime_api.h
+505 -26
View File
@@ -1482,18 +1482,18 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name);
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName);
hipError_t hipGetSymbolSize(size_t* size, const void* symbolName);
hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol);
hipError_t hipGetSymbolSize(size_t* size, const void* symbol);
hipError_t hipMemcpyToSymbol(const void* symbol, const void* src,
size_t sizeBytes, size_t offset __dparm(0),
hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src,
hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src,
size_t sizeBytes, size_t offset,
hipMemcpyKind kind, hipStream_t stream __dparm(0));
hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName,
hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol,
size_t sizeBytes, size_t offset __dparm(0),
hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost));
hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol,
size_t sizeBytes, size_t offset,
hipMemcpyKind kind,
hipStream_t stream __dparm(0));
@@ -1933,6 +1933,15 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent);
*/
hipError_t hipFreeArray(hipArray* array);
/**
* @brief Frees a mipmapped array on the device
*
* @param[in] mipmappedArray - Pointer to mipmapped array to free
*
* @return #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray);
/**
* @brief Allocate an array on the device.
*
@@ -1947,6 +1956,39 @@ hipError_t hipFreeArray(hipArray* array);
hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc,
struct hipExtent extent, unsigned int flags);
/**
* @brief Allocate a mipmapped array on the device
*
* @param[out] mipmappedArray - Pointer to allocated mipmapped array in device memory
* @param[in] desc - Requested channel format
* @param[in] extent - Requested allocation size (width field in elements)
* @param[in] numLevels - Number of mipmap levels to allocate
* @param[in] flags - Flags for extensions
*
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryAllocation
*/
hipError_t hipMallocMipmappedArray(
hipMipmappedArray_t *mipmappedArray,
const struct hipChannelFormatDesc* desc,
struct hipExtent extent,
unsigned int numLevels,
unsigned int flags __dparm(0));
/**
* @brief Gets a mipmap level of a HIP mipmapped array
*
* @param[out] levelArray - Returned mipmap level HIP array
* @param[in] mipmappedArray - HIP mipmapped array
* @param[in] level - Mipmap level
*
* @return #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipGetMipmappedArrayLevel(
hipArray_t *levelArray,
hipMipmappedArray_const_t mipmappedArray,
unsigned int level);
/**
* @brief Copies data between host and device.
*
@@ -2159,6 +2201,31 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p);
*/
hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms* p, hipStream_t stream __dparm(0));
/**
* @brief Copies data between host and device.
*
* @param[in] pCopy 3D memory copy parameters
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
* #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
*
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol,
* hipMemcpyAsync
*/
hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy);
/**
* @brief Copies data between host and device asynchronously.
*
* @param[in] pCopy 3D memory copy parameters
* @param[in] stream Stream to use
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
* #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
*
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol,
* hipMemcpyAsync
*/
hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream);
// doxygen end Memory
/**
* @}
@@ -2957,17 +3024,6 @@ hipError_t hipModuleOccupancyMaxPotentialBlockSizeWithFlags(int* gridSize, int*
hipFunction_t f, size_t dynSharedMemPerBlk,
int blockSizeLimit, unsigned int flags);
/**
* @brief Returns occupancy for a device function.
*
* @param [out] numBlocks Returned occupancy
* @param [in] func Kernel function for which occupancy is calulated
* @param [in] blockSize Block size the kernel is intended to be launched with
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
*/
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk);
/**
* @brief Returns occupancy for a device function.
*
@@ -2979,6 +3035,29 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk);
/**
* @brief Returns occupancy for a device function.
*
* @param [out] numBlocks Returned occupancy
* @param [in] f Kernel function(hipFunction_t) for which occupancy is calulated
* @param [in] blockSize Block size the kernel is intended to be launched with
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
* @param [in] flags Extra flags for occupancy calculation (only default supported)
*/
hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
/**
* @brief Returns occupancy for a device function.
*
* @param [out] numBlocks Returned occupancy
* @param [in] func Kernel function for which occupancy is calulated
* @param [in] blockSize Block size the kernel is intended to be launched with
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
*/
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk);
/**
* @brief Returns occupancy for a device function.
*
@@ -2992,18 +3071,20 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags __dparm(hipOccupancyDefault));
/**
* @brief Returns occupancy for a device function.
* @brief determine the grid and block sizes to achieves maximum occupancy for a kernel
*
* @param [out] numBlocks Returned occupancy
* @param [in] f Kernel function(hipFunction_t) for which occupancy is calulated
* @param [in] blockSize Block size the kernel is intended to be launched with
* @param [out] gridSize minimum grid size for maximum potential occupancy
* @param [out] blockSize block size for maximum potential occupancy
* @param [in] f kernel function for which occupancy is calulated
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
* @param [in] flags Extra flags for occupancy calculation (only default supported)
* @param [in] blockSizeLimit the maximum block size for the kernel, use 0 for no limit
*
* @returns hipSuccess, hipInvalidDevice, hipErrorInvalidValue
*/
hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
const void* f, size_t dynSharedMemPerBlk,
int blockSizeLimit);
#if __HIP_VDI__ && !defined(__HCC__)
/**
* @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched
* on respective streams before enqueuing any other work on the specified streams from any other threads
@@ -3018,7 +3099,6 @@ hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags);
#endif
// doxygen end Version Management
/**
@@ -3260,6 +3340,206 @@ hipError_t hipLaunchKernel(const void* function_address,
size_t sharedMemBytes __dparm(0),
hipStream_t stream __dparm(0));
#if __HIP_VDI__
hipError_t hipBindTexture(
size_t* offset,
const textureReference* tex,
const void* devPtr,
const hipChannelFormatDesc* desc,
size_t size = UINT_MAX);
hipError_t hipBindTexture2D(
size_t* offset,
const textureReference* tex,
const void* devPtr,
const hipChannelFormatDesc* desc,
size_t width,
size_t height,
size_t pitch);
hipError_t hipBindTextureToArray(
const textureReference* tex,
hipArray_const_t array,
const hipChannelFormatDesc* desc);
hipError_t hipBindTextureToMipmappedArray(
const textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc);
hipError_t hipGetTextureAlignmentOffset(
size_t* offset,
const textureReference* texref);
hipError_t hipGetTextureReference(
const textureReference** texref,
const void* symbol);
hipError_t hipUnbindTexture(const textureReference* tex);
hipError_t hipCreateTextureObject(
hipTextureObject_t* pTexObject,
const hipResourceDesc* pResDesc,
const hipTextureDesc* pTexDesc,
const hipResourceViewDesc* pResViewDesc);
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
hipError_t hipGetChannelDesc(
hipChannelFormatDesc* desc,
hipArray_const_t array);
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 hipTexRefGetAddress(
hipDeviceptr_t* dev_ptr,
const textureReference* texRef);
hipError_t hipTexRefGetAddressMode(
hipTextureAddressMode* pam,
const textureReference* texRef,
int dim);
hipError_t hipTexRefGetFilterMode(
hipTextureFilterMode* pfm,
const textureReference* texRef);
hipError_t hipTexRefGetFlags(
unsigned int* pFlags,
const textureReference* texRef);
hipError_t hipTexRefGetFormat(
hipArray_Format* pFormat,
int* pNumChannels,
const textureReference* texRef);
hipError_t hipTexRefGetMaxAnisotropy(
int* pmaxAnsio,
const textureReference* texRef);
hipError_t hipTexRefGetMipmapFilterMode(
hipTextureFilterMode* pfm,
const textureReference* texRef);
hipError_t hipTexRefGetMipmapLevelBias(
float* pbias,
const textureReference* texRef);
hipError_t hipTexRefGetMipmapLevelClamp(
float* pminMipmapLevelClamp,
float* pmaxMipmapLevelClamp,
const textureReference* texRef);
hipError_t hipTexRefGetMipMappedArray(
hipMipmappedArray_t* pArray,
const textureReference* texRef);
hipError_t hipTexRefSetAddress(
size_t* ByteOffset,
textureReference* texRef,
hipDeviceptr_t dptr,
size_t bytes);
hipError_t hipTexRefSetAddress2D(
textureReference* texRef,
const HIP_ARRAY_DESCRIPTOR* desc,
hipDeviceptr_t dptr,
size_t Pitch);
hipError_t hipTexRefSetAddressMode(
textureReference* texRef,
int dim,
hipTextureAddressMode am);
hipError_t hipTexRefSetArray(
textureReference* tex,
hipArray_const_t array,
unsigned int flags);
hipError_t hipTexRefSetBorderColor(
textureReference* texRef,
float* pBorderColor);
hipError_t hipTexRefSetFilterMode(
textureReference* texRef,
hipTextureFilterMode fm);
hipError_t hipTexRefSetFlags(
textureReference* texRef,
unsigned int Flags);
hipError_t hipTexRefSetFormat(
textureReference* texRef,
hipArray_Format fmt,
int NumPackedComponents);
hipError_t hipTexRefSetMaxAnisotropy(
textureReference* texRef,
unsigned int maxAniso);
hipError_t hipTexRefSetMipmapFilterMode(
textureReference* texRef,
hipTextureFilterMode fm);
hipError_t hipTexRefSetMipmapLevelBias(
textureReference* texRef,
float bias);
hipError_t hipTexRefSetMipmapLevelClamp(
textureReference* texRef,
float minMipMapLevelClamp,
float maxMipMapLevelClamp);
hipError_t hipTexRefSetMipmappedArray(
textureReference* texRef,
hipMipmappedArray* mipmappedArray,
unsigned int Flags);
hipError_t hipMipmappedArrayCreate(
hipMipmappedArray_t* pHandle,
HIP_ARRAY3D_DESCRIPTOR* pMipmappedArrayDesc,
unsigned int numMipmapLevels);
hipError_t hipMipmappedArrayDestroy(
hipMipmappedArray_t hMipmappedArray);
hipError_t hipMipmappedArrayGetLevel(
hipArray_t* pLevelArray,
hipMipmappedArray_t hMipMappedArray,
unsigned int level);
hipError_t hipTexObjectCreate(
hipTextureObject_t* pTexObject,
const HIP_RESOURCE_DESC* pResDesc,
const HIP_TEXTURE_DESC* pTexDesc,
const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
hipError_t hipTexObjectDestroy(
hipTextureObject_t texObject);
hipError_t hipTexObjectGetResourceDesc(
HIP_RESOURCE_DESC* pResDesc,
hipTextureObject_t texObject);
hipError_t hipTexObjectGetResourceViewDesc(
HIP_RESOURCE_VIEW_DESC* pResViewDesc,
hipTextureObject_t texObject);
hipError_t hipTexObjectGetTextureDesc(
HIP_TEXTURE_DESC* pTexDesc,
hipTextureObject_t texObject);
#endif
/**
* @}
*/
@@ -3269,6 +3549,60 @@ hipError_t hipLaunchKernel(const void* function_address,
} /* extern "c" */
#endif
#if defined(__cplusplus) && !defined(__HCC__) && defined(__clang__) && defined(__HIP__)
template <typename T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
T f, size_t dynSharedMemPerBlk = 0, int blockSizeLimit = 0) {
return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, reinterpret_cast<const void*>(f),dynSharedMemPerBlk,blockSizeLimit);
}
template <typename T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeWithFlags(int* gridSize, int* blockSize,
T f, size_t dynSharedMemPerBlk = 0, int blockSizeLimit = 0, unsigned int flags = 0 ) {
return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, reinterpret_cast<const void*>(f),dynSharedMemPerBlk,blockSizeLimit);
}
#endif // defined(__cplusplus) && !defined(__HCC__) && defined(__clang__) && defined(__HIP__)
#if defined(__cplusplus) && !defined(__HCC__)
template <typename T>
hipError_t hipGetSymbolAddress(void** devPtr, const T &symbol) {
return ::hipGetSymbolAddress(devPtr, (const void *)&symbol);
}
template <typename T>
hipError_t hipGetSymbolSize(size_t* size, const T &symbol) {
return ::hipGetSymbolSize(size, (const void *)&symbol);
}
template <typename T>
hipError_t hipMemcpyToSymbol(const T& symbol, const void* src, size_t sizeBytes,
size_t offset __dparm(0),
hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)) {
return ::hipMemcpyToSymbol((const void*)&symbol, src, sizeBytes, offset, kind);
}
template <typename T>
hipError_t hipMemcpyToSymbolAsync(const T& symbol, const void* src, size_t sizeBytes, size_t offset,
hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
return ::hipMemcpyToSymbolAsync((const void*)&symbol, src, sizeBytes, offset, kind, stream);
}
template <typename T>
hipError_t hipMemcpyFromSymbol(void* dst, const T &symbol,
size_t sizeBytes, size_t offset __dparm(0),
hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)) {
return ::hipMemcpyFromSymbol(dst, (const void*)&symbol, sizeBytes, offset, kind);
}
template <typename T>
hipError_t hipMemcpyFromSymbolAsync(void* dst, const T& symbol, size_t sizeBytes, size_t offset,
hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
return ::hipMemcpyFromSymbolAsync(dst, (const void*)&symbol, sizeBytes, offset, kind, stream);
}
#endif
#if USE_PROF_API
#include <hip/hcc_detail/hip_prof_str.h>
#endif
@@ -3307,12 +3641,16 @@ inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
class TlsData;
#if !__HIP_VDI__
hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
const hipChannelFormatDesc* desc, size_t size = UINT_MAX);
#endif
#if !__HIP_VDI__
hipError_t ihipBindTextureImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, size_t* offset,
const void* devPtr, const struct hipChannelFormatDesc* desc,
size_t size, textureReference* tex);
#endif
/*
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture
@@ -3329,11 +3667,13 @@ hipError_t ihipBindTextureImpl(TlsData *tls, int dim, enum hipTextureReadMode re
* @param[in] size - Size of the memory area pointed to by devPtr
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
**/
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex, const void* devPtr,
const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) {
return ihipBindTextureImpl(nullptr, dim, readMode, offset, devPtr, &desc, size, &tex);
}
#endif
/*
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture
@@ -3349,81 +3689,114 @@ hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex,
* @param[in] size - Size of the memory area pointed to by devPtr
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
**/
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex, const void* devPtr,
size_t size = UINT_MAX) {
return ihipBindTextureImpl(nullptr, dim, readMode, offset, devPtr, &(tex.channelDesc), size, &tex);
}
#endif
// C API
#if !__HIP_VDI__
hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr,
const hipChannelFormatDesc* desc, size_t width, size_t height,
size_t pitch);
#endif
#if !__HIP_VDI__
hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size_t* offset,
const void* devPtr, const struct hipChannelFormatDesc* desc,
size_t width, size_t height, textureReference* tex, size_t pitch);
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& tex,
const void* devPtr, size_t width, size_t height, size_t pitch) {
return ihipBindTexture2DImpl(dim, readMode, offset, devPtr, &(tex.channelDesc), width, height,
&tex);
}
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& tex,
const void* devPtr, const struct hipChannelFormatDesc& desc,
size_t width, size_t height, size_t pitch) {
return ihipBindTexture2DImpl(dim, readMode, offset, devPtr, &desc, width, height, &tex);
}
#endif
// C API
#if !__HIP_VDI__
hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
const hipChannelFormatDesc* desc);
#endif
#if !__HIP_VDI__
hipError_t ihipBindTextureToArrayImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode,
hipArray_const_t array,
const struct hipChannelFormatDesc& desc,
textureReference* tex);
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex, hipArray_const_t array) {
return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, tex.channelDesc, &tex);
}
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex, hipArray_const_t array,
const struct hipChannelFormatDesc& desc) {
return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, desc, &tex);
}
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
hipArray_const_t array,
const struct hipChannelFormatDesc* desc) {
return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, *desc, tex);
}
#endif
// C API
#if !__HIP_VDI__
hipError_t hipBindTextureToMipmappedArray(const textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc);
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
hipMipmappedArray_const_t mipmappedArray) {
return hipSuccess;
}
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc& desc) {
return hipSuccess;
}
#endif
#if __HIP_VDI__ && !defined(__HCC__)
template <typename F>
inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) {
return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize,(hipFunction_t)kernel, dynSharedMemPerBlk, blockSizeLimit);
}
template <class T>
inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,
void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) {
@@ -3453,15 +3826,22 @@ inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchPara
*
* @return #hipSuccess
**/
#if !__HIP_VDI__
hipError_t hipUnbindTexture(const textureReference* tex);
#endif
#if !__HIP_VDI__
extern hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject);
#endif
#if !__HIP_VDI__
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipUnbindTexture(struct texture<T, dim, readMode>& tex) {
return ihipUnbindTextureImpl(tex.textureObject);
}
#endif
#if !__HIP_VDI__
hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array);
hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref);
hipError_t hipGetTextureReference(const textureReference** texref, const void* symbol);
@@ -3499,11 +3879,110 @@ hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex);
hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc,
hipDeviceptr_t devPtr, size_t pitch);
#endif
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
#if __HIP_VDI__
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTexture(
size_t *offset,
const struct texture<T, dim, readMode> &tex,
const void *devPtr,
size_t size = UINT_MAX)
{
return hipBindTexture(offset, tex, devPtr, tex.channelDesc, size);
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTexture(
size_t *offset,
const struct texture<T, dim, readMode> &tex,
const void *devPtr,
const struct hipChannelFormatDesc &desc,
size_t size = UINT_MAX)
{
return hipBindTexture(offset, &tex, devPtr, &desc, size);
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTexture2D(
size_t *offset,
const struct texture<T, dim, readMode> &tex,
const void *devPtr,
size_t width,
size_t height,
size_t pitch)
{
return hipBindTexture2D(offset, &tex, devPtr, &tex.channelDesc, width, height, pitch);
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTexture2D(
size_t *offset,
const struct texture<T, dim, readMode> &tex,
const void *devPtr,
const struct hipChannelFormatDesc &desc,
size_t width,
size_t height,
size_t pitch)
{
return hipBindTexture2D(offset, &tex, devPtr, &desc, width, height, pitch);
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTextureToArray(
const struct texture<T, dim, readMode> &tex,
hipArray_const_t array)
{
struct cudaChannelFormatDesc desc;
hipError_t err = hipGetChannelDesc(&desc, array);
return (err == hipSuccess) ? hipBindTextureToArray(tex, array, desc) : err;
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTextureToArray(
const struct texture<T, dim, readMode> &tex,
hipArray_const_t array,
const struct hipChannelFormatDesc &desc)
{
return hipBindTextureToArray(&tex, array, &desc);
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTextureToMipmappedArray(
const struct texture<T, dim, readMode> &tex,
hipMipmappedArray_const_t mipmappedArray)
{
struct hipChannelFormatDesc desc;
hipArray_t levelArray;
hipError_t err = hipGetMipmappedArrayLevel(&levelArray, mipmappedArray, 0);
if (err != hipSuccess) {
return err;
}
err = hipGetChannelDesc(&desc, levelArray);
return (err == hipSuccess) ? hipBindTextureToMipmappedArray(tex, mipmappedArray, desc) : err;
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipBindTextureToMipmappedArray(
const struct texture<T, dim, readMode> &tex,
hipMipmappedArray_const_t mipmappedArray,
const struct cudaChannelFormatDesc &desc)
{
return hipBindTextureToMipmappedArray(&tex, mipmappedArray, &desc);
}
template<class T, int dim, enum hipTextureReadMode readMode>
static inline hipError_t hipUnbindTexture(
const struct texture<T, dim, readMode> &tex)
{
return hipUnbindTexture(&tex);
}
#endif
// doxygen end Texture
/**
* @}
include/hip/hcc_detail/hip_texture_types.h
+4 -2
View File
@@ -57,25 +57,27 @@ struct __HIP_TEXTURE_ATTRIB texture : public textureReference {
texture(int norm = 0, enum hipTextureFilterMode fMode = hipFilterModePoint,
enum hipTextureAddressMode aMode = hipAddressModeClamp) {
normalized = norm;
readMode = hipReadModeNormalizedFloat;
readMode = mode;
filterMode = fMode;
addressMode[0] = aMode;
addressMode[1] = aMode;
addressMode[2] = aMode;
channelDesc = hipCreateChannelDesc<T>();
sRGB = 0;
textureObject = nullptr;
}
texture(int norm, enum hipTextureFilterMode fMode, enum hipTextureAddressMode aMode,
struct hipChannelFormatDesc desc) {
normalized = norm;
readMode = hipReadModeNormalizedFloat;
readMode = mode;
filterMode = fMode;
addressMode[0] = aMode;
addressMode[1] = aMode;
addressMode[2] = aMode;
channelDesc = desc;
sRGB = 0;
textureObject = nullptr;
}
};
include/hip/hcc_detail/hiprtc.h
+4
View File
@@ -28,7 +28,9 @@ extern "C" {
#include <stdlib.h>
#if !defined(_WIN32)
#pragma GCC visibility push (default)
#endif
enum hiprtcResult {
HIPRTC_SUCCESS = 0,
@@ -81,7 +83,9 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
#if !defined(_WIN32)
#pragma GCC visibility pop
#endif
#ifdef __cplusplus
}
include/hip/hcc_detail/ockl_image.h
+135
View File
@@ -0,0 +1,135 @@
/*
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.
*/
#pragma once
#include <hip/hip_vector_types.h>
extern "C" {
#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(4)))
__device__ float4::Native_vec_ __ockl_image_load_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, int c);
__device__ float4::Native_vec_ __ockl_image_load_1Db(unsigned int ADDRESS_SPACE_CONSTANT*i, int c);
__device__ float4::Native_vec_ __ockl_image_load_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_load_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_load_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_load_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_load_CM(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, int f);
__device__ float4::Native_vec_ __ockl_image_load_CMa(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int f);
__device__ float4::Native_vec_ __ockl_image_load_lod_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, int c, int l);
__device__ float4::Native_vec_ __ockl_image_load_lod_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, int l);
__device__ float4::Native_vec_ __ockl_image_load_lod_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, int l);
__device__ float4::Native_vec_ __ockl_image_load_lod_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int l);
__device__ float4::Native_vec_ __ockl_image_load_lod_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int l);
__device__ float4::Native_vec_ __ockl_image_load_lod_CM(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, int f, int l);
__device__ float4::Native_vec_ __ockl_image_load_lod_CMa(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int f, int l);
__device__ void __ockl_image_store_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, int c, float4::Native_vec_ p);
__device__ void __ockl_image_store_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, float4::Native_vec_ p);
__device__ void __ockl_image_store_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, float4::Native_vec_ p);
__device__ void __ockl_image_store_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
__device__ void __ockl_image_store_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
__device__ void __ockl_image_store_CM(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
__device__ void __ockl_image_store_CMa(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, int c, int l, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, int l, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, int l, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int l, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int l, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_CM(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int l, float4::Native_vec_ p);
__device__ void __ockl_image_store_lod_CMa(unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, int l, float4::Native_vec_ p);
__device__ float4::Native_vec_ __ockl_image_sample_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float c);
__device__ float4::Native_vec_ __ockl_image_sample_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_sample_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_sample_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_sample_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_sample_CM(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_sample_CMa(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_sample_grad_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float c, float dx, float dy);
__device__ float4::Native_vec_ __ockl_image_sample_grad_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c, float dx, float dy);
__device__ float4::Native_vec_ __ockl_image_sample_grad_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c, float2::Native_vec_ dx, float2::Native_vec_ dy);
__device__ float4::Native_vec_ __ockl_image_sample_grad_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float2::Native_vec_ dx, float2::Native_vec_ dy);
__device__ float4::Native_vec_ __ockl_image_sample_grad_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float4::Native_vec_ dx, float4::Native_vec_ dy);
__device__ float4::Native_vec_ __ockl_image_sample_lod_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float c, float l);
__device__ float4::Native_vec_ __ockl_image_sample_lod_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c, float l);
__device__ float4::Native_vec_ __ockl_image_sample_lod_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c, float l);
__device__ float4::Native_vec_ __ockl_image_sample_lod_2Da(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float l);
__device__ float4::Native_vec_ __ockl_image_sample_lod_3D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float l);
__device__ float4::Native_vec_ __ockl_image_sample_lod_CM(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float l);
__device__ float4::Native_vec_ __ockl_image_sample_lod_CMa(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float l);
__device__ float4::Native_vec_ __ockl_image_gather4r_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_gather4g_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_gather4b_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
__device__ float4::Native_vec_ __ockl_image_gather4a_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
};
include/hip/hcc_detail/texture_fetch_functions.h
+386
View File
@@ -0,0 +1,386 @@
/*
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.
*/
#pragma once
#if defined(__cplusplus)
#include <hip/hip_vector_types.h>
#include <hip/texture_types.h>
#include <hip/hcc_detail/ockl_image.h>
#include <type_traits>
#define TEXTURE_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
template<typename T>
struct __hip_is_tex_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_tex_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_tex_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
};
template<typename T>
struct __hip_is_tex_normalized_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_tex_normalized_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
};
template <
typename T,
hipTextureReadMode readMode,
typename Enable = void>
struct __hip_tex_ret
{
static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
};
template <
typename T,
hipTextureReadMode readMode>
using __hip_tex_ret_t = typename __hip_tex_ret<T, readMode, bool>::type;
template <typename T>
struct __hip_tex_ret<
T,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
{
using type = T;
};
template<
typename T,
unsigned int rank>
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeElementType>, rank>;
};
template<typename T>
struct __hip_tex_ret<
T,
hipReadModeNormalizedFloat,
typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
{
using type = float;
};
template<
typename T,
unsigned int rank>
struct __hip_tex_ret<
HIP_vector_type<T, rank>,
hipReadModeNormalizedFloat,
typename std::enable_if<__hip_is_tex_normalized_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<__hip_tex_ret_t<T, hipReadModeNormalizedFloat>, rank>;
};
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> t, int x)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_load_1Db(i, x);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1D(texture<T, hipTextureType1D, readMode> t, float x)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_1D(i, s, x);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2D(texture<T, hipTextureType2D, readMode> t, float x, float y)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayered(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayered(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3D(texture<T, hipTextureType3D, readMode> t, float x, float y, float z)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemap(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLod(texture<T, hipTextureType1D, readMode> t, float x, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLod(texture<T, hipTextureType2D, readMode> t, float x, float y, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredLod(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_1Da(i, s, float2(x, layer).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredLod(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_2Da(i, s, float4(x, y, layer, 0.0f).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DLod(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLod(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayered(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredLod(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float level)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
// return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return {};
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
// return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
return {};
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DGrad(texture<T, hipTextureType1D, readMode> t, float x, float dPdx, float dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DGrad(texture<T, hipTextureType2D, readMode> t, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_2D(i, s, float2(x, y).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1DLayeredGrad(texture<T, hipTextureType1DLayered, readMode> t, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex2DLayeredGrad(texture<T, hipTextureType2DLayered, readMode> t, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_2Da(i, s, float4(x, y, layer, 0.0f).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex3DGrad(texture<T, hipTextureType3D, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
auto tmp = __ockl_image_sample_grad_3D(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
return *reinterpret_cast<__hip_tex_ret_t<T, readMode>*>(&tmp);
}
template <
typename T,
hipTextureReadMode readMode,
typename Enable = void>
struct __hip_tex2dgather_ret
{
static_assert(std::is_same<Enable, void>::value, "Invalid channel type!");
};
template <
typename T,
hipTextureReadMode readMode>
using __hip_tex2dgather_ret_t = typename __hip_tex2dgather_ret<T, readMode, bool>::type;
template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<T>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
template<
typename T,
unsigned int rank>
struct __hip_tex2dgather_ret<
HIP_vector_type<T, rank>,
hipReadModeElementType,
typename std::enable_if<__hip_is_tex_channel_type<HIP_vector_type<T, rank>>::value, bool>::type>
{
using type = HIP_vector_type<T, 4>;
};
template <typename T>
struct __hip_tex2dgather_ret<
T,
hipReadModeNormalizedFloat,
typename std::enable_if<__hip_is_tex_normalized_channel_type<T>::value, bool>::type>
{
using type = float4;
};
template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_tex2dgather_ret_t<T, readMode> tex2Dgather(texture<T, hipTextureType2D, readMode> t, float x, float y, int comp=0)
{
TEXTURE_PARAMETERS_INIT;
switch (comp) {
case 1: {
auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
}
case 2: {
auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
}
case 3: {
auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
}
default: {
auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
return *reinterpret_cast<__hip_tex2dgather_ret_t<T, readMode>*>(&tmp);
}
}
return {};
}
#endif
include/hip/hcc_detail/texture_indirect_functions.h
+501
View File
@@ -0,0 +1,501 @@
/*
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.
*/
#pragma once
#if defined(__cplusplus)
#include <hip/hip_vector_types.h>
#include <hip/hip_texture_types.h>
#include <hip/hcc_detail/ockl_image.h>
#include <type_traits>
#define TEXTURE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
template<typename T>
struct __hip_is_itex_channel_type
{
static constexpr bool value =
std::is_same<T, char>::value ||
std::is_same<T, unsigned char>::value ||
std::is_same<T, short>::value ||
std::is_same<T, unsigned short>::value ||
std::is_same<T, int>::value ||
std::is_same<T, unsigned int>::value ||
std::is_same<T, float>::value;
};
template<
typename T,
unsigned int rank>
struct __hip_is_itex_channel_type<HIP_vector_type<T, rank>>
{
static constexpr bool value =
__hip_is_itex_channel_type<T>::value &&
((rank == 1) ||
(rank == 2) ||
(rank == 4));
};
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_load_1Db(i, x);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1Dfetch(T *ptr, hipTextureObject_t textureObject, int x)
{
*ptr = tex1Dfetch<T>(textureObject, x);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1D(hipTextureObject_t textureObject, float x)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1D(i, s, x);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1D(T *ptr, hipTextureObject_t textureObject, float x)
{
*ptr = tex1D<T>(textureObject, x);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2D(hipTextureObject_t textureObject, float x, float y)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2D(T *ptr, hipTextureObject_t textureObject, float x, float y)
{
*ptr = tex2D<T>(textureObject, x, y);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3D(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_3D(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3D(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = tex3D<T>(textureObject, x, y, z);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayered(hipTextureObject_t textureObject, float x, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayered(T *ptr, hipTextureObject_t textureObject, float x, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, layer);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayered(hipTextureObject_t textureObject, float x, float y, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer)
{
*ptr = tex1DLayered<T>(textureObject, x, y, layer);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemap(hipTextureObject_t textureObject, float x, float y, float z)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_CM(i, s, float4(x, y, z, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemap(T *ptr, hipTextureObject_t textureObject, float x, float y, float z)
{
*ptr = texCubemap<T>(textureObject, x, y, z);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayered(hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_CMa(i, s, float4(x, y, z, layer).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayered(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, z, layer);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2Dgather(hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
TEXTURE_OBJECT_PARAMETERS_INIT
switch (comp) {
case 1: {
auto tmp = __ockl_image_gather4r_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
break;
}
case 2: {
auto tmp = __ockl_image_gather4g_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
break;
}
case 3: {
auto tmp = __ockl_image_gather4b_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
break;
}
default: {
auto tmp = __ockl_image_gather4a_2D(i, s, float2(x, y).data);
return *reinterpret_cast<T*>(&tmp);
break;
}
};
return {};
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2Dgather(T *ptr, hipTextureObject_t textureObject, float x, float y, int comp = 0)
{
*ptr = texCubemapLayered<T>(textureObject, x, y, comp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLod(hipTextureObject_t textureObject, float x, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_1D(i, s, x, level);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLod(T *ptr, hipTextureObject_t textureObject, float x, float level)
{
*ptr = tex1DLod<T>(textureObject, x, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLod(hipTextureObject_t textureObject, float x, float y, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_2D(i, s, float2(x, y).data, level);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float level)
{
*ptr = tex2DLod<T>(textureObject, x, y, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3DLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_3D(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3DLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = tex3DLod<T>(textureObject, x, y, z, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_1Da(i, s, float2(x, layer).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, int layer, float level)
{
*ptr = tex1DLayeredLod<T>(textureObject, x, layer, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_2Da(i, s, float4(x, y, layer, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
*ptr = tex2DLayeredLod<T>(textureObject, x, y, layer, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLod(hipTextureObject_t textureObject, float x, float y, float z, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_CM(i, s, float4(x, y, z, 0.0f).data, level);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float level)
{
*ptr = texCubemapLod<T>(textureObject, x, y, z, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
// return *reinterpret_cast<T*>(&tmp);
return {};
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapGrad<T>(textureObject, x, y, z, dPdx, dPdy);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayeredLod(hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_lod_CMa(i, s, float4(x, y, z, layer).data, level);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayeredLod(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float level)
{
*ptr = texCubemapLayeredLod<T>(textureObject, x, y, z, layer, level);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DGrad(hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_1D(i, s, x, dPdx, dPdy);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DGrad(T *ptr, hipTextureObject_t textureObject, float x, float dPdx, float dPdy)
{
*ptr = tex1DGrad<T>(textureObject, x, dPdx, dPdy);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DGrad(hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_2D(i, s, float2(x, y).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float2 dPdx, float2 dPdy)
{
*ptr = tex2DGrad<T>(textureObject, x, y, dPdx, dPdy);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_3D(i, s, float4(x, y, z, 0.0f).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex3DGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
*ptr = tex3DGrad<T>(textureObject, x, y, z, dPdx, dPdy);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex1DLayeredGrad(hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_1Da(i, s, float2(x, layer).data, dPdx, dPdy);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex1DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, int layer, float dPdx, float dPdy)
{
*ptr = tex1DLayeredGrad<T>(textureObject, x, layer, dPdx, dPdy);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T tex2DLayeredGrad(hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
auto tmp = __ockl_image_sample_grad_2Da(i, s, float4(x, y, layer, 0.0f).data, float2(dPdx.x, dPdx.y).data, float2(dPdy.x, dPdy.y).data);
return *reinterpret_cast<T*>(&tmp);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void tex2DLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, int layer, float2 dPdx, float2 dPdy)
{
*ptr = tex2DLayeredGrad<T>(textureObject, x, y, layer, dPdx, dPdy);
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, float4(x, y, z, layer).data, float4(dPdx.x, dPdx.y, dPdx.z, 0.0f).data, float4(dPdy.x, dPdy.y, dPdy.z, 0.0f).data);
// return *reinterpret_cast<T*>(&tmp);
return {};
}
template <
typename T,
typename std::enable_if<__hip_is_itex_channel_type<T>::value>::type* = nullptr>
static __device__ void texCubemapLayeredGrad(T *ptr, hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
*ptr = texCubemapLayeredGrad<T>(textureObject, x, y, z, layer, dPdx, dPdy);
}
#endif
include/hip/hip_ext.h
+3 -1
View File
@@ -109,7 +109,9 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipEvent_t stopEvent = nullptr)
__attribute__((deprecated("use hipExtModuleLaunchKernel instead")));
#if !__HIP_VDI__ && defined(__cplusplus)
//#if !__HIP_VDI__ && defined(__cplusplus)
#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
//kernel_descriptor and hip_impl::make_kernarg are in "grid_launch_GGL.hpp"
namespace hip_impl {
inline
include/hip/hip_runtime_api.h
+2 -3
View File
@@ -336,7 +336,7 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributeTexturePitchAlignment, ///<Pitch alignment requirement for 2D texture references bound to pitched memory;
hipDeviceAttributeKernelExecTimeout, ///<Run time limit for kernels executed on the device
hipDeviceAttributeCanMapHostMemory, ///<Device can map host memory into device address space
hipDeviceAttributeEccEnabled, ///<Device has ECC support enabled
hipDeviceAttributeEccEnabled, ///<Device has ECC support enabled
hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc, ///< Supports cooperative launch on multiple
///devices with unmatched functions
@@ -344,9 +344,8 @@ typedef enum hipDeviceAttribute_t {
///devices with unmatched grid dimensions
hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim, ///< Supports cooperative launch on multiple
///devices with unmatched block dimensions
hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem, ///< Supports cooperative launch on multiple
hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem ///< Supports cooperative launch on multiple
///devices with unmatched shared memories
} hipDeviceAttribute_t;
enum hipComputeMode {