SWDEV-245533 - Implementing Mipmap array APIs
Change-Id: I222c8d0fcbd50cdefd702462e689463bb36366d2
[ROCm/hip commit: 9c07201c6b]
This commit is contained in:
kjayapra-amd
committed by
Karthik Jayaprakash
Karthik Jayaprakash
parent
2ec091086a
commit
06e4e7ee8f
@@ -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;
|
||||
|
||||
|
||||
@@ -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<amd::Memory>(image);
|
||||
*mipmapped_array_pptr = new hipMipmappedArray();
|
||||
(*mipmapped_array_pptr)->data = reinterpret_cast<void*>(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<cl_mem>(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<cl_mem>(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<amd::Memory>(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<cl_mem>((*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));
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,161 @@
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <vector>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
|
||||
// Height Width Vector
|
||||
std::vector<unsigned int> hw_vector = {2048, 1024, 512, 256, 64};
|
||||
std::vector<unsigned int> 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<float>(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<int>(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"<<std::endl;
|
||||
#endif
|
||||
|
||||
if (testResult) {
|
||||
passed();
|
||||
} else {
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user