HIP Texture Support

src/hip_memory.cpp

@@ -27,10 +27,6 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 #include "hip_hcc_internal.h"
 #include "trace_helper.h"
-#include "hip/hcc_detail/hip_texture.h"
-#include <hc_am.hpp>
-
-
 
 // Internal HIP APIS:
 namespace hip_internal {
@@ -377,10 +373,11 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
     auto ctx = ihipGetTlsDefaultCtx();
 
     *array = (hipArray*)malloc(sizeof(hipArray));
+    array[0]->type  = flags;
     array[0]->width = width;
     array[0]->height = height;
-
-    array[0]->f = desc->f;
+    array[0]->depth = 1;
+    array[0]->desc = *desc;
 
     void ** ptr = &array[0]->data;
 
@@ -418,6 +415,62 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
     return ihipLogStatus(hip_status);
 }
 
+hipError_t hipMalloc3DArray(hipArray_t *array,
+                            const struct hipChannelFormatDesc* desc,
+                            struct hipExtent extent,
+                            unsigned int flags)
+{
+    HIP_INIT();
+    HIP_SET_DEVICE();
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+
+    *array = (hipArray*)malloc(sizeof(hipArray));
+    array[0]->type  = flags;
+    array[0]->width = extent.width;
+    array[0]->height = extent.height;
+    array[0]->depth = extent.depth;
+    array[0]->desc = *desc;
+
+    void ** ptr = &array[0]->data;
+
+    if (ctx) {
+        const unsigned am_flags = 0;
+        const size_t size = extent.width*extent.height*extent.depth;
+
+        size_t allocSize = 0;
+        switch(desc->f) {
+            case hipChannelFormatKindSigned:
+                allocSize = size * sizeof(int);
+                break;
+            case hipChannelFormatKindUnsigned:
+                allocSize = size * sizeof(unsigned int);
+                break;
+            case hipChannelFormatKindFloat:
+                allocSize = size * sizeof(float);
+                break;
+            case hipChannelFormatKindNone:
+                allocSize = size * sizeof(size_t);
+                break;
+            default:
+                hip_status = hipErrorUnknown;
+                break;
+        }
+        *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0);
+
+        if (size && (*ptr == NULL)) {
+            hip_status = hipErrorMemoryAllocation;
+        }
+
+    } else {
+        hip_status = hipErrorMemoryAllocation;
+    }
+
+    //return ihipLogStatus(hip_status);
+    return hip_status;
+}
+
 hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr)
 {
     HIP_INIT_API(flagsPtr, hostPtr);
@@ -857,7 +910,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
 
     size_t byteSize;
     if(dst) {
-        switch(dst[0].f) {
+        switch(dst[0].desc.f) {
             case hipChannelFormatKindSigned:
                 byteSize = sizeof(int);
                 break;
@@ -918,6 +971,56 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
     return ihipLogStatus(e);
 }
 
+hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
+{
+    HIP_INIT_SPECIAL_API((TRACE_MCMD), p);
+
+    hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
+
+    hc::completion_future marker;
+
+    hipError_t e = hipSuccess;
+
+    size_t byteSize;
+    if(p) {
+        switch(p->dstArray->desc.f) {
+            case hipChannelFormatKindSigned:
+                byteSize = sizeof(int);
+                break;
+            case hipChannelFormatKindUnsigned:
+                byteSize = sizeof(unsigned int);
+                break;
+            case hipChannelFormatKindFloat:
+                byteSize = sizeof(float);
+                break;
+            case hipChannelFormatKindNone:
+                byteSize = sizeof(size_t);
+                break;
+            default:
+                byteSize = 0;
+                break;
+        }
+    } else {
+        return ihipLogStatus(hipErrorUnknown);
+    }
+
+    try {
+        for (int i = 0; i < p->extent.depth; i++) {
+            for(int j = 0; j < p->extent.height; j++) {
+                // TODO: p->srcPos or p->dstPos are not 0.
+                unsigned char* src = (unsigned char*)p->srcPtr.ptr + i*p->srcPtr.ysize*p->srcPtr.pitch + j*p->srcPtr.pitch;
+                unsigned char* dst = (unsigned char*)p->dstArray->data + i*p->dstArray->height*p->dstArray->width*byteSize + j*p->dstArray->width*byteSize;
+                stream->locked_copySync(dst, src, p->extent.width*byteSize, p->kind);
+			}
+        }
+    }
+    catch (ihipException ex) {
+        e = ex._code;
+    }
+
+    return ihipLogStatus(e);
+}
+
 // TODO - make member function of stream?
 template <typename T>
 void
@@ -1006,7 +1109,7 @@ hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t s
     return ihipLogStatus(e);
 };
 
-hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
+hipError_t hipMemset(void* dst, int value, size_t sizeBytes)
 {
     HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
 
@@ -1058,6 +1161,58 @@ hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
     return ihipLogStatus(e);
 }
 
+hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height)
+{
+   HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, pitch, value, width, height);
+
+    hipError_t e = hipSuccess;
+
+    hipStream_t stream = hipStreamNull;
+    // TODO - call an ihip memset so HIP_TRACE is correct.
+    stream =  ihipSyncAndResolveStream(stream);
+
+    if (stream) {
+        auto crit = stream->lockopen_preKernelCommand();
+
+        hc::completion_future cf ;
+
+        size_t sizeBytes = pitch * height;
+        if ((sizeBytes & 0x3) == 0) {
+            // use a faster dword-per-workitem copy:
+            try {
+                value = value & 0xff;
+                uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
+                ihipMemsetKernel<uint32_t> (stream, crit, static_cast<uint32_t*> (dst), value32, sizeBytes/sizeof(uint32_t), &cf);
+            }
+            catch (std::exception &ex) {
+                e = hipErrorInvalidValue;
+            }
+        } else {
+            // use a slow byte-per-workitem copy:
+            try {
+                ihipMemsetKernel<char> (stream, crit, static_cast<char*> (dst), value, sizeBytes, &cf);
+            }
+            catch (std::exception &ex) {
+                e = hipErrorInvalidValue;
+            }
+        }
+        // TODO - is hipMemset supposed to be async?
+        cf.wait();
+
+        stream->lockclose_postKernelCommand("hipMemset", &crit->_av);
+
+        if (HIP_LAUNCH_BLOCKING) {
+            tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset in %s.\n", __func__, ToString(stream).c_str());
+            cf.wait();
+            tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed in %s.\n", __func__, ToString(stream).c_str());
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+}
+
 hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char  value, size_t sizeBytes )
 {
     HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes);
@@ -1108,7 +1263,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char  value, size_t sizeByte
     return ihipLogStatus(e);
 }
 
