 	return (isect->objid >= 0);
 }
 
+#define ENABLE_PLANE_LOOP_SBVH
+
 template <idaten::IntersectType Type>
 AT_CUDA_INLINE __device__ bool intersectSBVH(
 	const Context* ctxt,

						transformedRay = r;
					}

					node = ctxt->nodes[(int)attrib.z];

					objid = (int)attrib.x;
					meshid = (int)attrib.w;

					toplayerHit = (int)node0.w;
					toplayerMiss = (int)node1.w;

					nodeid = 0;

					continue;
				}
			}
			else if (attrib.y >= 0) {
				int primidx = (int)attrib.y;
				aten::PrimitiveParamter prim;
				prim.v0 = ((aten::vec4*)ctxt->prims)[primidx * aten::PrimitiveParamter_float4_size + 0];
				prim.v1 = ((aten::vec4*)ctxt->prims)[primidx * aten::PrimitiveParamter_float4_size + 1];

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

				bool isIntersect = (Type == idaten::IntersectType::Any
					? isHit
					: isectTmp.t < isect->t);

				if (isIntersect) {
					*isect = isectTmp;
					isect->objid = objid;
					isect->primid = primidx;
					isect->mtrlid = prim.mtrlid;

					isect->meshid = prim.gemoid;
					isect->meshid = (isect->meshid < 0 ? meshid : isect->meshid);

					t_max = isect->t;

					if (Type == idaten::IntersectType::Closer
						|| Type == idaten::IntersectType::Any)
					{
						return true;
					}
				}
			}
		}
		else {
			//isHit = aten::aabb::hit(r, boxmin, boxmax, t_min, t_max, &t);
			isHit = hitAABB(transformedRay.org, transformedRay.dir, boxmin, boxmax, t_min, t_max, &t);
		}

		if (isHit) {


		bool isHit = false;

#ifdef ENABLE_PLANE_LOOP_SBVH
		if (attrib.x >= 0) {
			// Leaf.
			const auto* s = &ctxt->shapes[(int)attrib.x];

			if (attrib.z >= 0) {	// exid
				//if (hitAABB(r.org, r.dir, boxmin, boxmax, t_min, t_max, &t))
				{
					if (s->mtxid >= 0) {
						auto mtxW2L = ctxt->matrices[s->mtxid * 2 + 1];
						transformedRay.dir = mtxW2L.applyXYZ(r.dir);
						transformedRay.dir = normalize(transformedRay.dir);
						transformedRay.org = mtxW2L.apply(r.org) + AT_MATH_EPSILON * transformedRay.dir;
					}
					else {
						transformedRay = r;
					}

					node = ctxt->nodes[(int)attrib.z];

					objid = (int)attrib.x;
					meshid = (int)attrib.w;

					toplayerHit = (int)node0.w;
					toplayerMiss = (int)node1.w;

					nodeid = 0;

					continue;
				}
			}
			else if (attrib.y >= 0) {
				int primidx = (int)attrib.y;
				aten::PrimitiveParamter prim;
				prim.v0 = ((aten::vec4*)ctxt->prims)[primidx * aten::PrimitiveParamter_float4_size + 0];
				prim.v1 = ((aten::vec4*)ctxt->prims)[primidx * aten::PrimitiveParamter_float4_size + 1];

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

				bool isIntersect = (Type == idaten::IntersectType::Any
					? isHit
					: isectTmp.t < isect->t);

				if (isIntersect) {
					*isect = isectTmp;
					isect->objid = objid;
					isect->primid = primidx;
					isect->mtrlid = prim.mtrlid;

					isect->meshid = prim.gemoid;
					isect->meshid = (isect->meshid < 0 ? meshid : isect->meshid);

					t_max = isect->t;

					if (Type == idaten::IntersectType::Closer
						|| Type == idaten::IntersectType::Any)
					{
						return true;
					}
				}
			}
		}
		else {
			//isHit = aten::aabb::hit(r, boxmin, boxmax, t_min, t_max, &t);
			isHit = hitAABB(transformedRay.org, transformedRay.dir, boxmin, boxmax, t_min, t_max, &t);
		}

		if (isHit) {
			nodeid = (int)node0.w;
		}
		else {
			nodeid = (int)node1.w;
		}

		if (nodeid < 0 && toplayerHit >= 0) {
			nodeid = isHit ? toplayerHit : toplayerMiss;
			toplayerHit = -1;
			toplayerMiss = -1;
			node = ctxt->nodes[0];
			transformedRay = r;
		}
#else
		if (attrib.x >= 0) {
			// Leaf.
			const auto* s = &ctxt->shapes[(int)attrib.x];

			if (attrib.z >= 0) {	// exid
									//if (aten::aabb::hit(r, boxmin, boxmax, t_min, t_max, &t)) {
				if (hitAABB(r.org, r.dir, boxmin, boxmax, t_min, t_max, &t)) {
					if (s->mtxid >= 0) {
						auto mtxW2L = ctxt->matrices[s->mtxid * 2 + 1];
