 			if (isectTmp.t < isect->t) {
 				*isect = isectTmp;
 				isect->objid = (int)attrib.x;
-
-				if (type == IntersectType::Closer) {
-					return true;
-				}
 			}
 		}
 		else {

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

				if (type == IntersectType::Closer) {
					return true;
				}
			}
		}
		else {

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

AT_CUDA_INLINE __device__ bool intersectBVHCloser(
	cudaTextureObject_t nodes,
	const Context* ctxt,
	const aten::ray r,
	float t_min, float t_max,
	aten::Intersection* isect)
{
	aten::Intersection isectTmp;

	real hitt = t_max;
	isect->t = t_max;

	int tmpexid = -1;

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

	float4 aabb[2];

	real t = AT_MATH_INF;

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

		auto boxmin = aten::vec3(aabb[0].x, aabb[0].y, aabb[0].z);
		auto boxmax = aten::vec3(aabb[1].x, aabb[1].y, aabb[1].z);

		bool isHit = false;

		if (attrib.x >= 0) {
			// Leaf.
			tmpexid = -1;

			const auto* s = &ctxt->shapes[(int)attrib.x];

			if (attrib.z >= 0) {	// exid
				aten::ray transformedRay;

				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;
				}

				isHit = intersectBVHCloserTriangles(ctxt->nodes[(int)attrib.z], ctxt, transformedRay, t_min, t_max, &isectTmp);
			}
			else {
				// TODO
				// Only sphere...
				//isHit = intersectShape(s, nullptr, ctxt, r, t_min, t_max, &recTmp, &recOptTmp);
				isectTmp.t = AT_MATH_INF;
				isHit = hitSphere(s, r, t_min, t_max, &isectTmp);
				isectTmp.mtrlid = s->mtrl.idx;
			}

			if (isectTmp.t < isect->t) {
				*isect = isectTmp;
				isect->objid = (int)attrib.x;
				return true;
			}
		}
		else {
			isHit = aten::aabb::hit(r, boxmin, boxmax, t_min, t_max, &t);
		}

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

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

AT_CUDA_INLINE __device__ bool intersectBVH(
	const Context* ctxt,
	const aten::ray& r,
