diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp index 89b06aca95..e6f3f2f622 100644 --- a/api/hip/hip_memory.cpp +++ b/api/hip/hip_memory.cpp @@ -500,13 +500,13 @@ hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc const cl_image_format image_format = { channelOrder, channelType }; // Dummy flags check - switch (flags) { + switch (flags) { case hipArrayLayered: + case hipArrayCubemap: case hipArraySurfaceLoadStore: case hipArrayTextureGather: assert(0 && "Unspported"); break; - case hipArrayCubemap: case hipArrayDefault: default: break; diff --git a/api/hip/hip_texture.cpp b/api/hip/hip_texture.cpp index d1c228707b..a4712de1cb 100644 --- a/api/hip/hip_texture.cpp +++ b/api/hip/hip_texture.cpp @@ -299,6 +299,8 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou amd::Memory* memory = nullptr; size_t offset = 0; + cl_mem_object_type clType; + switch (pResDesc->resType) { case hipResourceTypeArray: { @@ -316,9 +318,23 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou case hipArrayTextureGather: case hipArrayDefault: default: - image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(), - CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat, - pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0); + switch(pResDesc->res.array.array->textureType) { + case hipTextureType3D: + clType = CL_MEM_OBJECT_IMAGE3D; + image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(), + clType, memory->getMemFlags(), imageFormat, + pResDesc->res.array.array->width, pResDesc->res.array.array->height, + pResDesc->res.array.array->depth, 0, 0); + break; + case hipTextureType2D: + clType = CL_MEM_OBJECT_IMAGE2D; + image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(), + clType, memory->getMemFlags(), imageFormat, + pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0); + break; + default: + break; + } break; } }