GPUMAT-访问自定义内核中的2个通道浮点数据
GpuMat - accessing 2 channel float data in custom kernel
据我所知,cv::cuda::PtrStep
用于将GpuMat
数据直接传递给自定义内核。我在此处找到了一个通道访问的示例,但是我的情况是2通道垫(CV_32FC2
)。在这种情况下,我试图实现复杂的绝对平方值,而复杂值的编码如:实际部分是第一平面,假想零件是给定Mat
的第二平面。
我尝试了:
__global__ void testKernel(const cv::cuda::PtrStepSz<cv::Vec2f> input, cv::cuda::PtrStepf output)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x <= input.cols - 1 && y <= input.rows - 1 && y >= 0 && x >= 0)
{
float val_re = input(x, y)[0];
float val_im = input(x, y) [1];
output(x, y) = val_re * val_re + val_im * val_im;
}
}
,但这会导致以下错误:
calling a __host__ function("cv::Vec<float, (int)2> ::operator []") from a __global__ function("gpuholo::testKernel") is not allowed
我明白了。[]
是__host__
限制功能,因为其cv::Vec2f
不是cv::cuda::Vec2f
(显然不存在)。但是我仍然真的很想访问数据。
是否还有其他机制可以访问设备侧面的2个通道数据,类似于Vec2f
?
我以将input
分为两个CV_32FC1
Mat
S的形式的解决方法,因此内核看起来像:
__global__ void testKernel(const cv::cuda::PtrStepSzf re, const cv::cuda::PtrStepSzf im, cv::cuda::PtrStepf output)
但是我想知道是否有"清洁剂"解决方案, Vec2f
-类似于一个。
您可以使用原始数据类型在自定义CUDA内核中访问GpuMat
的数据。例如float2
CUDA运行时提供的类型可以用作cv::Vec2f
的部分替换。这是一个示例代码,演示了用于访问GpuMat
数据的原始数据类型的使用情况。
#include <iostream>
#include <cuda_runtime.h>
#include <opencv2/opencv.hpp>
using std::cout;
using std::endl;
__global__ void kernel_absolute(float2* src, float* dst, int rows, int cols, int iStep, int oStep)
{
int i = blockIdx.y * blockDim.y + threadIdx.y; //Row number
int j = blockIdx.x * blockDim.x + threadIdx.x; //Column number
if (i<rows && j<cols)
{
/* Compute linear index from 2D indices */
int tidIn = i * iStep + j;
int tidOut = i * oStep + j;
/* Read input value */
float2 input = src[tidIn];
/* Calculate absolute value */
float output = sqrtf(input.x * input.x + input.y * input.y);
/* Write output value */
dst[tidOut] = output;
}
}
int main(int argc, char** argv)
{
/* Example to compute absolute value of each element of a complex matrix */
int rows = 10;
int cols = 10;
int input_data_type = CV_32FC2; //input is complex
int output_data_type = CV_32FC1; //output is real
/* Create input matrix on host */
cv::Mat input = cv::Mat::zeros(rows,cols,input_data_type) + cv::Vec2f(1,1) /* Initial value is (1,1) */;
/* Display input */
cout<<input<<endl;
/* Create input matrix on device */
cv::cuda::GpuMat input_d;
/* Copy from host to device */
input_d.upload(input);
/* Create output matrix on device */
cv::cuda::GpuMat output_d(rows,cols, output_data_type);
/* Compute element step value of input and output */
int iStep = input_d.step / sizeof(float2);
int oStep = output_d.step / sizeof(float);
/* Choose appropriate block size */
dim3 block(8,8);
/* Compute grid size using input size and block size */
dim3 grid ( (cols + block.x -1)/block.x, (rows + block.y -1)/block.y );
/* Launch CUDA kernel to compute absolute value */
kernel_absolute<<<grid, block>>>( reinterpret_cast<float2*>(input_d.data), reinterpret_cast<float*>(output_d.data), rows, cols, iStep, oStep );
/* Check kernel launch errors */
assert( cudaSuccess == cudaDeviceSynchronize() );
cv::Mat output;
/* Copy results from device to host */
output_d.download(output);
/* Display output */
cout<<endl<<output<<endl;
return 0;
}
用cuda 8.0上的Ubuntu上的以下命令进行了编译和测试:
nvcc -o complect.cu -ark = sm_61 -l/usr/local/lib -lopencv_core
如果要与内核单一输入一起工作,则可以将2频道图像弄平到1频道图像。
// test image
Mat h_mat(Size(50,50),CV_32FC2,Scalar(0.0));
// Mat::reshape takes number of channels and rows, for your example 1,1
Mat h_mat_flat = h_mat.reshape(1,1);
// to upload to gpu
GpuMat d_mat_flat(h_mat_flat.size(), h_mat_flat.type());
d_mat_flat.upload(h_mat_flat);
现在,您可以将d_mat_flat
作为PtrStepSzf
。
相关文章:
- 防止主数据类型C++的隐式转换
- 用于访问容器<T>数据成员的正确 API
- 嵌套在类中时无法设置成员数据
- 使用流处理接收到的数据
- 静态数据成员的问题-修复链接错误会导致编译器错误
- 处理小于cpu数据总线的数据类型.(c++转换为机器代码)
- 在cuda线程之间共享大量常量数据
- C++将文本文件中的数据读取到结构数组中
- 如何在C++中序列化结构数据
- 深度缓冲区未填充阴影贴图渲染通道中的数据
- OpenCV-浮动数据类型和强度范围在0-255范围内的RGB通道
- 通过用于UV4L数据通道的unix域套接字发送数据
- GPUMAT-访问自定义内核中的2个通道浮点数据
- 将一个通道数据复制到OpenCV中的另一个通道
- 打包左通道和右通道数据
- 在传输中未收到数据包而没有通道数据包丢失的原因
- 在双通道矩阵中插入数据
- 带alpha的WebP会丢失颜色通道数据
- 当"修改器"列表中有2个以上的修改器时,3ds Max C++修改器不会接收贴图通道数据
- 使用IOBluetooth框架的音频数据污染了我的RFCOMM通道