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

Texture driver APIs support

Browse Source
This commit is contained in:
Rahul Garg
2017-11-09 22:10:55 +05:30
parent 31bcb59f62
commit ef09c4918d
8 changed files with 4114 additions and 17 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/driver_types.h
+46
View File
@@ -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;
include/hip/hcc_detail/hip_runtime_api.h
+17 -7
View File
@@ -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
/**
include/hip/hcc_detail/texture_functions.h
+3800 -1
View File
File diff suppressed because it is too large Load Diff
include/hip/hcc_detail/texture_types.h
+2
View File
@@ -93,6 +93,8 @@ struct textureReference
float maxMipmapLevelClamp;
hipTextureObject_t textureObject;
int numChannels;
enum hipArray_Format format;
};
/**
include/hip/hip_runtime_api.h
+3 -1
View File
@@ -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
};
src/hip_memory.cpp
+132 -6
View File
@@ -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<hsa_agent_t*>(acc.get_hsa_agent());
size_t allocGranularity = 0;
hsa_amd_memory_pool_t *allocRegion = static_cast<hsa_amd_memory_pool_t*>(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);
}
src/hip_module.cpp
+23 -1
View File
@@ -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<std::string, uintptr_t> coGlobals;
//TODO Use Pool APIs from HCC to get memory regions.
#include <cassert>
@@ -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<void*>(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<void*>(it->second);
return ihipLogStatus(ret);
}
}
src/hip_texture.cpp
+91 -1
View File
@@ -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);
}