AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

SWDEV-368553 – Add missing mipmap Apis

Bladeren bron 
Add missing mipmap Apis’ implementation.
Fix some bugs of mimpmap apis.
Use hipmipmappedArray to differentiate cuda
and driver apis on Nvidia.
Change-Id: I6079d9f3b2ddf4e42b9a6f7f3902322cfca02cfd


[ROCm/clr commit: f03c11491b]
This commit is contained in:
taosang2
2023-06-01 18:31:48 -04:00
gecommit door Tao Sang
bovenliggende c14a017d77
commit 9aa3a87b66
5 gewijzigde bestanden met toevoegingen van 115 en 43 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+22 -5
Bestand weergeven
@@ -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));
}
projects/clr/hipamd/src/hip_memory.cpp
+41 -24
Bestand weergeven
@@ -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<hipMipmappedArray_t>(mipmappedArray),
level));
}
projects/clr/hipamd/src/hip_texture.cpp
+43 -8
Bestand weergeven
@@ -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<cl_mem>(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);
projects/clr/rocclr/platform/memory.cpp
+5 -4
Bestand weergeven
@@ -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
projects/clr/rocclr/platform/memory.hpp
+4 -2
Bestand weergeven
@@ -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.