 	int shapeid;
 };
 
-__device__ bool intersectBVH(
+__device__ bool intersectBVHTriangles(
 	cudaTextureObject_t nodes,
-	BVHCandidate* candidates,
-	int& candidateNum,
 	const Context* ctxt,
 	const aten::ray r,
 	float t_min, float t_max,

}

__device__ bool intersectBVH(
	cudaTextureObject_t nodes,
	const Context* ctxt,
	const aten::ray r,
	float t_min, float t_max,
	aten::Intersection* isect,
	IntersectType type = IntersectType::Closest)
{
	aten::Intersection isectTmp;

	int nodeid = 0;
	float4 node;	// x:left, y:right
	float4 attrib;	// x:shapeid, y:primid, z:nestid

	float4 _boxmin;
	float4 _boxmax;

	aten::vec3 boxmin;
	aten::vec3 boxmax;

	isect->t = t_max;

	for (;;) {
		if (nodeid < 0) {
			break;
		}

		node = tex1Dfetch<float4>(nodes, 4 * nodeid + 0);	// x : hit, y: miss
		attrib = tex1Dfetch<float4>(nodes, 4 * nodeid + 1);	// x : shapeid, y : primgid, z : exid
		_boxmin = tex1Dfetch<float4>(nodes, 4 * nodeid + 2);
		_boxmax = tex1Dfetch<float4>(nodes, 4 * nodeid + 3);

		boxmin = aten::vec3(_boxmin.x, _boxmin.y, _boxmin.z);
		boxmax = aten::vec3(_boxmax.x, _boxmax.y, _boxmax.z);

		bool isHit = false;

		if (attrib.x >= 0 || attrib.y >= 0) {
			const auto& prim = ctxt->prims[(int)attrib.y];

			isectTmp.t = AT_MATH_INF;
			isHit = hitTriangle(&prim, ctxt, r, &isectTmp);
			isectTmp.mtrlid = prim.mtrlid;

			if (isectTmp.t < isect->t) {
				*isect = isectTmp;
				isect->objid = (int)attrib.x;
				isect->primid = (int)attrib.y;

				if (type == IntersectType::Closer) {
					return true;
				}
			}
		}
		else {
			isHit = aten::aabb::hit(r, boxmin, boxmax, t_min, t_max);
		}

		if (isHit) {
			nodeid = (int)node.x;
		}
		else {
			nodeid = (int)node.y;
		}
	}

	return (isect->objid >= 0);
}


__device__ bool intersectBVH(
	const Context* ctxt,
	const aten::ray& r,
	aten::Intersection* isect)

	float t_min = AT_MATH_EPSILON;
	float t_max = AT_MATH_INF;

	bool isHit = intersectBVH(
		ctxt->nodes[0],
		ctxt,
		r,
		t_min, t_max,
		isect);

	return isHit;
}

