diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index f2033be250..75ca1c1367 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -253,7 +253,8 @@ inline static CUresourcetype hipResourcetype_enumToCUresourcetype( #define hipTexRef CUtexref #define hiparray CUarray -typedef CUmipmappedArray hipMipmappedArray_t; +typedef CUmipmappedArray hipmipmappedArray; +typedef cudaMipmappedArray_t hipMipmappedArray_t; #define HIP_TRSA_OVERRIDE_FORMAT CU_TRSA_OVERRIDE_FORMAT #define HIP_TRSF_READ_AS_INTEGER CU_TRSF_READ_AS_INTEGER @@ -1579,22 +1580,38 @@ inline static hipError_t hipFreeArray(hipArray* array) { return hipCUDAErrorTohipError(cudaFreeArray(array)); } -inline static hipError_t hipMipmappedArrayCreate(hipMipmappedArray_t* pHandle, +inline static hipError_t hipMipmappedArrayCreate(hipmipmappedArray* pHandle, HIP_ARRAY3D_DESCRIPTOR* pMipmappedArrayDesc, unsigned int numMipmapLevels) { return hipCUResultTohipError(cuMipmappedArrayCreate(pHandle, pMipmappedArrayDesc, numMipmapLevels)); } -inline static hipError_t hipMipmappedArrayDestroy(hipMipmappedArray_t hMipmappedArray) { +inline static hipError_t hipMipmappedArrayDestroy(hipmipmappedArray hMipmappedArray) { return hipCUResultTohipError(cuMipmappedArrayDestroy(hMipmappedArray)); } -inline static hipError_t hipMipmappedArrayGetLevel(hipArray_t* pLevelArray, - hipMipmappedArray_t hMipMappedArray, +inline static hipError_t hipMipmappedArrayGetLevel(hiparray* pLevelArray, + hipmipmappedArray hMipMappedArray, unsigned int level) { return hipCUResultTohipError(cuMipmappedArrayGetLevel((CUarray*)pLevelArray, hMipMappedArray, level)); } +inline static hipError_t hipMallocMipmappedArray(hipMipmappedArray_t* pHandle, + const hipChannelFormatDesc* desc, hipExtent extent, + unsigned int numLevels, unsigned int flags = 0) { + return hipCUDAErrorTohipError(cudaMallocMipmappedArray(pHandle, desc, extent, numLevels, flags)); +} + +inline static hipError_t hipFreeMipmappedArray(hipMipmappedArray_t hMipmappedArray) { + return hipCUDAErrorTohipError(cudaFreeMipmappedArray(hMipmappedArray)); +} + +inline static hipError_t hipGetMipmappedArrayLevel(hipArray_t* pLevelArray, + hipMipmappedArray_t hMipMappedArray, + unsigned int level) { + return hipCUDAErrorTohipError(cudaGetMipmappedArrayLevel(pLevelArray, hMipMappedArray, level)); +} + inline static hipError_t hipHostGetDevicePointer(void** devPtr, void* hostPtr, unsigned int flags) { return hipCUDAErrorTohipError(cudaHostGetDevicePointer(devPtr, hostPtr, flags)); } diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 0288f87bb4..cf48db9062 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -3948,30 +3948,6 @@ hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy) { HIP_RETURN(ihipMemcpyParam3D(&desc, nullptr)); } -hipError_t hipMallocMipmappedArray(hipMipmappedArray_t *mipmappedArray, - const hipChannelFormatDesc* desc, - hipExtent extent, - unsigned int numLevels, - unsigned int flags) { - HIP_INIT_API(hipMallocMipmappedArray, mipmappedArray, desc, extent, numLevels, flags); - CHECK_STREAM_CAPTURE_SUPPORTED(); - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray) { - HIP_INIT_API(hipFreeMipmappedArray, mipmappedArray); - CHECK_STREAM_CAPTURE_SUPPORTED(); - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t hipGetMipmappedArrayLevel(hipArray_t *levelArray, - hipMipmappedArray_const_t mipmappedArray, - unsigned int level) { - HIP_INIT_API(hipGetMipmappedArrayLevel, levelArray, mipmappedArray, level); - - HIP_RETURN(hipErrorNotSupported); -} - hipError_t ihipMipmapArrayCreate(hipMipmappedArray_t* mipmapped_array_pptr, HIP_ARRAY3D_DESCRIPTOR* mipmapped_array_desc_ptr, unsigned int num_mipmap_levels) { @@ -4029,6 +4005,7 @@ hipError_t ihipMipmapArrayCreate(hipMipmappedArray_t* mipmapped_array_pptr, (*mipmapped_array_pptr)->max_mipmap_level = num_mipmap_levels; (*mipmapped_array_pptr)->flags = mipmapped_array_desc_ptr->Flags; (*mipmapped_array_pptr)->format = mipmapped_array_desc_ptr->Format; + (*mipmapped_array_pptr)->num_channels = mipmapped_array_desc_ptr->NumChannels; return hipSuccess; } @@ -4100,6 +4077,9 @@ hipError_t ihipMipmappedArrayGetLevel(hipArray_t* level_array_pptr, (*level_array_pptr)->isDrv = 0; (*level_array_pptr)->textureType = 0; + amd::ScopedLock lock(hip::hipArraySetLock); + hip::hipArraySet.insert(*level_array_pptr); + return hipSuccess; } @@ -4127,3 +4107,40 @@ hipError_t hipMipmappedArrayGetLevel(hipArray_t* level_array_pptr, HIP_RETURN(ihipMipmappedArrayGetLevel(level_array_pptr, mipmapped_array_ptr, mip_level)); } +hipError_t hipMallocMipmappedArray(hipMipmappedArray_t *mipmappedArray, + const hipChannelFormatDesc* desc, + hipExtent extent, + unsigned int numLevels, + unsigned int flags) { + HIP_INIT_API(hipMallocMipmappedArray, mipmappedArray, desc, extent, numLevels, flags); + if (mipmappedArray == nullptr || desc == nullptr) { + return hipErrorInvalidValue; + } + CHECK_STREAM_CAPTURE_SUPPORTED(); + HIP_ARRAY3D_DESCRIPTOR allocateArray = {extent.width, + extent.height, + extent.depth, + hip::getArrayFormat(*desc), + hip::getNumChannels(*desc), + flags}; + if(!hip::CheckArrayFormat(*desc)) { + return hipErrorInvalidValue; + } + HIP_RETURN(ihipMipmapArrayCreate(mipmappedArray, &allocateArray, numLevels)); +} + +hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray) { + HIP_INIT_API(hipFreeMipmappedArray, mipmappedArray); + CHECK_STREAM_CAPTURE_SUPPORTED(); + HIP_RETURN(ihipMipmappedArrayDestroy(mipmappedArray)); +} + +hipError_t hipGetMipmappedArrayLevel(hipArray_t *levelArray, + hipMipmappedArray_const_t mipmappedArray, + unsigned int level) { + HIP_INIT_API(hipGetMipmappedArrayLevel, levelArray, mipmappedArray, level); + CHECK_STREAM_CAPTURE_SUPPORTED(); + HIP_RETURN(ihipMipmappedArrayGetLevel(levelArray, + const_cast(mipmappedArray), + level)); +} diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index 8c44373901..0f7d45dac7 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/projects/clr/hipamd/src/hip_texture.cpp @@ -139,25 +139,22 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, return hipErrorInvalidValue; } - // Mipmaps are currently not supported. - if (pResDesc->resType == hipResourceTypeMipmappedArray) { - return hipErrorNotSupported; - } // We don't program the max_ansio_ratio field in the the HW sampler SRD. if (pTexDesc->maxAnisotropy != 0) { return hipErrorNotSupported; } // We don't program the lod_bias field in the HW sampler SRD. if (pTexDesc->mipmapLevelBias != 0) { + LogError("mipmapLevelBias not supported!"); return hipErrorNotSupported; } // We don't program the min_lod field in the HW sampler SRD. if (pTexDesc->minMipmapLevelClamp != 0) { - return hipErrorNotSupported; + LogInfo("minMipmapLevelClamp ignored!"); } // We don't program the max_lod field in the HW sampler SRD. if (pTexDesc->maxMipmapLevelClamp != 0) { - return hipErrorNotSupported; + LogInfo("maxMipmapLevelClamp ignored!"); } // TODO ROCclr assumes all dimensions have the same addressing mode. @@ -242,12 +239,50 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, if (image == nullptr) { return hipErrorInvalidValue; } + } else if (image->parent()) { + image->retain(); // Because it will be released as a view in ihipDestroyTextureObject() } break; } - case hipResourceTypeMipmappedArray: - return hipErrorInvalidValue; + case hipResourceTypeMipmappedArray: { + cl_mem memObj = reinterpret_cast(pResDesc->res.array.array->data); + if (!is_valid(memObj)) { + return hipErrorInvalidValue; + } + image = as_amd(memObj)->asImage(); + hipTextureReadMode readMode = pTexDesc->readMode; + // 32-bit integer format will not be promoted, regardless of whether or not + // this hipTextureDesc::readMode is set hipReadModeNormalizedFloat is specified. + if ((pResDesc->res.array.array->Format == HIP_AD_FORMAT_SIGNED_INT32) || + (pResDesc->res.array.array->Format == HIP_AD_FORMAT_UNSIGNED_INT32)) { + readMode = hipReadModeElementType; + } + + // We need to create an image view if the user requested to use normalized pixel values, + // due to already having the image created with a different format. + if ((pResViewDesc != nullptr) || (readMode == hipReadModeNormalizedFloat) || + (pTexDesc->sRGB == 1)) { + // TODO ROCclr currently right now can only change the format of the image. + const cl_channel_order channelOrder = (pResViewDesc != nullptr) + ? hip::getCLChannelOrder(hip::getNumChannels(pResViewDesc->format), pTexDesc->sRGB) + : hip::getCLChannelOrder(pResDesc->res.mipmap.mipmap->num_channels, pTexDesc->sRGB); + const cl_channel_type channelType = (pResViewDesc != nullptr) + ? hip::getCLChannelType(hip::getArrayFormat(pResViewDesc->format), readMode) + : hip::getCLChannelType(pResDesc->res.mipmap.mipmap->format, readMode); + const amd::Image::Format imageFormat(cl_image_format{channelOrder, channelType}); + if (!imageFormat.isValid()) { + return hipErrorInvalidValue; + } + + image = image->createView(*hip::getCurrentDevice()->asContext(), imageFormat, nullptr, 0, 0, + true); + if (image == nullptr) { + return hipErrorInvalidValue; + } + } + break; + } case hipResourceTypeLinear: { const cl_channel_order channelOrder = hip::getCLChannelOrder(hip::getNumChannels(pResDesc->res.linear.desc), pTexDesc->sRGB); const cl_channel_type channelType = hip::getCLChannelType(hip::getArrayFormat(pResDesc->res.linear.desc), pTexDesc->readMode); diff --git a/projects/clr/rocclr/platform/memory.cpp b/projects/clr/rocclr/platform/memory.cpp index b71f3af15d..65e7f4881c 100644 --- a/projects/clr/rocclr/platform/memory.cpp +++ b/projects/clr/rocclr/platform/memory.cpp @@ -567,14 +567,15 @@ void Pipe::initDeviceMemory() { #define GETMIPDIM(dim, mip) (((dim >> mip) > 0) ? (dim >> mip) : 1) -Image::Image(const Format& format, Image& parent, uint baseMipLevel, cl_mem_flags flags) +Image::Image(const Format& format, Image& parent, uint baseMipLevel, cl_mem_flags flags, + bool isMipmapView) : Memory(parent, flags, 0, parent.getWidth() * parent.getHeight() * parent.getDepth() * format.getElementSize()), impl_(format, Coord3D(parent.getWidth() * parent.getImageFormat().getElementSize() / format.getElementSize(), parent.getHeight(), parent.getDepth()), parent.getRowPitch(), parent.getSlicePitch(), parent.getBytePitch()), - mipLevels_(1), + mipLevels_(isMipmapView ? parent.getMipLevels() : 1), baseMipLevel_(baseMipLevel) { if (baseMipLevel > 0) { impl_.region_.c[0] = GETMIPDIM(parent.getWidth(), baseMipLevel) * @@ -1191,10 +1192,10 @@ bool Image::Format::isSupported(const Context& context, cl_mem_object_type image // ================================================================================================ Image* Image::createView(const Context& context, const Format& format, device::VirtualDevice* vDev, - uint baseMipLevel, cl_mem_flags flags) { + uint baseMipLevel, cl_mem_flags flags, bool createMipmapView) { // Find the image dimensions and create a corresponding object - Image* view = new (context) Image(format, *this, baseMipLevel, flags); + Image* view = new (context) Image(format, *this, baseMipLevel, flags, createMipmapView); if (view != nullptr) { // Set GPU virtual device for this view diff --git a/projects/clr/rocclr/platform/memory.hpp b/projects/clr/rocclr/platform/memory.hpp index 1ea2ff0e4f..42a6dc949d 100644 --- a/projects/clr/rocclr/platform/memory.hpp +++ b/projects/clr/rocclr/platform/memory.hpp @@ -540,7 +540,8 @@ class Image : public Memory { uint baseMipLevel_; //!< The base mip level for a view protected: - Image(const Format& format, Image& parent, uint baseMipLevel = 0, cl_mem_flags flags = 0); + Image(const Format& format, Image& parent, uint baseMipLevel = 0, cl_mem_flags flags = 0, + bool isMipmapView = false); ///! Initializes the device memory array which is nested // after'Image' object in memory layout. @@ -596,7 +597,8 @@ class Image : public Memory { const Format& format, //!< The new format for a view device::VirtualDevice* vDev, //!< Virtual device object uint baseMipLevel = 0, //!< Base mip level for a view - cl_mem_flags flags = 0 //!< Memory allocation flags + cl_mem_flags flags = 0, //!< Memory allocation flags + bool createMipmapView = false //!< To create mipmap view based on this image ); //! Returns the impl for this image.