 //#pragma optimize( "", off)
 
 __global__ void computeSkinning(
-	uint32_t indexNum,
+	uint32_t vtxNum,
 	const aten::SkinningVertex* __restrict__ vertices,
-	const uint32_t* __restrict__ indices,
 	const aten::mat4* __restrict__ matrices,
 	aten::vec4* dstPos,
 	aten::vec4* dstNml)
 {
 	const auto idx = blockIdx.x * blockDim.x + threadIdx.x;
 
-	if (idx >= indexNum) {
+	if (idx >= vtxNum) {
 		return;
 	}
 
-	const auto vtxIdx = indices[idx];
-	const auto* vtx = &vertices[vtxIdx];
+	const auto* vtx = &vertices[idx];
 
 	aten::vec4 srcPos = vtx->position;
 	aten::vec4 srcNml = aten::vec4(vtx->normal, 0);

}

__global__ void computeSkinningWithTriangles(
	uint32_t triNum,
	const aten::SkinningVertex* __restrict__ vertices,
	aten::PrimitiveParamter* triangles,
	const aten::mat4* __restrict__ matrices,
	int indexOffset,
	aten::vec4* dstPos,
	aten::vec4* dstNml)
{
	const auto idx = blockIdx.x * blockDim.x + threadIdx.x;

	if (idx >= triNum) {
		return;
	}

	auto* tri = &triangles[idx];

#pragma unroll
	for (int t = 0; t < 3; t++) {
		const auto vtxIdx = tri->idx[t];
		const auto* vtx = &vertices[vtxIdx];

		aten::vec4 srcPos = vtx->position;
		aten::vec4 srcNml = aten::vec4(vtx->normal, 0);

		aten::vec4 resultPos(0);
		aten::vec4 resultNml(0);

		for (int i = 0; i < 4; i++) {
			int idx = int(vtx->blendIndex[i]);
			float weight = vtx->blendWeight[i];

			aten::mat4 mtx = matrices[idx];

			resultPos += weight * mtx * vtx->position;
			resultNml += weight * mtx * srcNml;
		}

		resultNml = normalize(resultNml);

		dstPos[vtxIdx] = aten::vec4(resultPos.x, resultPos.y, resultPos.z, vtx->uv[0]);
		dstNml[vtxIdx] = aten::vec4(resultNml.x, resultNml.y, resultNml.z, vtx->uv[1]);
	}

	{
		const auto& v0 = dstPos[tri->idx[0]];
		const auto& v1 = dstPos[tri->idx[1]];
		const auto& v2 = dstPos[tri->idx[2]];

		auto a = v1 - v0;
		auto b = v2 - v0;
		
		tri->area = cross(a, b).length();

		tri->idx[0] += indexOffset;
		tri->idx[1] += indexOffset;
		tri->idx[2] += indexOffset;
	}
}

// NOTE

		{
			auto willComputeWithTriangles = m_triangles.num() > 0;

			const auto vtxNum = m_vertices.num();

			dim3 block(512);
			dim3 grid((vtxNum + block.x - 1) / block.x);

			if (willComputeWithTriangles) {
				computeSkinningWithTriangles << <grid, block >> > (
					vtxNum,
					m_vertices.ptr(),
					m_matrices.ptr(),
					dstPos, dstNml);

				checkCudaKernel(computeSkinningWithTriangles);

				const auto triNum = m_triangles.num();

				grid = dim3((triNum + block.x - 1) / block.x);

				setTriangleParam << <grid, block >> > (
					triNum,
					m_triangles.ptr(),
					indexOffset,
					dstPos);

				checkCudaKernel(setTriangleParam);
			}
			else {
				computeSkinning << <grid, block >> > (
					vtxNum,
					m_vertices.ptr(),
					m_matrices.ptr(),
					dstPos, dstNml);