-hipError_t hipMemGetInfo  (size_t *free, size_t *total)
+hipError_t hipMemGetInfo(size_t *free, size_t *total)
 {
     HIP_INIT_API(free, total);
 

src/hip_texture.cpp

@@ -0,0 +1,668 @@
+
+#include <map>
+
+#include <string.h>
+
+#include "hsa/hsa.h"
+#include "hsa/hsa_ext_amd.h"
+
+#include "hip/hip_runtime.h"
+#include "hip_hcc_internal.h"
+#include "trace_helper.h"
+
+#include "hip_texture.h"
+
+static std::map<hipTextureObject_t, hipTexture*> textureHash;
+
+void saveTextureInfo(const hipTexture* pTexture,
+                     const hipResourceDesc* pResDesc,
+                     const hipTextureDesc* pTexDesc,
+                     const hipResourceViewDesc* pResViewDesc)
+{
+    if (pResDesc != nullptr) {
+        memcpy((void*)&(pTexture->resDesc), (void*)pResDesc, sizeof(hipResourceDesc));
+    }
+
+    if (pTexDesc != nullptr) {
+        memcpy((void*)&(pTexture->texDesc), (void*)pTexDesc, sizeof(hipTextureDesc));
+    }
+
+    if (pResViewDesc != nullptr) {
+        memcpy((void*)&(pTexture->resViewDesc), (void*)pResViewDesc, sizeof(hipResourceViewDesc));
+    }
+}
+
+void getChannelOrderAndType(const hipChannelFormatDesc& desc,
+                            enum hipTextureReadMode readMode,
+                            hsa_ext_image_channel_order_t& channelOrder,
+                            hsa_ext_image_channel_type_t& channelType)
+{
+    if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w != 0) {
+        channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
+    } else if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w == 0) {
+        channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
+    } else if (desc.x != 0 && desc.y != 0 && desc.z == 0 && desc.w == 0) {
+        channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
+    } else if (desc.x != 0 && desc.y == 0 && desc.z == 0 && desc.w == 0) {
+        channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
+    } else {
+    }
+
+    switch (desc.f) {
+    case hipChannelFormatKindUnsigned:
+        switch(desc.x) {
+        case 32:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
+            break;
+        case 16:
+            channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 :
+                                                                   HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
+            break;
+        case 8:
+            channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 :
+                                                                   HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
+            break;
+        default:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
+        }
+        break;
+    case hipChannelFormatKindSigned:
+        switch(desc.x) {
+        case 32:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
+            break;
+        case 16:
+            channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 :
+                                                                   HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
+            break;
+        case 8:
+            channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 :
+                                                                   HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
+            break;
+        default:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
+        }
+        break;
+    case hipChannelFormatKindFloat:
+        switch(desc.x) {
+        case 32:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
+            break;
+        case 16:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
+            break;
+        case 8:
+            break;
+        default:
+            channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
+        }
+        break;
+    case hipChannelFormatKindNone:
+    default:
+        break;
+    }
+}
+
+void fillSamplerDescriptor(hsa_ext_sampler_descriptor_t& samplerDescriptor,
+                           enum hipTextureAddressMode addressMode,
+                           enum hipTextureFilterMode filterMode,
+                           int normalizedCoords)
+{
+    if (normalizedCoords) {
+        samplerDescriptor.coordinate_mode = HSA_EXT_SAMPLER_COORDINATE_MODE_NORMALIZED;
+    } else {
+        samplerDescriptor.coordinate_mode = HSA_EXT_SAMPLER_COORDINATE_MODE_UNNORMALIZED;
+    }
+
+    switch (filterMode) {
+    case hipFilterModePoint:
+        samplerDescriptor.filter_mode = HSA_EXT_SAMPLER_FILTER_MODE_NEAREST;
+        break;
+    case hipFilterModeLinear:
+        samplerDescriptor.filter_mode = HSA_EXT_SAMPLER_FILTER_MODE_LINEAR;
+        break;
+    }
+
+    switch (addressMode) {
+    case hipAddressModeWrap:
+        samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_REPEAT;
+        break;
+    case hipAddressModeClamp:
+        samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE;
+        break;
+    case hipAddressModeMirror:
+        samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT;
+        break;
+    case hipAddressModeBorder:
+        samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_BORDER;
+        break;
+    }
+}
+
+bool getHipTextureObject(hipTextureObject_t* pTexObject,
+                         hsa_ext_image_t& image,
+                         hsa_ext_sampler_t sampler)
+{
+    unsigned int* texSRD;
+    hipMalloc((void **) &texSRD, HIP_TEXTURE_OBJECT_SIZE_DWORD * 4);
+    hipMemcpy(texSRD, (void *)image.handle, HIP_IMAGE_OBJECT_SIZE_DWORD * 4, hipMemcpyDeviceToDevice);
+    hipMemcpy(texSRD + HIP_SAMPLER_OBJECT_OFFSET_DWORD, (void *)sampler.handle, HIP_SAMPLER_OBJECT_SIZE_DWORD * 4, hipMemcpyDeviceToDevice);
+    *pTexObject = (hipTextureObject_t) texSRD;
+
+#ifdef DEBUG
+    unsigned int* srd = (unsigned int*) malloc(HIP_TEXTURE_OBJECT_SIZE_DWORD * 4);
+    hipMemcpy(srd, texSRD, HIP_TEXTURE_OBJECT_SIZE_DWORD * 4, hipMemcpyDeviceToHost);
+    printf("New SRD: \n");
+    for (int i = 0; i < HIP_TEXTURE_OBJECT_SIZE_DWORD; i++) {
+        printf("SRD[%d]: %x\n", i, srd[i]);
+    }
+    printf("\n");
+#endif
+    return true;
+}
+
+// Texture Object APIs
+hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
+                                  const hipResourceDesc* pResDesc,
+                                  const hipTextureDesc* pTexDesc,
+                                  const hipResourceViewDesc* pResViewDesc)
+{
+    HIP_INIT_API(pTexObject, pResDesc, pTexDesc, pResViewDesc);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hc::accelerator acc = ctx->getDevice()->_acc;
+        auto device = ctx->getWriteableDevice();
+
+        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+        hipTexture* pTexture = (hipTexture*) malloc(sizeof(hipTexture));
+        if (pTexture != nullptr) {
+            memset(pTexture, 0, sizeof(hipTexture));
+            saveTextureInfo(pTexture, pResDesc, pTexDesc, pResViewDesc);
+        }
+
+        hsa_ext_image_descriptor_t imageDescriptor;
+        hsa_ext_image_channel_order_t channelOrder;
+        hsa_ext_image_channel_type_t channelType;
+        void* devPtr = nullptr;
+
+        switch (pResDesc->resType) {
+        case hipResourceTypeArray:
+            devPtr = pResDesc->res.array.array->data;
+            imageDescriptor.width = pResDesc->res.array.array->width;
+            imageDescriptor.height = pResDesc->res.array.array->height;
+            switch (pResDesc->res.array.array->type) {
+            case hipArrayLayered:
+                imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA;
+                imageDescriptor.depth = 0;
+                imageDescriptor.array_size = pResDesc->res.array.array->depth;
+                break;
+            case hipArrayCubemap:
+                imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D;
+                imageDescriptor.depth = pResDesc->res.array.array->depth;
+                imageDescriptor.array_size = 0;
+                break;
+            case hipArraySurfaceLoadStore:
+            case hipArrayTextureGather:
+            case hipArrayDefault:
+            default:
+                imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
+                imageDescriptor.depth = 0;
+                imageDescriptor.array_size = 0;
+                break;
+            }
+            getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode, channelOrder, channelType);
+            break;
+        case hipResourceTypeMipmappedArray:
+            devPtr = pResDesc->res.mipmap.mipmap->data;
+            imageDescriptor.width = pResDesc->res.mipmap.mipmap->width;
+            imageDescriptor.height = pResDesc->res.mipmap.mipmap->height;
+            imageDescriptor.depth = pResDesc->res.mipmap.mipmap->depth;
+            imageDescriptor.array_size = 0;
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
+            getChannelOrderAndType(pResDesc->res.mipmap.mipmap->desc, pTexDesc->readMode, channelOrder, channelType);
+            break;
+        case hipResourceTypeLinear:
+            devPtr = pResDesc->res.linear.devPtr;
+            imageDescriptor.width = pResDesc->res.linear.sizeInBytes;
+            imageDescriptor.height = 1;
+            imageDescriptor.depth = 0;
+            imageDescriptor.array_size = 0;
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_1D; // ? HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR
+            getChannelOrderAndType(pResDesc->res.linear.desc, pTexDesc->readMode, channelOrder, channelType);
+            break;
+        case hipResourceTypePitch2D:
+            devPtr = pResDesc->res.pitch2D.devPtr;
+            imageDescriptor.width = pResDesc->res.pitch2D.width;
+            imageDescriptor.height = pResDesc->res.pitch2D.height;
+            imageDescriptor.depth = 0;
+            imageDescriptor.array_size = 0;
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
+            getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, channelOrder, channelType);
+            break;
+        default:
+            break;
+        }
+
+        imageDescriptor.format.channel_order = channelOrder;
+        imageDescriptor.format.channel_type = channelType;
+
+        hsa_ext_sampler_descriptor_t samplerDescriptor;
+        fillSamplerDescriptor(samplerDescriptor, pTexDesc->addressMode[0], pTexDesc->filterMode, pTexDesc->normalizedCoords);
+
+        hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+        if (HSA_STATUS_SUCCESS != hsa_ext_image_create(*agent, &imageDescriptor, devPtr, permission, &(pTexture->image)) ||
+            HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
+            return ihipLogStatus(hipErrorRuntimeOther);
+        }
+
+        getHipTextureObject(pTexObject, pTexture->image, pTexture->sampler);
+
+        textureHash[*pTexObject] = pTexture;
+    }
+
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject)
+{
+    HIP_INIT_API(textureObject);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hc::accelerator acc = ctx->getDevice()->_acc;
+        auto device = ctx->getWriteableDevice();
+
+        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+        hipTexture* pTexture = textureHash[textureObject];
+        if (pTexture != nullptr) {
+            hsa_ext_image_destroy(*agent, pTexture->image);
+            hsa_ext_sampler_destroy(*agent, pTexture->sampler);
+            free(pTexture);
+            textureHash.erase(textureObject);
+        }
+    }
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTextureObject_t textureObject)
+{
+    HIP_INIT_API(pResDesc, textureObject);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hipTexture* pTexture = textureHash[textureObject];
+        if (pTexture != nullptr && pResDesc != nullptr) {
+            memcpy((void*)pResDesc, (void*)&(pTexture->resDesc), sizeof(hipResourceDesc));
+        }
+    }
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, hipTextureObject_t textureObject)
+{
+    HIP_INIT_API(pResViewDesc, textureObject);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hipTexture* pTexture = textureHash[textureObject];
+        if (pTexture != nullptr && pResViewDesc != nullptr) {
+            memcpy((void*)pResViewDesc, (void*)&(pTexture->resViewDesc), sizeof(hipResourceViewDesc));
+        }
+    }
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject)
+{
+    HIP_INIT_API(pTexDesc, textureObject);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hipTexture* pTexture = textureHash[textureObject];
+        if (pTexture != nullptr && pTexDesc != nullptr) {
+            memcpy((void*)pTexDesc, (void*)&(pTexture->texDesc), sizeof(hipTextureDesc));
+        }
+    }
+    return ihipLogStatus(hip_status);
+}
+
+// Texture Reference APIs
+hipError_t ihipBindTextureImpl(int dim,
+                               enum hipTextureReadMode readMode,
+                               size_t *offset,
+                               const void *devPtr,
+                               const struct hipChannelFormatDesc& desc,
+                               size_t size,
+                               enum hipTextureAddressMode addressMode,
+                               enum hipTextureFilterMode filterMode,
+                               int normalizedCoords,
+                               hipTextureObject_t& textureObject)
+{
+    HIP_INIT_API();
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hc::accelerator acc = ctx->getDevice()->_acc;
+        auto device = ctx->getWriteableDevice();
+
+        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+        hipTexture* pTexture = (hipTexture*) malloc(sizeof(hipTexture));
+        if (pTexture != nullptr) {
+            memset(pTexture, 0, sizeof(hipTexture));
+        }
+
+        hsa_ext_image_descriptor_t imageDescriptor;
+
+        assert(dim == hipTextureType1D);
+
+        imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_1D;
+        imageDescriptor.width = size;
+        imageDescriptor.height = 1;
+        imageDescriptor.depth = 1;
+        imageDescriptor.array_size = 0;
+
+        hsa_ext_image_channel_order_t channelOrder;
+        hsa_ext_image_channel_type_t channelType;
+        getChannelOrderAndType(desc, readMode, channelOrder, channelType);
+        imageDescriptor.format.channel_order = channelOrder;
+        imageDescriptor.format.channel_type = channelType;
+
+        hsa_ext_sampler_descriptor_t samplerDescriptor;
+        fillSamplerDescriptor(samplerDescriptor, addressMode, filterMode, normalizedCoords);
+
+        hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+
+        if (HSA_STATUS_SUCCESS != hsa_ext_image_create(*agent, &imageDescriptor, devPtr, permission, &(pTexture->image)) ||
+            HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
+            return ihipLogStatus(hipErrorRuntimeOther);
+        }
+        getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler);
+        textureHash[textureObject] = pTexture;
+    }
+
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipBindTexture(size_t* offset,
+                          textureReference* tex,
+                          const void* devPtr,
+                          const hipChannelFormatDesc* desc,
+                          size_t size)
+{
+    // TODO: hipReadModeElementType is default.
+    return ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType,
+                               offset, devPtr, *desc, size,
+                               tex->addressMode[0], tex->filterMode, tex->normalized,
+                               tex->textureObject);
+}
+
+hipError_t ihipBindTexture2DImpl(int dim,
+                                 enum hipTextureReadMode readMode,
+                                 size_t *offset,
+                                 const void *devPtr,
+                                 const struct hipChannelFormatDesc& desc,
+                                 size_t width,
+                                 size_t height,
+                                 enum hipTextureAddressMode addressMode,
+                                 enum hipTextureFilterMode filterMode,
+                                 int normalizedCoords,
+                                 hipTextureObject_t& textureObject)
+{
+    HIP_INIT_API();
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hc::accelerator acc = ctx->getDevice()->_acc;
+        auto device = ctx->getWriteableDevice();
+
+        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+        hipTexture* pTexture = (hipTexture*) malloc(sizeof(hipTexture));
+        if (pTexture != nullptr) {
+            memset(pTexture, 0, sizeof(hipTexture));
+        }
+
+        hsa_ext_image_descriptor_t imageDescriptor;
+
+        assert(dim == hipTextureType2D);
+
+        imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
+        imageDescriptor.width = width;
+        imageDescriptor.height = height;
+        imageDescriptor.depth = 1;
+        imageDescriptor.array_size = 0;
+
+        hsa_ext_image_channel_order_t channelOrder;
+        hsa_ext_image_channel_type_t channelType;
+        getChannelOrderAndType(desc, readMode, channelOrder, channelType);
+        imageDescriptor.format.channel_order = channelOrder;
+        imageDescriptor.format.channel_type = channelType;
+
+        hsa_ext_sampler_descriptor_t samplerDescriptor;
+        fillSamplerDescriptor(samplerDescriptor, addressMode, filterMode, normalizedCoords);
+
+        hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+
+        if (HSA_STATUS_SUCCESS != hsa_ext_image_create(*agent, &imageDescriptor, devPtr, permission, &(pTexture->image)) ||
+            HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
+            return ihipLogStatus(hipErrorRuntimeOther);
+        }
+        getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler);
+        textureHash[textureObject] = pTexture;
+    }
+
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipBindTexture2D(size_t* offset,
+                            textureReference* tex,
+                            const void* devPtr,
+                            const hipChannelFormatDesc* desc,
+                            size_t width,
+                            size_t height,
+                            size_t pitch)
+{
+    // TODO: hipReadModeElementType is default.
+	return ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType,
+                                 offset, devPtr, *desc, width, height,
+                                 tex->addressMode[0], tex->filterMode, tex->normalized,
+                                 tex->textureObject);
+}
+
+hipError_t ihipBindTextureToArrayImpl(int dim,
+                                      enum hipTextureReadMode readMode,
+                                      hipArray_const_t array,
+                                      const struct hipChannelFormatDesc& desc,
+                                      enum hipTextureAddressMode addressMode,
+                                      enum hipTextureFilterMode filterMode,
+                                      int normalizedCoords,
+                                      hipTextureObject_t& textureObject)
+{
+	HIP_INIT_API();
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hc::accelerator acc = ctx->getDevice()->_acc;
+        auto device = ctx->getWriteableDevice();
+
+        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+        hipTexture* pTexture = (hipTexture*) malloc(sizeof(hipTexture));
+        if (pTexture != nullptr) {
+            memset(pTexture, 0, sizeof(hipTexture));
+        }
+
+        hsa_ext_image_descriptor_t imageDescriptor;
+
+        imageDescriptor.width = array->width;
+        imageDescriptor.height = array->height;
+        imageDescriptor.depth = array->depth;
+        imageDescriptor.array_size = 0;
+
+        switch (dim) {
+        case hipTextureType1D:
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_1D;
+            imageDescriptor.height = 1;
+            imageDescriptor.depth = 1;
+            break;
+        case hipTextureType2D:
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
+            imageDescriptor.depth = 1;
+            break;
+        case hipTextureType3D:
+        case hipTextureTypeCubemap:
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D;
+            break;
+        case hipTextureType1DLayered:
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_1DA;
+            imageDescriptor.height = 1;
+            imageDescriptor.array_size = array->height;
+            break;
+        case hipTextureType2DLayered:
+            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA;
+            imageDescriptor.depth = 1;
+            imageDescriptor.array_size = array->depth;
+            break;
+        case hipTextureTypeCubemapLayered:
+        default:
+            break;
+        }
+
+        hsa_ext_image_channel_order_t channelOrder;
+        hsa_ext_image_channel_type_t channelType;
+        getChannelOrderAndType(desc, readMode, channelOrder, channelType);
+        imageDescriptor.format.channel_order = channelOrder;
+        imageDescriptor.format.channel_type = channelType;
+
+        hsa_ext_sampler_descriptor_t samplerDescriptor;
+        fillSamplerDescriptor(samplerDescriptor, addressMode, filterMode, normalizedCoords);
+
+        hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+
+        if (HSA_STATUS_SUCCESS != hsa_ext_image_create(*agent, &imageDescriptor, array->data, permission, &(pTexture->image)) ||
+            HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
+            return ihipLogStatus(hipErrorRuntimeOther);
+        }
+        getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler);
+        textureHash[textureObject] = pTexture;
+    }
+
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipBindTextureToArray(textureReference* tex,
+                                 hipArray_const_t array,
+                                 const hipChannelFormatDesc* desc)
+{
+    // TODO: hipReadModeElementType is default.
+    return ihipBindTextureToArrayImpl(hipTextureType2D, hipReadModeElementType,
+                                      array, *desc,
+                                      tex->addressMode[0], tex->filterMode, tex->normalized,
+                                      tex->textureObject);
+}
+
+hipError_t hipBindTextureToMipmappedArray(textureReference* tex,
+                                          hipMipmappedArray_const_t mipmappedArray,
+                                          const hipChannelFormatDesc* desc)
+{
+	return hipSuccess;
+}
+
+hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject)
+{
+    HIP_INIT_API();
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        hc::accelerator acc = ctx->getDevice()->_acc;
+        auto device = ctx->getWriteableDevice();
+
+        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+		hipTexture* pTexture = textureHash[textureObject];
+        if (pTexture != nullptr) {
+            hsa_ext_image_destroy(*agent, pTexture->image);
+            hsa_ext_sampler_destroy(*agent, pTexture->sampler);
+            free(pTexture);
+            textureHash.erase(textureObject);
+        }
+    }
+
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipUnbindTexture(const textureReference* tex)
+{
+    return ihipUnbindTextureImpl(tex->textureObject);
+}
+
+hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array)
+{
+    HIP_INIT_API(desc, array);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+        *desc = array->desc;
+    }
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* tex)
+{
+    HIP_INIT_API(offset, tex);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+    }
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipGetTextureReference(const textureReference** tex, const void* symbol)
+{
+    HIP_INIT_API(tex, symbol);
+    HIP_SET_DEVICE();
+
+    hipError_t  hip_status = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx) {
+    }
+    return ihipLogStatus(hip_status);
+}

src/hip_texture.h

@@ -0,0 +1,36 @@
+/*
+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.
+*/
+
+#ifndef HIP_INCLUDE_HCC_DETAIL_HIP_TEXTURE_H
+#define HIP_INCLUDE_HCC_DETAIL_HIP_TEXTURE_H
+
+#include <hip/hcc_detail/texture_types.h>
+
+struct hipTexture {
+    hipResourceDesc resDesc;
+    hipTextureDesc texDesc;
+    hipResourceViewDesc resViewDesc;
+    hsa_ext_image_t image;
+    hsa_ext_sampler_t sampler;
+};
+
+#endif