18 changed files with 480 additions and 68 deletions
@ -0,0 +1,29 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/11. |
|||
// |
|||
|
|||
#ifndef NURBSEVALUATOR_BVH_CUH |
|||
#define NURBSEVALUATOR_BVH_CUH |
|||
|
|||
#include "../aabb.cuh" |
|||
|
|||
struct BVHNode { |
|||
AABB bounds; // AABB包围盒 |
|||
int firstChild, level, childNum = 4; // 第一个孩子节点的下标，该结点所在层次，孩子个数 |
|||
// 曲面片u, v的范围 |
|||
float u0, u1, v0, v1; |
|||
}; |
|||
|
|||
class BVH { |
|||
public: |
|||
int maxLevel, size; |
|||
BVHNode* nodes = nullptr; |
|||
|
|||
__host__ void printQuadTree(); |
|||
}; |
|||
|
|||
__global__ void |
|||
buildSurfaceBvh(const float *k, int level, const float *evaluatedPoints, float d_lastKnot_u, float d_lastKnot_v, |
|||
int d_sampleCnt_u, int d_sampleCnt_v, BVHNode *BVH); |
|||
|
|||
#endif //NURBSEVALUATOR_BVH_CUH |
|||
@ -0,0 +1,28 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/11. |
|||
// |
|||
|
|||
#ifndef NURBSEVALUATOR_AABB_CUH |
|||
#define NURBSEVALUATOR_AABB_CUH |
|||
|
|||
#include "vec.cuh" |
|||
#include "cstdio" |
|||
|
|||
class AABB { |
|||
public: |
|||
// 边界 |
|||
FVec3 pMin, pMax; |
|||
__device__ __host__ AABB(); |
|||
__device__ __host__ explicit AABB(const FVec3& p); |
|||
__device__ __host__ AABB(const FVec3& p1, const FVec3& p2); |
|||
|
|||
// aabb包围盒合并操作，包围盒和点 |
|||
__device__ __host__ AABB Union(const FVec3& p); |
|||
// aabb包围盒合并操作，两个包围盒 |
|||
__device__ __host__ AABB Union(const AABB& b2); |
|||
// 判断两个aabb包围盒是否重叠 |
|||
__device__ __host__ bool IsOverlap(const AABB& b2); |
|||
}; |
|||
|
|||
|
|||
#endif //NURBSEVALUATOR_AABB_CUH |
|||
@ -0,0 +1,19 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/11. |
|||
// |
|||
|
|||
#ifndef NURBSEVALUATOR_VEC_CUH |
|||
#define NURBSEVALUATOR_VEC_CUH |
|||
#include "cuda_runtime.h" |
|||
|
|||
class FVec3 { |
|||
public: |
|||
float x, y, z; |
|||
__device__ __host__ FVec3(float x_, float y_, float z_); |
|||
__device__ __host__ FVec3(FVec3 const &fVec3); |
|||
|
|||
__device__ __host__ FVec3(); |
|||
}; |
|||
|
|||
|
|||
#endif //NURBSEVALUATOR_VEC_CUH |
|||
@ -0,0 +1,137 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/11. |
|||
// |
|||
|
|||
#include "../../../include/device/Nurbs/bvh.cuh" |
|||
/** |
|||
* 当前层的第一个节点的坐标 |
|||
* @param layer 层号。根节点为第一层 |
|||
*/ |
|||
__device__ int getStartIdxOfLayerN(int layer) { |
|||
if (layer == 1) return 0; |
|||
int num = 1; |
|||
// 找规律得出 |
|||
for (int i = 2; i < layer; i++) { |
|||
num = num * 4 + 1; |
|||
} |
|||
return num; |
|||
} |
|||
|
|||
/** |
|||
* 根据u、v参数域中的小矩形的位置，判断它在BVH叶节点层（也就是最后一层）中的位置 |
|||
* @param ix |
|||
* @param iy |
|||
* @return |
|||
*/ |
|||
__device__ int getChildNodeIdx(int ix, int iy) { |
|||
int idx = 0; |
|||
while(ix > 0) { |
|||
int log2XCeil = int(__log2f(ix)); |
|||
idx += int(pow(4, log2XCeil)); |
|||
ix -= int(pow(2, log2XCeil)); |
|||
} |
|||
while(iy > 0) { |
|||
int log2XCeil = int(__log2f(iy)); |
|||
idx += 2 * int(pow(4, log2XCeil)); |
|||
iy -= int(pow(2, log2XCeil)); |
|||
} |
|||
return idx; |
|||
} |
|||
|
|||
|
|||
__global__ void |
|||
buildSurfaceBvh(const float *k, int level, const float *evaluatedPoints, float lastKnot_u, float lastKnot_v, |
|||
int sampleCnt_u, int sampleCnt_v, BVHNode *BVH) { |
|||
// 假设还是二维grid和二维的block |
|||
int ix = blockIdx.x * blockDim.x + threadIdx.x; |
|||
int iy = blockIdx.y * blockDim.y + threadIdx.y; |
|||
if (ix >= sampleCnt_u - 1 || iy >= sampleCnt_v - 1) { |
|||
return; |
|||
} |
|||
float singleStepVal_u = lastKnot_u / (sampleCnt_u - 1); |
|||
float singleStepVal_v = lastKnot_v / (sampleCnt_v - 1); |
|||
|
|||
float u = singleStepVal_u * ix, v = singleStepVal_v * iy; |
|||
|
|||
int bias = getChildNodeIdx(ix, iy); // 每层的节点在总数组中的编号相对该层第一个节点的偏置 |
|||
for (int step = 1; step <= sampleCnt_u - 1 || step <= sampleCnt_v - 1; step = step * 2, --level) { |
|||
|
|||
float step_u = min(step, sampleCnt_u - 1); |
|||
float step_v = min(step, sampleCnt_v - 1); |
|||
|
|||
int baseIdx = getStartIdxOfLayerN(level); // 当前层的第一个节点的坐标 |
|||
int idx = baseIdx + bias; |
|||
|
|||
int tmpIdx = (ix * sampleCnt_v + iy) * 3; |
|||
// AABB bound(FVec3(evaluatedPoints[tmpIdx] + *k, evaluatedPoints[tmpIdx + 1] + *k, |
|||
// evaluatedPoints[tmpIdx + 2] + *k)); // 最初只包含u，v点 |
|||
AABB bound(FVec3(evaluatedPoints[tmpIdx], evaluatedPoints[tmpIdx + 1], |
|||
evaluatedPoints[tmpIdx + 2])); // 最初只包含u，v点 |
|||
if(idx == 75) printf("75: %g, %g, %g\n", evaluatedPoints[tmpIdx], evaluatedPoints[tmpIdx + 1], evaluatedPoints[tmpIdx + 2]); |
|||
if (step == 1) { |
|||
// 叶子节点 |
|||
BVH[idx].u0 = u; |
|||
BVH[idx].u1 = u + step_u * singleStepVal_u; |
|||
BVH[idx].v0 = v; |
|||
BVH[idx].v1 = v + step_v * singleStepVal_v; |
|||
BVH[idx].level = level; |
|||
// int tmpIdx1 = ((ix + 1) * sampleCnt_v + iy) * 3; |
|||
// bound = bound.Union(FVec3(evaluatedPoints[tmpIdx1] + *k, evaluatedPoints[tmpIdx1 + 1] + *k, |
|||
// evaluatedPoints[tmpIdx1 + 2] + *k)); |
|||
// int tmpIdx2 = (ix * sampleCnt_v + iy + 1) * 3; |
|||
// bound = bound.Union(FVec3(evaluatedPoints[tmpIdx2] + *k, evaluatedPoints[tmpIdx2 + 1] + *k, |
|||
// evaluatedPoints[tmpIdx2 + 2] + *k)); |
|||
// int tmpIdx3 = ((ix + 1) * sampleCnt_v + iy + 1) * 3; |
|||
// bound = bound.Union(FVec3(evaluatedPoints[tmpIdx3] + *k, evaluatedPoints[tmpIdx3 + 1] + *k, |
|||
// evaluatedPoints[tmpIdx3 + 2] + *k)); |
|||
int tmpIdx1 = ((ix + 1) * sampleCnt_v + iy) * 3; |
|||
bound = bound.Union(FVec3(evaluatedPoints[tmpIdx1], evaluatedPoints[tmpIdx1 + 1], |
|||
evaluatedPoints[tmpIdx1 + 2])); |
|||
int tmpIdx2 = (ix * sampleCnt_v + iy + 1) * 3; |
|||
bound = bound.Union(FVec3(evaluatedPoints[tmpIdx2], evaluatedPoints[tmpIdx2 + 1], |
|||
evaluatedPoints[tmpIdx2 + 2])); |
|||
int tmpIdx3 = ((ix + 1) * sampleCnt_v + iy + 1) * 3; |
|||
bound = bound.Union(FVec3(evaluatedPoints[tmpIdx3], evaluatedPoints[tmpIdx3 + 1], |
|||
evaluatedPoints[tmpIdx3 + 2])); |
|||
|
|||
BVH[idx].bounds = bound; |
|||
BVH[idx].childNum = -1; |
|||
|
|||
|
|||
} else { |
|||
if (ix % step == 0 || iy % step == 0) { |
|||
// 非叶子节点 |
|||
BVH[idx].u0 = u; |
|||
BVH[idx].u1 = u + step_u * singleStepVal_u; |
|||
BVH[idx].v0 = v; |
|||
BVH[idx].v1 = v + step_v * singleStepVal_v; |
|||
BVH[idx].level = level; |
|||
|
|||
int firstChild = 4 * idx + 1; |
|||
bound = bound.Union(BVH[firstChild].bounds); |
|||
bound = bound.Union(BVH[firstChild + 1].bounds); |
|||
bound = bound.Union(BVH[firstChild + 2].bounds); |
|||
bound = bound.Union(BVH[firstChild + 3].bounds); |
|||
BVH[idx].firstChild = firstChild; |
|||
BVH[idx].bounds = bound; |
|||
} |
|||
} |
|||
|
|||
bias /= 4; |
|||
|
|||
__syncthreads(); |
|||
} |
|||
} |
|||
|
|||
__host__ void BVH::printQuadTree() { |
|||
|
|||
for (int l = 1, levelNodesCnt = 1, baseIdx = 0; l <= maxLevel; l++, levelNodesCnt *= 4, baseIdx = baseIdx * 4 + 1) { |
|||
for (int biasIdx = 0; biasIdx < levelNodesCnt; biasIdx++) { |
|||
int idx = baseIdx + biasIdx; |
|||
auto pMax = nodes[idx].bounds.pMax; |
|||
auto pMin = nodes[idx].bounds.pMin; |
|||
printf("<<%g, %g, %g>,<%g, %g, %g>> ", pMin.x, pMin.y, pMin.z, pMax.x, pMax.y, pMax.z); |
|||
} |
|||
printf("\n =============$$$$$$$============= \n"); |
|||
} |
|||
} |
|||
@ -1,9 +1,9 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/9. |
|||
// |
|||
#include "../../include/device/NurbsCommon.cuh" |
|||
#include "../../include/device/NurbsCurve.cuh" |
|||
#include "../../include/utils.h" |
|||
#include "../../../include/device/Nurbs/nurbs_common.cuh" |
|||
#include "../../../include/device/Nurbs/nurbs_curve.cuh" |
|||
#include "../../../include/utils.h" |
|||
|
|||
__global__ void |
|||
NurbsCurve::g_evaluate(float *res, const float *NTexture, const float *d_points, const int d_pointsCnt, |
|||
@ -0,0 +1,42 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/11. |
|||
// |
|||
|
|||
#include "../../include/device/aabb.cuh" |
|||
|
|||
#define D_DBL_MAX 1.7976931348623158e+308 |
|||
#define D_FLT_MAX 3.402823466e+38F |
|||
|
|||
__device__ __host__ AABB::AABB() { |
|||
float minNum = -D_FLT_MAX; |
|||
float maxNum = D_FLT_MAX; |
|||
pMin = FVec3(minNum, minNum, minNum); |
|||
pMax = FVec3(maxNum, maxNum, maxNum); |
|||
} |
|||
|
|||
__device__ __host__ AABB::AABB(const FVec3 &p) : pMin(p), pMax(p) {} |
|||
|
|||
__device__ __host__ AABB::AABB(const FVec3 &p1, const FVec3 &p2) { |
|||
pMin = FVec3(fmin(p1.x, p2.x), fmin(p1.y, p2.y), fmin(p1.z, p2.z)); |
|||
pMax = FVec3(fmax(p1.x, p2.x), fmax(p1.y, p2.y), fmax(p1.z, p2.z)); |
|||
} |
|||
|
|||
__device__ __host__ AABB AABB::Union(const FVec3 &p) { |
|||
return {FVec3(fmin(pMin.x, p.x), fmin(pMin.y, p.y), fmin(pMin.z, p.z)), |
|||
FVec3(fmax(pMax.x, p.x), fmax(pMax.y, p.y), fmax(pMax.z, p.z))}; |
|||
} |
|||
|
|||
__device__ __host__ AABB AABB::Union(const AABB &b2) { |
|||
return {FVec3(fmin(pMin.x, b2.pMin.x), fmin(pMin.y, b2.pMin.y), fmin(pMin.z, b2.pMin.z)), |
|||
FVec3(fmax(pMax.x, b2.pMax.x), fmax(pMax.y, b2.pMax.y), fmax(pMax.z, b2.pMax.z))}; |
|||
} |
|||
|
|||
__device__ __host__ bool AABB::IsOverlap(const AABB& b2) { |
|||
if (pMin.x > b2.pMax.x) return false; |
|||
if (b2.pMin.x > pMax.x) return false; |
|||
if (pMin.y > b2.pMax.y) return false; |
|||
if (b2.pMin.y > pMax.y) return false; |
|||
if (pMin.z > b2.pMax.z) return false; |
|||
if (b2.pMin.z > pMax.z) return false; |
|||
return true; |
|||
} |
|||
@ -0,0 +1,13 @@ |
|||
// |
|||
// Created by 14727 on 2022/12/11. |
|||
// |
|||
|
|||
#include "../../include/device/vec.cuh" |
|||
|
|||
__device__ __host__ FVec3::FVec3(float x_, float y_, float z_) : x(x_), y(y_), z(z_) {} |
|||
|
|||
__device__ __host__ FVec3::FVec3(const FVec3 &fVec3) : x(static_cast<float>(fVec3.x)), |
|||
y(static_cast<float>(fVec3.y)), |
|||
z(static_cast<float>(fVec3.z)) {} |
|||
|
|||
__device__ __host__ FVec3::FVec3():x(0.), y(0.), z(0.) {} |
|||
Loading…
Reference in new issue