第1章 GPU图像数据结构
参考OpenCV中Mat和GpuMat的设计,对当前Image类设计了GPU版本,即GPUImage。
1.1. GPU图像头
设计图像头。
struct GPUImageHeader
{
int32_t nWidth = 0; //宽度
int32_t nHeight = 0; //高度
int16_t nChannel = 0; //通道数
int32_t nRefCount = 0; //引用计数
void* pImage = nullptr;
};
其中指针指向的是GPU上的地址。引用计数用来实现浅拷贝以及显存的自动释放管理。不像OpenCV和Nppi,分通道和深度的做法。GPUImage的nChannel可以理解为通道*深度。
1.2. GPU图像数据结构接口
GPUImage用来管理CudaMalloc申请的显存,在GBox设备上,使用CudaMallocManaged来申请的GPU和CPU共用的地址。绑定功能暂不实现。
class BASE_EXPORT GPUImage
{
public:
GPUImage();
GPUImage(int nW, int nH, int nD); //create
~GPUImage();
GPUImage(const GPUImage& m);
GPUImage& operator=(const GPUImage& m);
void Create(int nW, int nH, int nC); //使用默认的内存布局
void Release();
//增加一个和常规图像的互
GPUImage(Image& m);
bool ToImage(Image& m);
bool FromImage(Image& m);
//Get
void* Ptr();
int Width();
int Height();
int Depth();
Rect Rect();
private:
GPUImageHeader* m_pHeader = nullptr;
};
1.2.1. 图像创建操作
图像创建操作:
void GPUImage::Create(int nW, int nH, int nC)
{
if (DebugCmd)
std::cout << "void Matrix::Create(int nW, int nH, int nC)" << std::endl;
//释放当前
Release();
m_pHeader = new GPUImageHeader();
m_pHeader->nWidth = nW;
m_pHeader->nHeight = nH;
m_pHeader->nChannel = nC;
size_t memSize = m_pHeader->nWidth * m_pHeader->nHeight * m_pHeader->nChannel + 31;
auto ret = cudaMalloc((void**)&(m_pHeader->pImage), memSize);
//如果ret不对,要抛异常
if (ret != cudaSuccess)
throw std::exception("GPUImage::FromImage(Image& m) error");
META_ATOMIC_ADD(&m_pHeader->nRefCount, 1);
if (DebugCmd)
std::cout << "m_Image.nRefCount:" << m_pHeader->nRefCount << std::endl;
}
采用cudamalloc申请图像空间。
1.2.2. 图像释放操作
图像释放操作:
void GPUImage::Release()
{
if (DebugCmd)
std::cout << "Release()" << std::endl;
if (m_pHeader)
{
assert(m_pHeader->nRefCount > 0);
if (META_ATOMIC_ADD(&m_pHeader->nRefCount, -1) == 1)
{
if (m_pHeader->pImage)
{
auto ret = cudaFree(m_pHeader->pImage);
//如果ret不对,要抛异常
if (ret != cudaSuccess)
throw std::exception("GPUImage::FromImage(Image& m) error");
}
m_pHeader->pImage = nullptr;
delete m_pHeader;
m_pHeader = nullptr;
if (DebugCmd)
std::cout << "Release::m_pHeader:delete" << std::endl;
}
else
{
if (DebugCmd)
std::cout << "Release::m_Image.nRefCount:" << m_pHeader->nRefCount << std::endl;
}
}
else
{
if (DebugCmd)
std::cout << "m_pHeader == nullptr" << std::endl;
}
}
释放时要判断引用计数,只有降为0,才进行显存的释放。
1.2.3. 图像拷贝构造和赋值构造
图像拷贝构造:
GPUImage::GPUImage(const GPUImage& m)
{
if (DebugCmd)
std::cout << "Matrix(const Matrix& m)" << std::endl;
if (m.m_pHeader)
{
META_ATOMIC_ADD(&m.m_pHeader->nRefCount, 1);
m_pHeader = m.m_pHeader;
if (DebugCmd)
std::cout << "m_Image.nRefCount:" << m_pHeader->nRefCount << std::endl;
}
else
{
m_pHeader = nullptr;
if (DebugCmd)
std::cout << "m_Image.nRefCount:0,m_pHeader == nullptr" << std::endl;
}
}
拷贝时,并不申请显存,而是将引用计算加1。
赋值构造:
GPUImage& GPUImage::operator=(const GPUImage& m)
{
if (DebugCmd)
std::cout << "Matrix& operator=(const Matrix& m)" << std::endl;
if (this->m_pHeader != m.m_pHeader)
{
if (m.m_pHeader)
{
META_ATOMIC_ADD(&m.m_pHeader->nRefCount, 1);
Release(); //释放当前
m_pHeader = m.m_pHeader;
if (DebugCmd)
std::cout << "m_Image.nRefCount:" << m_pHeader->nRefCount << std::endl;
}
else
{
m_pHeader = nullptr;
if (DebugCmd)
std::cout << "m_Image.nRefCount:0,m_pHeader == nullptr" << std::endl;
}
}
else
{
if (DebugCmd)
std::cout << "this->m_pHeader == m.m_pHeader" << std::endl;
}
return *this;
}
赋值构造时,要释放被赋值图像,,然后把赋值图像引用计数加1。
1.3. GPU图像数据结构引用计数测试
将上述DebugCmd打开,测试如下代码:
void test1(MetaCore::GPUImage imA)
{
std::cout << "期望引用计数会减1" << std::endl;
return;
}
void test2(MetaCore::GPUImage& imB)
{
std::cout << "期望不发生任何事情" << std::endl;
return;
}
MetaCore::GPUImage test3()
{