From 06e4e7ee8f5f94e650c7b3e5a7e2456b29327a46 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 22 Mar 2021 15:52:46 -0400 Subject: [PATCH] SWDEV-245533 - Implementing Mipmap array APIs Change-Id: I222c8d0fcbd50cdefd702462e689463bb36366d2 [ROCm/hip commit: 9c07201c6b7bc7d823fa72133685633fd799aaeb] --- .../hip/include/hip/amd_detail/driver_types.h | 20 +- projects/hip/rocclr/hip_memory.cpp | 240 ++++++++++++++---- .../src/texture/hipTextureMipmapObj2D.cpp | 161 ++++++++++++ 3 files changed, 367 insertions(+), 54 deletions(-) create mode 100644 projects/hip/tests/src/texture/hipTextureMipmapObj2D.cpp diff --git a/projects/hip/include/hip/amd_detail/driver_types.h b/projects/hip/include/hip/amd_detail/driver_types.h index 26a48cabc5..4de4e3ee7e 100644 --- a/projects/hip/include/hip/amd_detail/driver_types.h +++ b/projects/hip/include/hip/amd_detail/driver_types.h @@ -116,14 +116,18 @@ typedef struct hipArray* hipArray_t; typedef hipArray_t hiparray; typedef const struct hipArray* hipArray_const_t; -// TODO: It needs to be modified since it was just copied from hipArray. -struct hipMipmappedArray { - void* data; // FIXME: generalize this - struct hipChannelFormatDesc desc; - unsigned int width; - unsigned int height; - unsigned int depth; -}; +typedef struct hipMipmappedArray { + void* data; + struct hipChannelFormatDesc desc; + unsigned int type; + unsigned int width; + unsigned int height; + unsigned int depth; + unsigned int min_mipmap_level; + unsigned int max_mipmap_level; + unsigned int flags; + enum hipArray_Format format; +} hipMipmappedArray; typedef struct hipMipmappedArray* hipMipmappedArray_t; diff --git a/projects/hip/rocclr/hip_memory.cpp b/projects/hip/rocclr/hip_memory.cpp index ae82572aad..18705e2932 100755 --- a/projects/hip/rocclr/hip_memory.cpp +++ b/projects/hip/rocclr/hip_memory.cpp @@ -619,6 +619,13 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder, return nullptr; } + bool mipMapSupport = true; + for (auto& dev : devices) { + if (!dev->settings().checkExtension(ClKhrMipMapImage)) { + mipMapSupport = false; + } + } + if (!amd::Image::validateDimensions(devices, imageType, imageWidth, @@ -629,6 +636,23 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder, return nullptr; } + if (numMipLevels > 0) { + if (mipMapSupport == true) { + size_t max_dim = std::max(std::max(imageWidth, imageHeight), imageDepth); + size_t mip_levels = 0; + for (mip_levels = 0; max_dim > 0; max_dim >>=1, mip_levels++); + // empty for loop + + if (mip_levels < numMipLevels) { + LogPrintfError("Invalid Mip Levels: %d", numMipLevels); + return nullptr; + } + } else { + LogPrintfError("Mipmap not supported on one of the devices, Mip Level: %d", numMipLevels); + return nullptr; + } + } + // TODO validate the image descriptor. amd::Image* image = nullptr; @@ -2529,52 +2553,6 @@ hipError_t hipMemcpyHtoAAsync(hipArray* dstArray, HIP_RETURN_DURATION(ihipMemcpyHtoA(srcHost, dstArray, {0, 0, 0}, {dstOffset, 0, 0}, {ByteCount, 1, 1}, 0, 0, stream, true)); } -hipError_t hipMipmappedArrayCreate(hipMipmappedArray_t* pHandle, - HIP_ARRAY3D_DESCRIPTOR* pMipmappedArrayDesc, - unsigned int numMipmapLevels) { - HIP_INIT_API(hipMipmappedArrayCreate, pHandle, pMipmappedArrayDesc, numMipmapLevels); - - HIP_RETURN(hipErrorNotSupported); -} - -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); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t hipMipmappedArrayDestroy(hipMipmappedArray_t hMipmappedArray) { - HIP_INIT_API(hipMipmappedArrayDestroy, hMipmappedArray); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray) { - HIP_INIT_API(hipFreeMipmappedArray, mipmappedArray); - - HIP_RETURN(hipErrorNotSupported); -} - -hipError_t hipMipmappedArrayGetLevel(hipArray_t* pLevelArray, - hipMipmappedArray_t hMipMappedArray, - unsigned int level) { - HIP_INIT_API(hipMipmappedArrayGetLevel, pLevelArray, hMipMappedArray, level); - - 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 hipMallocHost(void** ptr, size_t size) { HIP_INIT_API(hipMallocHost, ptr, size); @@ -2596,3 +2574,173 @@ 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); + + HIP_RETURN(hipErrorNotSupported); +} + +hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray) { + HIP_INIT_API(hipFreeMipmappedArray, mipmappedArray); + + 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) { + + const cl_channel_order channel_order = hip::getCLChannelOrder( + mipmapped_array_desc_ptr->NumChannels, 0); + const cl_channel_type channel_type = hip::getCLChannelType(mipmapped_array_desc_ptr->Format, + hipReadModeElementType); + const cl_mem_object_type image_type = hip::getCLMemObjectType(mipmapped_array_desc_ptr->Width, + mipmapped_array_desc_ptr->Height, + mipmapped_array_desc_ptr->Depth, + mipmapped_array_desc_ptr->Flags); + + // Create a new amd::Image with mipmap + amd::Image* image = ihipImageCreate(channel_order, + channel_type, + image_type, + mipmapped_array_desc_ptr->Width, + mipmapped_array_desc_ptr->Height, + mipmapped_array_desc_ptr->Depth, + mipmapped_array_desc_ptr->Depth, + 0 /* row pitch */, + 0 /* slice pitch */, + num_mipmap_levels, + nullptr /* buffer */); + + if (image == nullptr) { + return hipErrorInvalidValue; + } + + cl_mem cl_mem_obj = as_cl(image); + *mipmapped_array_pptr = new hipMipmappedArray(); + (*mipmapped_array_pptr)->data = reinterpret_cast(cl_mem_obj); + + (*mipmapped_array_pptr)->desc = hip::getChannelFormatDesc( + mipmapped_array_desc_ptr->NumChannels, + mipmapped_array_desc_ptr->Format); + (*mipmapped_array_pptr)->type = image_type; + (*mipmapped_array_pptr)->width = mipmapped_array_desc_ptr->Width; + (*mipmapped_array_pptr)->height = mipmapped_array_desc_ptr->Height; + (*mipmapped_array_pptr)->depth = mipmapped_array_desc_ptr->Depth; + (*mipmapped_array_pptr)->min_mipmap_level = 0; + (*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; + + return hipSuccess; +} + +hipError_t ihipMipmappedArrayDestroy(hipMipmappedArray_t mipmapped_array_ptr) { + + if (mipmapped_array_ptr == nullptr) { + return hipErrorInvalidValue; + } + + cl_mem mem_obj = reinterpret_cast(mipmapped_array_ptr->data); + if (is_valid(mem_obj) == false) { + return hipErrorInvalidValue; + } + + for (auto& dev : g_devices) { + dev->NullStream()->finish(); + } + + as_amd(mem_obj)->release(); + + delete mipmapped_array_ptr; + + return hipSuccess; +} + +hipError_t ihipMipmappedArrayGetLevel(hipArray_t* level_array_pptr, + hipMipmappedArray_t mipmapped_array_ptr, + unsigned int mip_level) { + + if (level_array_pptr == nullptr || mipmapped_array_ptr == nullptr) { + return hipErrorInvalidValue; + } + + // Convert the raw data to amd::Image + cl_mem cl_mem_obj = reinterpret_cast(mipmapped_array_ptr->data); + if (is_valid(cl_mem_obj) == false) { + return hipErrorInvalidValue; + } + + amd::Image* image = as_amd(cl_mem_obj)->asImage(); + if (image == nullptr) { + return hipErrorInvalidValue; + } + + // Create new hip Array parameter and create an image view with new mip level. + (*level_array_pptr) = new hipArray(); + (*level_array_pptr)->data = as_cl(image->createView(image->getContext(), + image->getImageFormat(), + NULL, mip_level, 0)); + + // Copy the new width, height & depth details of the flag to hipArray. + cl_mem cl_mip_mem_obj = reinterpret_cast((*level_array_pptr)->data); + if (is_valid(cl_mem_obj) == false) { + return hipErrorInvalidValue; + } + + // Fill the hip_array info from newly created amd::Image's view + amd::Image* mipmap_image = as_amd(cl_mip_mem_obj)->asImage(); + (*level_array_pptr)->width = mipmap_image->getWidth(); + (*level_array_pptr)->height = mipmap_image->getHeight(); + (*level_array_pptr)->depth = mipmap_image->getDepth(); + + const cl_mem_object_type image_type = hip::getCLMemObjectType((*level_array_pptr)->width, + (*level_array_pptr)->height, + (*level_array_pptr)->depth, + mipmapped_array_ptr->flags); + (*level_array_pptr)->type = image_type; + (*level_array_pptr)->Format = mipmapped_array_ptr->format; + (*level_array_pptr)->desc = mipmapped_array_ptr->desc; + (*level_array_pptr)->NumChannels = hip::getNumChannels((*level_array_pptr)->desc); + (*level_array_pptr)->isDrv = 0; + (*level_array_pptr)->textureType = 0; + + return hipSuccess; +} + +hipError_t hipMipmappedArrayCreate(hipMipmappedArray_t* mipmapped_array_pptr, + HIP_ARRAY3D_DESCRIPTOR* mipmapped_array_desc_ptr, + unsigned int num_mipmap_levels) { + HIP_INIT_API(hipMipmappedArrayCreate, mipmapped_array_pptr, mipmapped_array_desc_ptr, + num_mipmap_levels); + + HIP_RETURN(ihipMipmapArrayCreate(mipmapped_array_pptr, mipmapped_array_desc_ptr, + num_mipmap_levels)); +} + +hipError_t hipMipmappedArrayDestroy(hipMipmappedArray_t mipmapped_array_ptr) { + HIP_INIT_API(hipMipmappedArrayDestroy, mipmapped_array_ptr); + + HIP_RETURN(ihipMipmappedArrayDestroy(mipmapped_array_ptr)); +} + +hipError_t hipMipmappedArrayGetLevel(hipArray_t* level_array_pptr, + hipMipmappedArray_t mipmapped_array_ptr, + unsigned int mip_level) { + HIP_INIT_API(hipMipmappedArrayGetLevel, level_array_pptr, mipmapped_array_ptr, mip_level); + + HIP_RETURN(ihipMipmappedArrayGetLevel(level_array_pptr, mipmapped_array_ptr, mip_level)); +} + diff --git a/projects/hip/tests/src/texture/hipTextureMipmapObj2D.cpp b/projects/hip/tests/src/texture/hipTextureMipmapObj2D.cpp new file mode 100644 index 0000000000..13428ca776 --- /dev/null +++ b/projects/hip/tests/src/texture/hipTextureMipmapObj2D.cpp @@ -0,0 +1,161 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include +#include +#include +#include + +#include +#include "test_common.h" + +// Height Width Vector +std::vector hw_vector = {2048, 1024, 512, 256, 64}; +std::vector mip_vector = {8, 4, 2, 1}; + +__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width, + int height, float level) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[y * width + x] = tex2DLod(textureObject, x, y, level); +} + +bool runMipMapTest(unsigned int width, unsigned int height, unsigned int mipmap_level) { + bool testResult = true; + + printf("Width: %u Height: %u mip: %u \n", width, height, mipmap_level); + + // Create new width & height to be tested + unsigned int orig_width = width; + unsigned int orig_height = height; + width /= pow(2, mipmap_level); + height /= pow(2, mipmap_level); + unsigned int size = width * height * sizeof(float); + + + float* hData = (float*)malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + hData[i * width + j] = i * width + j; + } + } + printf("hData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hData[i]); + if (i % width == 0) { + printf("\n"); + } + } + printf("\n"); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + HIP_ARRAY3D_DESCRIPTOR mipmapped_array_desc; + memset(&mipmapped_array_desc, 0x00, sizeof(HIP_ARRAY3D_DESCRIPTOR)); + mipmapped_array_desc.Width = orig_width; + mipmapped_array_desc.Height = orig_height; + mipmapped_array_desc.Depth = 0; + mipmapped_array_desc.Format = HIP_AD_FORMAT_FLOAT; + mipmapped_array_desc.NumChannels = ((channelDesc.x != 0) + (channelDesc.y != 0) + + (channelDesc.z != 0) + (channelDesc.w != 0)); + mipmapped_array_desc.Flags = 0; + + + hipMipmappedArray* mip_array_ptr; + hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc, 2 * mipmap_level); + + hipArray *hipArray = nullptr; + HIPCHECK(hipMipmappedArrayGetLevel(&hipArray, mip_array_ptr, mipmap_level)); + HIPCHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice)); + + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Specify texture object parameters + hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = hipAddressModeWrap; + texDesc.addressMode[1] = hipAddressModeWrap; + texDesc.filterMode = hipFilterModePoint; + texDesc.readMode = hipReadModeElementType; + texDesc.normalizedCoords = 0; + + // Create texture object + hipTextureObject_t textureObject = 0; + hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + + float* dData = NULL; + hipMalloc((void**)&dData, size); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, + width, height, (2 * mipmap_level)); + + hipDeviceSynchronize(); + + float* hOutputData = (float*)malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + printf("dData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hOutputData[i]); + if (i % width == 0) { + printf("\n"); + } + } + printf("\n"); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i * width + j] != hOutputData[i * width + j]) { + printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], + hOutputData[i * width + j]); + testResult = false; + break; + } + } + } + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); + return testResult; +} + + +bool runTest(int argc, char** argv) { + bool testResult = true; + + for (auto& hw: hw_vector) { + for (auto& mip: mip_vector) { + if ((hw / static_cast(pow (2,(mip * 2)))) > 0) { + testResult |= runMipMapTest(hw, hw, mip); + } + } + } + + printf("\n"); + return testResult; +} + +int main(int argc, char** argv) { + bool testResult = true; + +#ifdef _WIN32 + testResult = runTest(argc, argv); +#else + std::cout<<"Mipmaps are Supported only on windows, skipping the test"<