 // https://github.com/leonardo-domingues/atrbvh
 // http://research.nvidia.com/sites/default/files/publications/karras2012hpg_paper.pdf
 
-__device__  int computeLongestCommonPrefix(
+__forceinline__ __device__ __host__ int computeLongestCommonPrefix(
 	const uint32_t* sortedKeys,
 	uint32_t numOfElems,
 	int index1, int index2,

	const auto lcp1 = computeLongestCommonPrefix(sortedKeys, numOfElems, i, i + 1, key1);
	const auto lcp2 = computeLongestCommonPrefix(sortedKeys, numOfElems, i, i - 1, key1);

	auto d = (lcp1 - lcp2) < 0 ? -1 : 1;

	// Compute upper bound for the length of the range
	const auto minLcp = computeLongestCommonPrefix(sortedKeys, numOfElems, i, i - d, key1);
	int lMax = 2;
	while (computeLongestCommonPrefix(sortedKeys, numOfElems, i, i + lMax * d, key1) > minLcp)


		// Compute bouding box.
		{
			dim3 block(128, 1, 1);
			dim3 grid((numberOfTris + block.x - 1) / block.x, 1, 1);

