 		float *pixel = (float *)(textureData + y * pitch) + 4 /*RGBA*/ * x;
 
 		// Write out pixel data
-			pixel[0] += pixelColor;
-			pixel[1] += pixelColor;
-			pixel[2] += pixelColor;
-			// Ignore alpha, since it's hardcoded to 1.0f in the display
-			// We have to use a RGBA format since CUDA-DirectX interop doesn't support R32G32B32_FLOAT
+		pixel[0] += pixelColor.x;
+		pixel[1] += pixelColor.y;
+		pixel[2] += pixelColor.z;
+		// Ignore alpha, since it's hardcoded to 1.0f in the display
+		// We have to use a RGBA format since CUDA-DirectX interop doesn't support R32G32B32_FLOAT
 	}
 }

__global__ void PathTraceKernel(unsigned char *textureData, uint width, uint height, size_t pitch, DeviceCamera *camera, Scene::Sphere *spheres, uint numSpheres, uint hashedFrameNumber) {
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	if (x >= width || y >= height) {
		return;
	}

	// Global threadId
	int threadId = (blockIdx.x + blockIdx.y * gridDim.x) * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;


	// Try to intersect with the spheres;
	float closestIntersection = FLT_MAX;
	float3 normal;
	for (uint i = 0; i < numSpheres; ++i) {
		float3 newNormal;
		float intersection = TestRaySphereIntersection(ray, spheres[i], newNormal);
		if (intersection > 0.0f && intersection < closestIntersection) {
			closestIntersection = intersection;
			normal = newNormal;
		}
	}

	float3 pixelColor;
	if (closestIntersection < FLT_MAX) {
		float attentuation = max(dot(normal, make_float3(0.70710678118f, 0.70710678118f, -0.70710678118f)), 0.0f);
		pixelColor = make_float3(0.846, 0.933, 0.949) * attentuation + make_float3(0.15f, 0.15f, 0.15f);
	} else {
		pixelColor = make_float3(0.0f, 0.0f, 0.0f);
	}

	if (x < width && y < height) {
