 	idaten::PathTracing::Path* paths,
 	int* hitindices,
 	int hitnum,
-	idaten::PathTracing::ShadowRay* shadowRays,
-	aten::ShapeParameter* shapes, int geomnum,
+	const idaten::PathTracing::ShadowRay* __restrict__ shadowRays,
+	const aten::ShapeParameter* __restrict__ shapes, int geomnum,
 	aten::MaterialParameter* mtrls,
-	aten::LightParameter* lights, int lightnum,
+	const aten::LightParameter* __restrict__ lights, int lightnum,
 	cudaTextureObject_t* nodes,
-	aten::PrimitiveParamter* prims,
+	const aten::PrimitiveParamter* __restrict__ prims,
 	cudaTextureObject_t vtxPos,
-	aten::mat4* matrices)
+	const aten::mat4* __restrict__ matrices)
 {
 	int idx = blockIdx.x * blockDim.x + threadIdx.x;
 

				texVtxPos, texVtxNml,
				m_mtxparams.ptr(),
				m_tex.ptr(),
				m_random.ptr());
		}
		else {
			shade<false> << <blockPerGrid, threadPerBlock >> > (


//#define ENABLE_SEPARATE_ALBEDO
//#define ENABLE_PROGRESSIVE

#define ENABLE_PERSISTENT_THREAD
#define SEPARATE_SHADOWRAY_HITTEST

__global__ void genPath(
	idaten::PathTracing::Path* paths,
