c++ - 在 CUDA 中计算网格的边界框

Question

我尝试在 CUDA 中为网格编写边界框程序

RT_PROGRAM void bounds (int primIdx, float result[6])
{
  int3 v_idx = index_buffer[primIdx];

  float3 v0 = vertex_buffer[ v_idx.x ];
  float3 v1 = vertex_buffer[ v_idx.y ];
  float3 v2 = vertex_buffer[ v_idx.z ];

  float3 min = fminf( fminf( v0, v1), v2 );  
  float3 max = fmaxf( fmaxf( v0, v1), v2 );

  Aabb* aabb = (Aabb*)result;
  aabb->m_min = min;
  aabb->m_max = max;
}

如果我使用以下行而不是最后两行，则不会发生错误:

aabb->m_min = make_float3(0);
aabb->m_max = make_float3(0);

在主机上，我使用这个初始代码:

Program mesh_bounds = g_ctx->createProgramFromPTXFile(ptx_path, "bounds");
Program mesh_intersect = g_ctx->createProgramFromPTXFile(ptx_path, "intersect2");

// create buffers
Buffer idx_buffer = g_ctx->createBuffer(RT_BUFFER_INPUT, RT_FORMAT_INT3, mesh->nfaces); // mesh->nfaces should always be 1 for triangles
Buffer vtx_buffer = g_ctx->createBuffer(RT_BUFFER_INPUT, RT_FORMAT_FLOAT3, mesh->nvertices);
Buffer nor_buffer = g_ctx->createBuffer(RT_BUFFER_INPUT, RT_FORMAT_FLOAT3, mesh->nfaces); // should only has 1 normal

// load materials
Material mat = g_ctx->createMaterial();

// create the Geometry
    Geometry geo = g_ctx->createGeometry();
    geo->setPrimitiveCount(mesh->nfaces); // mesh->nfaces should be 1

// set both programs to geometry
geo->setBoundingBoxProgram(mesh_bounds);
geo->setIntersectionProgram(mesh_intersect);

// set buffers
geo["vertex_buffer"]->setBuffer(vtx_buffer);
geo["normal_buffer"]->setBuffer(nor_buffer);
geo["index_buffer"]->setBuffer(idx_buffer);

// create Geometry Instance
GeometryInstance inst = g_ctx->createGeometryInstance();
inst->setGeometry(geo);
inst->setMaterialCount(1);
inst->setMaterial(0, mat);

    GeometryGroup grp = g_ctx->createGeometryGroup();
grp->setChildCount(1);
grp->setChild(0, inst);
    grp->setAcceleration(g_ctx->createAcceleration("Bvh", "Bvh"));
    g_meshes.push_back(grp);

启动脚本后，我收到以下错误消息:

***ERROR***: Unknown error (Details: Function "_rtContextLaunch2D" caught exception:
Encountered a CUDA error: driver(). cuEventSynchronize(m_event) returned (999): 
Unknown, [6619195])

有人有想法吗？

Answer 1

OP已经 self 回答了这个问题。不过，最好说 Thrust 已经有了 Bounding Box 构造的示例，尽管仅限于 2D，请参阅 Thrust Bounding Box Example .

下面，我提供了该示例的 3D 扩展，它使用 float3 而不是其中使用的自定义 point2d 结构。

#include <thrust/transform_reduce.h>
#include <thrust/device_vector.h>
#include <thrust/pair.h>
#include <thrust/random.h>
#include <thrust/extrema.h>

/***********************/
/* BOUNDING BOX STRUCT */
/***********************/
struct bbox
{
    float3 lower_left, upper_right;

    // --- Empty box constructor
    __host__ __device__ bbox() {}

    // --- Construct a box from a single point
    __host__ __device__ bbox(const float3 &point) : lower_left(point), upper_right(point) {}

    // --- Construct a box from a pair of points
    __host__ __device__ bbox(const float3 &ll, const float3 &ur) : lower_left(ll), upper_right(ur) {}

};

/*********************************/
/* BOUNDING BOX REDUCTION STRUCT */
/*********************************/
// --- Reduce a pair of bounding boxes (a, b) to a bounding box containing a and b
struct bbox_reduction : public thrust::binary_function<bbox, bbox, bbox>
{
    __host__ __device__ bbox operator()(bbox a, bbox b)
    {
        // --- Lower left corner
        float3 ll = make_float3(thrust::min(a.lower_left.x, b.lower_left.x), thrust::min(a.lower_left.y, b.lower_left.y), thrust::min(a.lower_left.z, b.lower_left.z));

        // --- Upper right corner
        float3 ur = make_float3(thrust::max(a.upper_right.x, b.upper_right.x), thrust::max(a.upper_right.y, b.upper_right.y), thrust::max(a.upper_right.z, b.upper_right.z));

        return bbox(ll, ur);
    }
};

/********/
/* MAIN */
/********/
int main(void)
{
    const size_t N = 40;
    thrust::default_random_engine rng;
    thrust::uniform_real_distribution<float> u01(0.0f, 1.0f);

    // --- Allocate space for 3D points
    thrust::device_vector<float3> d_points(N);

    // --- Generate random 3D points in the unit cube
    for (size_t i = 0; i < N; i++)
    {
        float x = u01(rng);
        float y = u01(rng);
        float z = u01(rng);
        d_points[i] = make_float3(x, y, z);
    }

    // --- The initial bounding box contains the first point of the point cloud
    bbox init = bbox(d_points[0], d_points[0]);

    // --- Binary reduction operation
    bbox_reduction binary_op;

    // --- Compute the bounding box for the point set
    bbox result = thrust::reduce(d_points.begin(), d_points.end(), init, binary_op);

    for (int k = 0; k < N; k++) {
        float3 temp = d_points[k];
        printf("%d %f %f %f\n", k, temp.x, temp.y, temp.z);
    }

    // --- Print output
    std::cout << "bounding box " << std::fixed;
    std::cout << "(" << result.lower_left.x << "," << result.lower_left.y << "," << result.lower_left.z << ") ";
    std::cout << "(" << result.upper_right.x << "," << result.upper_right.y << "," << result.upper_right.z << ")" << std::endl;

    return 0;
}