 {
 	const auto ix = blockIdx.x * blockDim.x + threadIdx.x;
 	const auto iy = blockIdx.y * blockDim.y + threadIdx.y;
-#else
-__host__ void raytracing(
-	int ix, int iy,
-	float4* p,
-	int width, int height,
-	Camera* camera,
-	Sphere* spheres, int num)
-{
-#endif
 
 	if (ix >= width && iy >= height) {
 		return;

	return isHit;
}

#if 1
__global__ void raytracing(
	float4* p,
	int width, int height,

	CameraSampleResult camsample;
	sampleCamera(&camsample, camera, s, t);

	aten::vec3 contrib(0);
	aten::vec3 throughput(1);

	int depth = 0;

	while (depth < 5) {
		aten::hitrecord rec;

		if (intersect(&camsample.r, &rec, &ctx)) {
			const aten::MaterialParameter& m = ctx.mtrls[rec.mtrlid];

			if (m.type.isEmissive) {
				auto emit = m.baseColor;
				contrib = throughput * emit;
				break;
			}
		}
		else {
			break;
		}

		depth++;
	}

	p[idx] = make_float4(contrib.x, contrib.y, contrib.z, 1);
}

static Sphere g_spheres[] = {
