基于CUDA实现立方体贴图 (Cubemaps) 转换为全景图 (Equirectangular Panorama)
生活随笔
收集整理的這篇文章主要介紹了
基于CUDA实现立方体贴图 (Cubemaps) 转换为全景图 (Equirectangular Panorama)
小編覺得挺不錯的,現在分享給大家,幫大家做個參考.
在立方體貼圖空間內發射光線(視線),計算球面光線(視線)會擊中哪個面的哪個像素的像素值,最終生成Equirectangular全景圖。
- InitSceneTexture():先獲取Cubemaps并將其綁定在GPU中
void InitSceneTexture() {char* path = "./Cubemaps";myEnvironment.setFileRoute(path);myEnvironment.getImageSource();int sceneLength = myEnvironment.Top.width;int sceneHeight = myEnvironment.Top.height;bind2DTopTexture(sceneLength, sceneLength);bind2DBottomTexture(sceneLength, sceneLength);bind2DLeftTexture(sceneLength, sceneLength);bind2DRightTexture(sceneLength, sceneLength);bind2DFrontTexture(sceneLength, sceneLength);bind2DBackTexture(sceneLength, sceneLength);
}
最后6個函數就是綁定紋理的過程,以top面為例,其它完整的代碼會放在后面
void bind2DTopTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Top.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texTopEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texTopEnvironment, arr_d, &texTopEnvironment.channelDesc, width, height, pitch);
}
- 需要注意的是:不要
cudaFree(arr_d )
- 根據視點viewPoint(默認是(0,0,0))和視線方向viewVector來確定擊中的像素。在這里我們立方體貼圖是設置的1000x1000的,所以中心位置是500
DEVICE Color renderScene(point& viewPoint, Vector& viewVector) {//float tx, ty, tz;bool positiveX, positiveY, positiveZ;float boxBoder = 500;insideIntersect_x(viewPoint, viewVector, boxBoder, tx, positiveX);insideIntersect_y(viewPoint, viewVector, boxBoder, ty, positiveY);insideIntersect_z(viewPoint, viewVector, boxBoder, tz, positiveZ);if (tx <= ty&&tx <= tz) {if (positiveX) {viewPoint += tx*viewVector;float xPos = viewPoint.z + boxBoder;float yPos = viewPoint.y + boxBoder;float4 hagles = tex2D(texRightEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}else {viewPoint += tx*viewVector;float xPos = boxBoder - viewPoint.z;float yPos = viewPoint.y + boxBoder;float4 hagles = tex2D(texLeftEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}}else if (ty <= tx&&ty <= tz) {if (positiveY) {viewPoint += ty*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder - viewPoint.z;float4 hagles = tex2D(texTopEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}else {viewPoint += ty*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder + viewPoint.z;float4 hagles = tex2D(texBottomEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}}else if (tz <= ty&&tz <= tx) {if (positiveZ) {viewPoint += tz*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder + viewPoint.y;float4 hagles = tex2D(texBackEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}else {viewPoint += tz*viewVector;float xPos = viewPoint.x + boxBoder;float yPos = viewPoint.y + boxBoder;float4 hagles = tex2D(texFrontEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}}
}
Color, point, Vector都是類似于cv::Point3f,不過需要基于HOST_AND_DEVICE重構數據結構。
下面這三個函數是為了確定擊中的是哪個面
insideIntersect_x(viewPoint, viewVector, boxBoder, tx, positiveX);insideIntersect_y(viewPoint, viewVector, boxBoder, ty, positiveY);insideIntersect_z(viewPoint, viewVector, boxBoder, tz, positiveZ);
完整代碼:
//irradiance.cu
__global__ void kernel8(int range_W, int range_H, cv::cuda::PtrStepSz<vec3f> outPano)
{point viewpoint = point(0, 0, 0);int X = blockIdx.x * blockDim.x + threadIdx.x;int Y = blockIdx.y * blockDim.y + threadIdx.y;if (X < range_W && Y < range_H) {vec3f& pixel = outPano(1000 - Y, X);float theta = 2 * PI * (float)X / (float)range_W - PI;float phi = PI * (float)Y / (float)range_H - PI / 2.0;Vector dir = Vector(sin(theta) * cos(phi), cos(phi) * cos(theta), sin(phi));Color hit_color = renderScene(viewpoint, dir);pixel[0] = hit_color.z * 255.0;pixel[1] = hit_color.y * 255.0;pixel[2] = hit_color.x * 255.0;}}cv::Mat Cube2Pano()
{InitSceneTexture();dim3 thread(32, 32);dim3 block((2000 + thread.x - 1) / thread.x, (1000 + thread.y - 1) / thread.y);cv::cuda::GpuMat device_panorama = cv::cuda::GpuMat(1000, 2000, CV_32FC3);kernel8 << <block, thread >> > (2000, 1000, device_panorama);cudaThreadSynchronize();cv::Mat host_pano;device_panorama.download(host_pano);return host_pano;
}
//irradiance.cuh
#pragma once
#include "Enviro.cuh"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"extern "C" cv::Mat Cube2Pano();
//Enviro.cuh
#pragma once
#include "string"
#include <opencv2/opencv.hpp>
#include "environment.h"
#define DEVICE __device__
#define HOD __host__ __device__
using namespace xcx;texture<float4, cudaTextureType2D, cudaReadModeElementType> texBottomEnvironment;
texture<float4, cudaTextureType2D, cudaReadModeElementType> texTopEnvironment;
texture<float4, cudaTextureType2D, cudaReadModeElementType> texLeftEnvironment;
texture<float4, cudaTextureType2D, cudaReadModeElementType> texRightEnvironment;
texture<float4, cudaTextureType2D, cudaReadModeElementType> texFrontEnvironment;
texture<float4, cudaTextureType2D, cudaReadModeElementType> texBackEnvironment;void bind2DTopTexture(int width, int height);
void bind2DBottomTexture(int width, int height);
void bind2DLeftTexture(int width, int height);
void bind2DRightTexture(int width, int height);
void bind2DFrontTexture(int width, int height);
void bind2DBackTexture(int width, int height);environment myEnvironment;void InitSceneTexture() {char* path = "./Image_out/segmentation";myEnvironment.setFileRoute(path);myEnvironment.getImageSource();int sceneLength = myEnvironment.Top.width;int sceneHeight = myEnvironment.Top.height;bind2DTopTexture(sceneLength, sceneLength);bind2DBottomTexture(sceneLength, sceneLength);bind2DLeftTexture(sceneLength, sceneLength);bind2DRightTexture(sceneLength, sceneLength);bind2DFrontTexture(sceneLength, sceneLength);bind2DBackTexture(sceneLength, sceneLength);}void bind2DTopTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Top.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texTopEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texTopEnvironment, arr_d, &texTopEnvironment.channelDesc, width, height, pitch);
}
void bind2DBottomTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Bottom.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texBottomEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texBottomEnvironment, arr_d, &texBottomEnvironment.channelDesc, width, height, pitch);
}
void bind2DLeftTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Left.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texLeftEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texLeftEnvironment, arr_d, &texLeftEnvironment.channelDesc, width, height, pitch);
}
void bind2DRightTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Right.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texRightEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texRightEnvironment, arr_d, &texRightEnvironment.channelDesc, width, height, pitch);
}
void bind2DFrontTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Front.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texFrontEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texFrontEnvironment, arr_d, &texFrontEnvironment.channelDesc, width, height, pitch);
}
void bind2DBackTexture(int width, int height) {size_t pitch, tex_ofs;float4*arr_d = 0;cudaMallocPitch((void**)&arr_d, &pitch, width * sizeof(float4), height);cudaMemcpy2D(arr_d, pitch, myEnvironment.Back.normData, width * sizeof(float4),width * sizeof(float4), height, cudaMemcpyHostToDevice);texBackEnvironment.normalized = false;cudaBindTexture2D(&tex_ofs, &texBackEnvironment, arr_d, &texBackEnvironment.channelDesc, width, height, pitch);
}DEVICE void insideIntersect_x(point& viewPoint, Vector& viewVector, float boxBoder, float& T, bool& Positive) {float t1, t2, txm, txi;if (viewVector.x != 0.0) {t1 = (boxBoder - viewPoint.x) / viewVector.x;t2 = (-boxBoder - viewPoint.x) / viewVector.x;txm = ((t1 < t2) ? t1 : t2);txi = ((t1 < t2) ? t2 : t1);if (t1 > t2)Positive = true;else Positive = false;}else {txm = 0;txi = 100000;Positive = true;}T = txi;
}
DEVICE void insideIntersect_y(point&viewPoint, Vector&viewVector, float boxBoder, float& T, bool& Positive) {float t1, t2, txm, txi;if (viewVector.y != 0.0) {t1 = (boxBoder - viewPoint.y) / viewVector.y;t2 = (-boxBoder - viewPoint.y) / viewVector.y;txm = ((t1 < t2) ? t1 : t2);txi = ((t1 < t2) ? t2 : t1);if (t1 > t2)Positive = true;else Positive = false;}else {txm = 0;txi = 100000;Positive = true;}T = txi;
}DEVICE void insideIntersect_z(point& viewPoint, xcx::Vector&viewVector, float boxBoder, float& T, bool& Positive) {float t1, t2, txm, txi;if (viewVector.z != 0.0) {t1 = (boxBoder - viewPoint.z) / viewVector.z;t2 = (-boxBoder - viewPoint.z) / viewVector.z;txm = ((t1 < t2) ? t1 : t2);txi = ((t1 < t2) ? t2 : t1);if (t1 > t2)Positive = true;else Positive = false;}else {txm = 0;txi = 100000;Positive = true;}T = txi;
}DEVICE Color renderScene(point& viewPoint, Vector& viewVector) {//float tx, ty, tz;bool positiveX, positiveY, positiveZ;float boxBoder = 500;insideIntersect_x(viewPoint, viewVector, boxBoder, tx, positiveX);insideIntersect_y(viewPoint, viewVector, boxBoder, ty, positiveY);insideIntersect_z(viewPoint, viewVector, boxBoder, tz, positiveZ);if (tx <= ty&&tx <= tz) {if (positiveX) {viewPoint += tx*viewVector;float xPos = viewPoint.z + boxBoder;float yPos = viewPoint.y + boxBoder;float4 hagles = tex2D(texRightEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}else {viewPoint += tx*viewVector;float xPos = boxBoder - viewPoint.z;float yPos = viewPoint.y + boxBoder;float4 hagles = tex2D(texLeftEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}}else if (ty <= tx&&ty <= tz) {if (positiveY) {viewPoint += ty*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder - viewPoint.z;float4 hagles = tex2D(texTopEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}else {viewPoint += ty*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder + viewPoint.z;float4 hagles = tex2D(texBottomEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}}else if (tz <= ty&&tz <= tx) {if (positiveZ) {viewPoint += tz*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder + viewPoint.y;float4 hagles = tex2D(texBackEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}else {viewPoint += tz*viewVector;float xPos = viewPoint.x + boxBoder;float yPos = viewPoint.y + boxBoder;float4 hagles = tex2D(texFrontEnvironment, (xPos + 0.5), (yPos + 0.5));return Color(hagles.x, hagles.y, hagles.z);}}
}DEVICE point positionatface(point& viewPoint, Vector& viewVector, int& faceIndex) {float tx, ty, tz;bool positiveX, positiveY, positiveZ;float boxBoder = 500;insideIntersect_x(viewPoint, viewVector, boxBoder, tx, positiveX);insideIntersect_y(viewPoint, viewVector, boxBoder, ty, positiveY);insideIntersect_z(viewPoint, viewVector, boxBoder, tz, positiveZ);if (tx <= ty&&tx <= tz) {if (positiveX) {viewPoint += tx*viewVector;float xPos = viewPoint.z + boxBoder;float yPos = viewPoint.y + boxBoder;point position = point(xPos, yPos);faceIndex = 1;return position;}else {viewPoint += tx*viewVector;float xPos = boxBoder - viewPoint.z;float yPos = viewPoint.y + boxBoder;point position = point(xPos, yPos);faceIndex = 3;return position;}}else if (ty <= tx&&ty <= tz) {if (positiveY) {viewPoint += ty*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder - viewPoint.z;point position = point(xPos, yPos);faceIndex = 4;return position;}else {viewPoint += ty*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder + viewPoint.z;point position = point(xPos, yPos);faceIndex = 5;return position;}}else if (tz <= ty&&tz <= tx) {if (positiveZ) {viewPoint += tz*viewVector;float xPos = boxBoder - viewPoint.x;float yPos = boxBoder + viewPoint.y;point position = point(xPos, yPos);faceIndex = 2;return position;}else {viewPoint += tz*viewVector;float xPos = viewPoint.x + boxBoder;float yPos = viewPoint.y + boxBoder;point position = point(xPos, yPos);faceIndex = 0;return position;}}
}
//environment.h
#pragma once
#include "./Util/xcx_geometry.h"
#include <string>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include "./OpenGLtools/stb_image.h"#define Fms_RGB 3
#define Fms_RGBA 4using std::string;class singleImage {public:string filePath;unsigned char* data;float4 *normData;int width, height, nrComponents;int format;void getImage(string path) {filePath = path;data = NULL;stbi_set_flip_vertically_on_load(true);data = stbi_load(path.c_str(), &width, &height, &nrComponents, 0);normData = (float4*)malloc(width*height * sizeof(float4));for (int i = 0; i < height; i++) {for (int j = 0; j < width; j++) {if (nrComponents == Fms_RGB) {normData[i*width + j].x = (float)data[(i*width + j) * 3 + 0] / 255.0;normData[i*width + j].y = (float)data[(i*width + j) * 3 + 1] / 255.0;normData[i*width + j].z = (float)data[(i*width + j) * 3 + 2] / 255.0;}else if (nrComponents == Fms_RGBA) {normData[i*width + j].x = (float)data[(i*width + j) * 4 + 0] / 255.0;normData[i*width + j].y = (float)data[(i*width + j) * 4 + 1] / 255.0;normData[i*width + j].z = (float)data[(i*width + j) * 4 + 2] / 255.0;}}}}singleImage() {}singleImage& operator = (const singleImage& Other) {width = Other.width;height = Other.height;data = Other.data;nrComponents = Other.nrComponents;format = Other.format;return *this;}singleImage(const singleImage& Other){*this = Other;};};
class environment {public:environment() {}void setFileRoute(string fileRoute) {filePath = fileRoute;}void getImageSource() {Bottom.getImage(filePath + "/bottom.jpg");Top.getImage(filePath + "/top.jpg");Left.getImage(filePath + "/left.jpg");Right.getImage(filePath + "/right.jpg");Front.getImage(filePath + "/front.jpg");Back.getImage(filePath + "/back.jpg");}environment& operator = (const environment& Other) {Bottom = Other.Bottom;Top = Other.Top;Front = Other.Front;Back = Other.Back;Left = Other.Left;Right = Other.Right;return *this;}environment(const environment& Other) {*this = Other;}
public:std::string filePath;singleImage Bottom;singleImage Top;singleImage Front;singleImage Back;singleImage Left;singleImage Right;};
總結
以上是生活随笔為你收集整理的基于CUDA实现立方体贴图 (Cubemaps) 转换为全景图 (Equirectangular Panorama)的全部內容,希望文章能夠幫你解決所遇到的問題。
- 上一篇: Python调整图像亮度和饱和度
- 下一篇: Tensorflow::Session