 {
 	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(
+	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;
+
 	real hitt = t_max;
 	isect->t = t_max;
 

	int shapeid;
};

__device__ bool intersectBVH(
	cudaTextureObject_t nodes,
	BVHCandidate* candidates,
	int& candidateNum,
	const Context* ctxt,
	const aten::ray r,
	float t_min, float t_max,

}

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