diff --git a/src/main.cu b/src/main.cu index 2fb4f75..15da3a1 100644 --- a/src/main.cu +++ b/src/main.cu @@ -56,13 +56,19 @@ __device__ vec3 color(const ray& r, hitable **world, curandState *local_rand_sta return vec3(0, 0, 0); } +__global__ void rand_init(curandState *rand_state) { + if (threadIdx.x == 0 && threadIdx.y == 0) { + curand_init(2002, 0, 0, rand_state); + } +} + __global__ void render_init(int max_x, int max_y, curandState *rand_state) { int x = threadIdx.x + blockIdx.x * blockDim.x; int y = threadIdx.y + blockIdx.y * blockDim.y; if ((x >= max_x) || (y >= max_y)) return; int pixel_idx = y*max_x + x; - curand_init(2002, pixel_idx, 0, &rand_state[pixel_idx]); + curand_init(2002+pixel_idx, 0, 0, &rand_state[pixel_idx]); } __global__ void render(vec3 *fb, @@ -93,31 +99,59 @@ __global__ void render(vec3 *fb, fb[pixel_idx] = col; } -__global__ void create_world(hitable **d_list, int d_list_size, hitable **d_world, camera **d_camera, int nx, int ny) { - if (threadIdx.x == 0 && blockIdx.y == 0) { - d_list[0] = new sphere(vec3( 0, 0 , -1), 0.5, new lambertian(vec3(0.1, 0.2, 0.5))); - d_list[1] = new sphere(vec3( 0, -100.5, -1), 100, new lambertian(vec3(0.8, 0.8, 0.0))); - d_list[2] = new sphere(vec3( 1, 0 , -1), 0.5, new metal(vec3(0.8, 0.6, 0.2), 0.0)); - d_list[3] = new sphere(vec3(-1, 0 , -1), 0.5, new dielectric(1.5)); - d_list[4] = new sphere(vec3(-1, 0 , -1), -0.45, new dielectric(1.5)); - *d_world = new hitable_list(d_list, d_list_size); +__global__ void create_world(hitable **d_list, int d_list_size, hitable **d_world, camera **d_camera, int nx, int ny, curandState *rand_state) { + if (!(threadIdx.x == 0 && blockIdx.y == 0)) return; - vec3 lookfrom(3,3,2); - vec3 lookat(0,0,-1); - float dist_to_focus = (lookfrom-lookat).length(); - float aperture = 2.0f; - *d_camera = new camera(lookfrom, lookat, - vec3(0,1,0), - 20.0, - float(nx)/float(ny), - aperture, - dist_to_focus); + curandState local_rand_state = *rand_state; + d_list[0] = new sphere(vec3(0, -1000, -1), 1000, new lambertian(vec3(0.5, 0.5, 0.5))); + + #define RND() curand_uniform(&local_rand_state) + + int idx = 1; + for (int a = -11; a < 11; a++) { + for (int b = -11; b < 11; b++) { + float choose_mat = RND(); + vec3 center(a + RND(), 0.2, b + RND()); + material *mat; + if (choose_mat < 0.8f) { + mat = new lambertian(vec3(RND()*RND(), RND()*RND(), RND()*RND())); + } else if (choose_mat < 0.95f) { + mat = new metal(vec3(0.5f*(1.0f+RND()), 0.5f*(1.0f+RND()), 0.5f*(1.0f+RND())), 0.5f*RND()); + } else { + mat = new dielectric(1.5); + } + d_list[idx++] = new sphere(center, 0.2, mat); + } } + + #undef RND + + d_list[idx++] = new sphere(vec3(0, 1,0), 1.0, new dielectric(1.5)); + d_list[idx++] = new sphere(vec3(-4, 1, 0), 1.0, new lambertian(vec3(0.4, 0.2, 0.1))); + d_list[idx++] = new sphere(vec3(4, 1, 0), 1.0, new metal(vec3(0.7, 0.6, 0.5), 0.0)); + *rand_state = local_rand_state; + + assert(d_list_size == idx); + *d_world = new hitable_list(d_list, d_list_size); + + vec3 lookfrom(13,2,3); + vec3 lookat(0,0,0); + float dist_to_focus = 10.0; // (lookfrom-lookat).length(); + float aperture = 0.1f; + *d_camera = new camera( + lookfrom, + lookat, + vec3(0,1,0), + 30.0, + float(nx)/float(ny), + aperture, + dist_to_focus + ); } __global__ void free_world(hitable **d_list, int d_list_size, hitable **d_world, camera **d_camera) { for (int i = 0; i < d_list_size; i++) { - delete &((sphere*)d_list[i])->mat; + delete ((sphere*)d_list[i])->mat; delete &d_list[i]; } delete &d_world; @@ -128,7 +162,7 @@ int main() { const char *image_filename = "out.ppm"; int nx = 1200; int ny = 600; - int ns = 100; + int ns = 20; int tx = 16; int ty = 16; @@ -145,16 +179,23 @@ int main() { // allocate random state curandState *d_rand_state; checkCudaErrors(cudaMalloc((void **)&d_rand_state, num_pixels*sizeof(curandState))); + curandState *d_rand_world_state; + checkCudaErrors(cudaMalloc((void **)&d_rand_world_state, num_pixels*sizeof(curandState))); + + // Initialize world rand state + rand_init<<<1,1>>>(d_rand_world_state); + checkCudaErrors(cudaGetLastError()); + checkCudaErrors(cudaDeviceSynchronize()); // populate world hitable **d_list; - int d_list_size = 5; + int d_list_size = 22*22 + 1 + 3; checkCudaErrors(cudaMalloc((void **)&d_list, d_list_size*sizeof(hitable *))); hitable **d_world; checkCudaErrors(cudaMalloc((void **)&d_world, sizeof(hitable *))); camera **d_camera; checkCudaErrors(cudaMalloc((void **)&d_camera, sizeof(camera *))); - create_world<<<1,1>>>(d_list, d_list_size, d_world, d_camera, nx, ny); + create_world<<<1,1>>>(d_list, d_list_size, d_world, d_camera, nx, ny, d_rand_world_state); checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize()); @@ -206,6 +247,7 @@ int main() { checkCudaErrors(cudaFree(d_world)); checkCudaErrors(cudaFree(d_camera)); checkCudaErrors(cudaFree(d_rand_state)); + checkCudaErrors(cudaFree(d_rand_world_state)); checkCudaErrors(cudaFree(fb)); cudaDeviceReset();