20 changed files with 643 additions and 158 deletions
@ -0,0 +1,42 @@ |
|||
#ifndef NURBSPERFORMER_LOOP_DETECTION_CUH |
|||
#define NURBSPERFORMER_LOOP_DETECTION_CUH |
|||
|
|||
#include "device/srf_mesh.cuh" |
|||
#include "map" |
|||
#include "set" |
|||
|
|||
class LoopDetection { |
|||
public: |
|||
SrfMesh *d_s_srfMesh; |
|||
SrfMesh *d_f_srfMesh; |
|||
|
|||
float s_paramRegionSize_u; |
|||
float s_paramRegionSize_v; |
|||
|
|||
int edgeSampleCnt; |
|||
|
|||
glm::vec<2, int> *d_selectedPointsIdx; |
|||
glm::vec2 *d_vectorFields; |
|||
int *d_rotationNumbers; |
|||
|
|||
LoopDetection(glm::vec3 *s_eval, glm::vec3 *f_eval, glm::vec3 *s_tan_u, glm::vec3 *f_tan_u, glm::vec3 *s_tan_v, |
|||
glm::vec3 *f_tan_v, glm::vec3 *s_norm, glm::vec3 *f_norm, int edgeSampleCnt, |
|||
float s_paramRegionSize_u, float s_paramRegionSize_v); |
|||
|
|||
__host__ void h_initOrientedDisAndVecFields( |
|||
const std::map<std::pair<int, int>, std::set<std::pair<int, int>>> &intersectBoxPairs); |
|||
__host__ void h_getRotationNumber(); |
|||
|
|||
~LoopDetection(); |
|||
}; |
|||
|
|||
__global__ void |
|||
g_initOrientedDisAndVecFields(SrfMesh *s_srfMesh, SrfMesh *f_srfMesh, int s_cellCnt, glm::vec<2, int> *s_cells, |
|||
glm::vec<2, int> **f_cells, const int *f_cellCnt, glm::vec<2, int> *selectedPointsIdx, |
|||
glm::vec2 *vectorFields); |
|||
|
|||
__global__ void |
|||
g_getRotationNumber(SrfMesh *s_srfMesh, float paramRegionSize_u, float paramRegionSize_v, glm::vec2 *vectorFields, |
|||
int *rotationNumbers); |
|||
|
|||
#endif //NURBSPERFORMER_LOOP_DETECTION_CUH |
|||
@ -0,0 +1,17 @@ |
|||
#ifndef NURBSPERFORMER_SRF_MESH_CUH |
|||
#define NURBSPERFORMER_SRF_MESH_CUH |
|||
|
|||
#include "glm/glm.hpp" |
|||
|
|||
class SrfMesh { |
|||
public: |
|||
int edgeSampleCnt; |
|||
glm::vec3 *evaluationRes; |
|||
glm::vec3 *tangent_u; |
|||
glm::vec3 *tangent_v; |
|||
glm::vec3 *normal; |
|||
SrfMesh(glm::vec3 *evalRes, glm::vec3 *tan_u, glm::vec3 *tan_v, glm::vec3 *nor, int edgeSampleCnt); |
|||
}; |
|||
|
|||
|
|||
#endif //NURBSPERFORMER_SRF_MESH_CUH |
|||
@ -0,0 +1,229 @@ |
|||
#include <vector> |
|||
#include "device/Nurbs/loop_detection.cuh" |
|||
#include "utils.h" |
|||
#include "cuda_runtime.h" |
|||
#include "device/device_utils.cuh" |
|||
|
|||
LoopDetection::LoopDetection(glm::vec3 *const s_eval, glm::vec3 *const f_eval, glm::vec3 *const s_tan_u, |
|||
glm::vec3 *const f_tan_u, glm::vec3 *const s_tan_v, glm::vec3 *const f_tan_v, |
|||
glm::vec3 *const s_norm, glm::vec3 *const f_norm, int edgeSampleCnt, |
|||
float s_paramRegionSize_u, float s_paramRegionSize_v) : |
|||
edgeSampleCnt(edgeSampleCnt), |
|||
s_paramRegionSize_u(s_paramRegionSize_u), |
|||
s_paramRegionSize_v(s_paramRegionSize_v) { |
|||
auto initSrfMesh = [](SrfMesh *&d_srfMesh, glm::vec3 *eval, glm::vec3 *tan_u, glm::vec3 *tan_v, glm::vec3 *norm, |
|||
int edgeSampleCnt) { |
|||
auto h_srfMesh = new SrfMesh(eval, tan_u, tan_v, norm, edgeSampleCnt); |
|||
safeCudaFree(d_srfMesh); |
|||
auto srfMeshSize = sizeof(SrfMesh); |
|||
cudaMalloc((void **) &d_srfMesh, srfMeshSize); |
|||
cudaMemcpy(d_srfMesh, h_srfMesh, srfMeshSize, cudaMemcpyHostToDevice); |
|||
safeFree(h_srfMesh); |
|||
}; |
|||
initSrfMesh(d_s_srfMesh, s_eval, s_tan_u, s_tan_v, s_norm, edgeSampleCnt); |
|||
initSrfMesh(d_f_srfMesh, f_eval, f_tan_u, f_tan_v, f_norm, edgeSampleCnt); |
|||
} |
|||
|
|||
|
|||
LoopDetection::~LoopDetection() { |
|||
safeCudaFree(d_selectedPointsIdx); |
|||
safeCudaFree(d_vectorFields); |
|||
safeCudaFree(d_s_srfMesh); |
|||
safeCudaFree(d_f_srfMesh); |
|||
safeCudaFree(d_rotationNumbers); |
|||
} |
|||
|
|||
__host__ void LoopDetection::h_initOrientedDisAndVecFields( |
|||
const std::map<std::pair<int, int>, std::set<std::pair<int, int>>> &intersectBoxPairs) { |
|||
int sampleCnt = edgeSampleCnt * edgeSampleCnt; |
|||
|
|||
safeCudaFree(d_selectedPointsIdx); |
|||
auto selectedPointsIdxVec = std::vector<glm::vec2>(sampleCnt, glm::vec<2, int>(-1, -1)); |
|||
auto h_selectedPointsIdx = new glm::vec<2, int>[sampleCnt]; |
|||
std::copy(selectedPointsIdxVec.begin(), selectedPointsIdxVec.end(), h_selectedPointsIdx); |
|||
int selectedPtsIdxBytes = sizeof(glm::vec<2, int>) * sampleCnt; |
|||
cudaMalloc((void**)&d_selectedPointsIdx, selectedPtsIdxBytes); |
|||
cudaMemcpy(d_selectedPointsIdx, h_selectedPointsIdx, selectedPtsIdxBytes, cudaMemcpyHostToDevice); |
|||
|
|||
safeCudaFree(d_vectorFields); |
|||
cudaMalloc((void**)&d_vectorFields, sizeof(glm::vec2) * sampleCnt); |
|||
|
|||
int mapCnt = intersectBoxPairs.size(); |
|||
|
|||
auto *h_s_cells = new glm::vec<2, int>[mapCnt]; |
|||
glm::vec<2, int> *d_s_cells; |
|||
auto **h_f_cells = new glm::vec<2, int> *[mapCnt]; |
|||
glm::vec<2, int> **d_f_cells; |
|||
auto *h_f_cellsCnt = new int[mapCnt]; |
|||
int *d_f_cellsCnt; |
|||
int i = 0; |
|||
for (const auto &boxPair: intersectBoxPairs) { |
|||
h_s_cells[i] = glm::vec<2, int>(boxPair.first.first, boxPair.first.second); |
|||
h_f_cellsCnt[i] = boxPair.second.size(); |
|||
auto h_f_cellsTmp = new glm::vec<2, int>[h_f_cellsCnt[i]]; |
|||
int j = 0; |
|||
for (auto box: boxPair.second) { |
|||
h_f_cellsTmp[j] = glm::vec<2, int>(box.first, box.second); |
|||
j++; |
|||
} |
|||
auto f_relatedCellsBytes = sizeof(glm::vec<2, int>) * h_f_cellsCnt[i]; |
|||
cudaMalloc((void **) &h_f_cells[i], f_relatedCellsBytes); |
|||
cudaMemcpy(h_f_cells[i], h_f_cellsTmp, f_relatedCellsBytes, cudaMemcpyHostToDevice); |
|||
free(h_f_cellsTmp); |
|||
i++; |
|||
} |
|||
|
|||
int s_cellsBytes = sizeof(glm::vec<2, int>) * mapCnt; |
|||
cudaMalloc((void **) &d_s_cells, s_cellsBytes); |
|||
cudaMemcpy(d_s_cells, h_s_cells, s_cellsBytes, cudaMemcpyHostToDevice); |
|||
|
|||
int f_cellsPtrBytes = sizeof(glm::vec<2, int>) * mapCnt; |
|||
cudaMalloc((void **) &d_f_cells, f_cellsPtrBytes); |
|||
cudaMemcpy(d_f_cells, h_f_cells, f_cellsPtrBytes, cudaMemcpyHostToDevice); |
|||
|
|||
int f_cellsCntBytes = sizeof(int) * mapCnt; |
|||
cudaMalloc((void **) &d_f_cellsCnt, f_cellsCntBytes); |
|||
cudaMemcpy(d_f_cellsCnt, h_f_cellsCnt, f_cellsCntBytes, cudaMemcpyHostToDevice); |
|||
|
|||
dim3 block(64, 64); |
|||
dim3 grid((mapCnt + block.x * block.y - 1) / (block.x * block.y)); |
|||
g_initOrientedDisAndVecFields <<< grid, block >>>(d_s_srfMesh, d_f_srfMesh, mapCnt, d_s_cells, d_f_cells, |
|||
d_f_cellsCnt, d_selectedPointsIdx, d_vectorFields); |
|||
|
|||
cudaFree(d_s_cells); |
|||
cudaFree(d_f_cells); |
|||
cudaFree(d_f_cellsCnt); |
|||
for (int k = 0; k < mapCnt; k++) { |
|||
cudaFree(h_f_cells[k]); |
|||
} |
|||
free(h_s_cells); |
|||
free(h_f_cells); |
|||
free(h_f_cellsCnt); |
|||
free(h_selectedPointsIdx); |
|||
} |
|||
|
|||
__host__ void LoopDetection::h_getRotationNumber() { |
|||
safeCudaFree(d_rotationNumbers); |
|||
cudaMalloc((void**)&d_rotationNumbers, sizeof(int) * edgeSampleCnt * edgeSampleCnt); |
|||
|
|||
dim3 block(32, 32); |
|||
dim3 grid((edgeSampleCnt + block.x - 1) / block.x, (edgeSampleCnt + block.y - 1) / block.y); |
|||
g_getRotationNumber<<<grid, block>>>(d_s_srfMesh, s_paramRegionSize_u, s_paramRegionSize_v, d_vectorFields, |
|||
d_rotationNumbers); |
|||
|
|||
} |
|||
|
|||
__device__ const int neighborIdxes[4][2] = {{1, 0}, |
|||
{0, 1}, |
|||
{-1, 0}, |
|||
{0, -1}}; |
|||
|
|||
__device__ const int dirs[2][2][2] = {{{-1, 1}, {1, 1}}, |
|||
{{-1, -1}, {1, -1}}}; // 必须是左上，右上，左下，右下的顺序 |
|||
|
|||
__global__ void |
|||
g_initOrientedDisAndVecFields(SrfMesh *s_srfMesh, SrfMesh *f_srfMesh, int s_cellCnt, glm::vec<2, int> *s_cells, |
|||
glm::vec<2, int> **f_cells, const int *f_cellCnt, glm::vec<2, int> *selectedPointsIdx, |
|||
glm::vec2 *vectorFields) { |
|||
// // 二维的grid和一维的block |
|||
// int idx = (blockIdx.y * gridDim.x + blockIdx.x) * blockDim.x * blockDim.y + threadIdx.x; |
|||
// 一维的grid和二维的block |
|||
int idx = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x; |
|||
if (idx >= s_cellCnt) return; |
|||
// printf("idx: %d\n", idx); |
|||
if (idx == 0) printf("<%d, %d>, %f\n", neighborIdxes[0][0], neighborIdxes[0][1], -D_FLT_MAX); |
|||
int s_edgeSampleCnt = s_srfMesh->edgeSampleCnt; |
|||
int f_edgeSampleCnt = f_srfMesh->edgeSampleCnt; |
|||
for (auto neighborIdx: neighborIdxes) { |
|||
for (int i = 0; i < 2; i++) { |
|||
for (int j = 0; j < 2; j++) { |
|||
float minDis = -D_FLT_MAX; |
|||
int minDisFIdx_u = -1, minDisFIdx_v = -1; |
|||
int idxI = s_cells[idx].x + i + neighborIdx[0]; |
|||
int idxJ = s_cells[idx].y + j + neighborIdx[1]; |
|||
if (idxI < 0 || idxI >= s_edgeSampleCnt || idxJ < 0 || idxJ >= s_edgeSampleCnt || |
|||
selectedPointsIdx[idxI * s_edgeSampleCnt + idxJ] != glm::vec<2, int>(2, 2)) { |
|||
// 尽量不重复计算 |
|||
// 但在多线程情况下，这样还是可以能重复计算 |
|||
continue; |
|||
} |
|||
for (int fi = 0; fi < f_cellCnt[idx]; fi++) { |
|||
for (int k = 0; k < 2; k++) { |
|||
for (int l = 0; l < 2; l++) { |
|||
for (int step_u = 0; step_u <= 2; step_u++) { |
|||
for (int step_v = 0; step_v + step_u <= 2; step_v++) { |
|||
int fIdx_u = f_cells[idx][fi].x + k + step_u * dirs[k][l][0]; |
|||
int fIdx_v = f_cells[idx][fi].y + l + step_v * dirs[k][l][1]; |
|||
if (fIdx_u < 0 || fIdx_u >= f_edgeSampleCnt || fIdx_v < 0 || |
|||
fIdx_v >= f_edgeSampleCnt) |
|||
continue; |
|||
auto dis = glm::distance(s_srfMesh->evaluationRes[idxI][idxJ], |
|||
f_srfMesh->evaluationRes[fIdx_u][fIdx_v]); |
|||
if (dis < minDis) { |
|||
minDis = dis; |
|||
minDisFIdx_u = fIdx_u; |
|||
minDisFIdx_v = fIdx_v; |
|||
} |
|||
} |
|||
} |
|||
} |
|||
} |
|||
} |
|||
selectedPointsIdx[idxI * s_edgeSampleCnt + idxJ] = glm::vec<2, int>(minDisFIdx_u, minDisFIdx_v); |
|||
auto n2 = f_srfMesh->normal[minDisFIdx_u * f_edgeSampleCnt + minDisFIdx_v]; |
|||
auto ps_pu = s_srfMesh->tangent_u[idxI * s_edgeSampleCnt + idxJ]; |
|||
auto ps_pv = s_srfMesh->tangent_v[idxI * s_edgeSampleCnt + idxJ]; |
|||
vectorFields[idxI * s_edgeSampleCnt + idxJ] = glm::vec2(glm::dot(n2, ps_pu), glm::dot(n2, ps_pv)); |
|||
} |
|||
} |
|||
} |
|||
} |
|||
|
|||
__global__ void |
|||
g_getRotationNumber(SrfMesh *s_srfMesh, float paramRegionSize_u, float paramRegionSize_v, glm::vec2 *vectorFields, |
|||
int *rotationNumbers) { |
|||
// 二维grid和二维的block |
|||
int ix = blockIdx.x * blockDim.x + threadIdx.x; |
|||
int iy = blockIdx.y * blockDim.y + threadIdx.y; |
|||
int s_edgeSampleCnt = s_srfMesh->edgeSampleCnt; |
|||
float F[2][2], G[2][2]; |
|||
for (int biasI = 0; biasI < 2; biasI++) { |
|||
for (int biasJ = 0; biasJ < 2; biasJ++) { |
|||
auto idxI = ix + biasI; |
|||
auto idxJ = iy + biasJ; |
|||
auto v = vectorFields[idxI * s_edgeSampleCnt + idxJ]; |
|||
auto vSquareSum = v.x * v.x + v.y * v.y; |
|||
auto pTheta_pChi = -v.y / vSquareSum; |
|||
auto pTheta_pPhi = v.x / vSquareSum; |
|||
// auto pChi_pu = |
|||
// (vectorFields[idxI + 1][idxJ].x - vectorFields[idxI - 1][idxJ].x) / (2 * uParamCellSize); |
|||
// auto pChi_pv = |
|||
// (vectorFields[idxI][idxJ + 1].x - vectorFields[idxI][idxJ - 1].x) / (2 * vParamCellSize); |
|||
glm::vec2 pV_pu; |
|||
glm::vec2 pV_pv; |
|||
if (idxI == 0) |
|||
pV_pu = (vectorFields[(idxI + 1) * s_edgeSampleCnt + idxJ] - |
|||
vectorFields[idxI * s_edgeSampleCnt + idxJ]) / paramRegionSize_u; |
|||
else if (idxI == s_edgeSampleCnt - 1) |
|||
pV_pu = (vectorFields[idxI * s_edgeSampleCnt + idxJ] - |
|||
vectorFields[(idxI - 1) * s_edgeSampleCnt + idxJ]) / paramRegionSize_u; |
|||
else |
|||
pV_pu = (vectorFields[(idxI + 1) * s_edgeSampleCnt + idxJ] - |
|||
vectorFields[(idxI - 1) * s_edgeSampleCnt + idxJ]) / (2 * paramRegionSize_u); |
|||
if (idxJ == 0) |
|||
pV_pv = (vectorFields[idxI * s_edgeSampleCnt + idxJ + 1] - |
|||
vectorFields[idxI * s_edgeSampleCnt + idxJ]) / paramRegionSize_v; |
|||
else if (idxJ == s_edgeSampleCnt - 1) |
|||
pV_pv = (vectorFields[idxI * s_edgeSampleCnt + idxJ] - |
|||
vectorFields[idxI * s_edgeSampleCnt + idxJ - 1]) / paramRegionSize_v; |
|||
else |
|||
pV_pv = (vectorFields[idxI * s_edgeSampleCnt + idxJ + 1] - |
|||
vectorFields[idxI * s_edgeSampleCnt + idxJ - 1]) / (2 * paramRegionSize_v); |
|||
F[biasI][biasJ] = pTheta_pChi * pV_pu.x + pTheta_pPhi * pV_pu.y; |
|||
G[biasI][biasJ] = pTheta_pChi * pV_pv.x + pTheta_pPhi * pV_pv.y; |
|||
} |
|||
} |
|||
rotationNumbers[ix * s_edgeSampleCnt + iy] = int( |
|||
round(((F[0][1] + F[1][1] - F[0][0] - F[1][0]) * paramRegionSize_u + |
|||
(G[0][0] + G[0][1] - G[1][0] - G[1][1]) * paramRegionSize_v)) / 2.); |
|||
} |
|||
@ -0,0 +1,4 @@ |
|||
#include "device/srf_mesh.cuh" |
|||
|
|||
SrfMesh::SrfMesh(glm::vec3 *evalRes, glm::vec3 *tan_u, glm::vec3 *tan_v, glm::vec3 *nor, int edgeSampleCnt) : |
|||
evaluationRes(evalRes), tangent_u(tan_u), tangent_v(tan_v), normal(nor), edgeSampleCnt(edgeSampleCnt) {} |
|||
Loading…
Reference in new issue