Skip to content

Commit

Permalink
basic material works
Browse files Browse the repository at this point in the history
  • Loading branch information
@netwarm007
netwarm007 committed Jan 4, 2023
1 parent 6c21df3 commit 963e61f
Show file tree
Hide file tree
Showing 5 changed files with 300 additions and 80 deletions.
157 changes: 130 additions & 27 deletions Asset/Shaders/CUDA/OptixTest.shader.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,17 +19,35 @@ static __forceinline__ __device__ float3 _f(vec3 v) {
return make_float3(v[0], v[1], v[2]);
}

struct MyTracePayload {
vec3 attenuation;
vec3 scatter_ray_origin;
vec3 scatter_ray_direction;
int max_depth;
bool done;
};

static __forceinline__ __device__ void trace(
OptixTraversableHandle handle,
vec3 ray_origin,
vec3 ray_direction,
float tmin,
float tmax,
vec3& prd ) {
unsigned int p0, p1, p2;
p0 = __float_as_uint(prd[0]);
p1 = __float_as_uint(prd[1]);
p0 = __float_as_uint(prd[2]);
MyTracePayload& prd ) {
unsigned int p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10;
#if 0 // we actually do not need pass-in anything, the payload is used only as a return value.
p0 = __float_as_uint(prd.attenuation[0]);
p1 = __float_as_uint(prd.attenuation[1]);
p2 = __float_as_uint(prd.attenuation[2]);
p3 = __float_as_uint(prd.scatter_ray_origin[0]);
p4 = __float_as_uint(prd.scatter_ray_origin[1]);
p5 = __float_as_uint(prd.scatter_ray_origin[2]);
p6 = __float_as_uint(prd.scatter_ray_direction[0]);
p7 = __float_as_uint(prd.scatter_ray_direction[1]);
p8 = __float_as_uint(prd.scatter_ray_direction[2]);
p9 = prd.max_depth;
p10 = (prd.done) ? 1 : 0;
#endif
optixTrace(
handle,
_f(ray_origin),
Expand All @@ -42,56 +60,102 @@ static __forceinline__ __device__ void trace(
0, // SBT offset
0, // SBT stride
0, // missSBTIndex
p0, p1, p2 );
prd[0] = __uint_as_float(p0);
prd[1] = __uint_as_float(p1);
prd[2] = __uint_as_float(p2);
p0, p1, p2,
p3, p4, p5,
p6, p7, p8,
p9, p10 );
prd.attenuation[0] = __uint_as_float(p0);
prd.attenuation[1] = __uint_as_float(p1);
prd.attenuation[2] = __uint_as_float(p2);
prd.scatter_ray_origin[0] = __uint_as_float(p3);
prd.scatter_ray_origin[1] = __uint_as_float(p4);
prd.scatter_ray_origin[2] = __uint_as_float(p5);
prd.scatter_ray_direction[0] = __uint_as_float(p6);
prd.scatter_ray_direction[1] = __uint_as_float(p7);
prd.scatter_ray_direction[2] = __uint_as_float(p8);
prd.max_depth = p9;
prd.done = (p10 == 1) ? true : false;
}


static __forceinline__ __device__ void setPayload( vec3 p )
static __forceinline__ __device__ void setPayload( const vec3& attenuation, bool is_scattered)
{
optixSetPayload_0( __float_as_uint( p[0] ) );
optixSetPayload_1( __float_as_uint( p[1] ) );
optixSetPayload_2( __float_as_uint( p[2] ) );
optixSetPayload_0( __float_as_uint( attenuation[0] ) );
optixSetPayload_1( __float_as_uint( attenuation[1] ) );
optixSetPayload_2( __float_as_uint( attenuation[2] ) );
optixSetPayload_10( is_scattered ? 0 : 1 );
}

static __forceinline__ __device__ void setPayload( const vec3& attenuation, const ray& ray_scattered, bool is_scattered, int max_depth )
{
auto orig = ray_scattered.getOrigin();
auto direct = ray_scattered.getDirection();
optixSetPayload_0( __float_as_uint( attenuation[0] ) );
optixSetPayload_1( __float_as_uint( attenuation[1] ) );
optixSetPayload_2( __float_as_uint( attenuation[2] ) );
optixSetPayload_3( __float_as_uint( orig[0] ) );
optixSetPayload_4( __float_as_uint( orig[1] ) );
optixSetPayload_5( __float_as_uint( orig[2] ) );
optixSetPayload_6( __float_as_uint( direct[0] ) );
optixSetPayload_7( __float_as_uint( direct[1] ) );
optixSetPayload_8( __float_as_uint( direct[2] ) );
optixSetPayload_9( max_depth );
optixSetPayload_10( is_scattered ? 0 : 1 );
}

static __forceinline__ __device__ vec3 getPayload()
static __forceinline__ __device__ MyTracePayload getPayload()
{
return vec3({
return MyTracePayload({
{
__uint_as_float( optixGetPayload_0() ),
__uint_as_float( optixGetPayload_1() ),
__uint_as_float( optixGetPayload_2() )
},
{
__uint_as_float( optixGetPayload_3() ),
__uint_as_float( optixGetPayload_4() ),
__uint_as_float( optixGetPayload_5() )
},
{
__uint_as_float( optixGetPayload_6() ),
__uint_as_float( optixGetPayload_7() ),
__uint_as_float( optixGetPayload_8() )
},
static_cast<int>(optixGetPayload_9()),
optixGetPayload_10() ? true : false
});
}

extern "C"
__global__ void __raygen__rg() {
uint3 launch_index = optixGetLaunchIndex();
RayGenData* rtData = (RayGenData*)optixGetSbtDataPointer();
// RayGenData* rtData = (RayGenData*)optixGetSbtDataPointer();

unsigned int i = launch_index.x;
unsigned int j = launch_index.y;
unsigned int pixel_index = j * params.image->Width + i;

curandStateMRG32k3a* local_rand_state = &params.rand_state[pixel_index];

int num_of_samples = rtData->num_of_samples;
int num_of_samples = params.num_of_samples;
vec3 col = {0.f, 0.f, 0.f};

for (int s = 0; s < num_of_samples; s++) {
float u = float(i + curand_uniform(local_rand_state)) / params.image->Width;
float v = float(j + curand_uniform(local_rand_state)) / params.image->Height;
ray r = params.cam->get_ray(u, v, local_rand_state);

vec3 attenuation = {1.0f, 1.0f, 1.0f};
trace( params.handle,
r.getOrigin(),
r.getDirection(),
0.00f, // tmin
FLT_MAX, // tmax
attenuation);
vec3 attenuation = {1.f, 1.f, 1.f};
MyTracePayload payload(attenuation, r.getOrigin(), r.getDirection(), true, params.max_depth);
do {
trace( params.handle,
payload.scatter_ray_origin,
payload.scatter_ray_direction,
0.001f, // tmin
FLT_MAX, // tmax
payload);

attenuation *= payload.attenuation;
} while (!payload.done);

col += attenuation;
}
Expand All @@ -110,12 +174,24 @@ __global__ void __miss__ms() {
float t = 0.5f * (unit_direction[1] + 1.0f);
vec3 c = (1.0f - t) * color({0, 0, 0}) + t * (missData->bg_color);

setPayload(c);
setPayload(c, false);
}

extern "C"
__global__ void __closesthit__ch() {
auto payload = getPayload();
if(payload.max_depth < 0) {
setPayload({0.f, 0.f, 0.f}, false);
}

uint3 launch_index = optixGetLaunchIndex();

float t_hit = optixGetRayTmax();
unsigned int i = launch_index.x;
unsigned int j = launch_index.y;
unsigned int pixel_index = j * params.image->Width + i;

curandStateMRG32k3a* local_rand_state = &params.rand_state[pixel_index];

const vec3 ray_orig = _V(optixGetWorldRayOrigin());
const vec3 ray_dir = _V(optixGetWorldRayDirection());
Expand All @@ -134,5 +210,32 @@ __global__ void __closesthit__ch() {
vec3 world_normal = _V(optixTransformNormalFromObjectToWorldSpace(_f(obj_normal)));
My::Normalize(world_normal);

setPayload(world_normal * 0.5f + 0.5f);
ray scattered;
ray r_in(ray_orig, ray_dir);
hit_record rec;
rec.set(t_hit, world_raypos, world_normal, My::DotProduct(ray_dir, world_normal) < 0, nullptr);
color attenuation;
HitGroupData* hg_data = reinterpret_cast<HitGroupData*>(optixGetSbtDataPointer());
bool b_scattered;
switch(hg_data->material_type) {
case Material::MAT_DIFFUSE:
{
b_scattered = lambertian::scatter_static(r_in, rec, attenuation, scattered, local_rand_state, hg_data->base_color);
break;
}
case Material::MAT_METAL:
{
b_scattered = metal::scatter_static(r_in, rec, attenuation, scattered, local_rand_state, hg_data->base_color, hg_data->fuzz);
break;
}
case Material::MAT_DIELECTRIC:
{
b_scattered = dielectric::scatter_static(r_in, rec, attenuation, scattered, local_rand_state, hg_data->ir);
break;
}
default: setPayload({0.8f, 0.3f, 0.9f}, r_in, false, payload.max_depth - 1);
return;
}

setPayload(attenuation, scattered, b_scattered, payload.max_depth - 1);
}
11 changes: 9 additions & 2 deletions Framework/GeomMath/geommath.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -367,6 +367,13 @@ __host__ __device__ Vector<T, N> operator*(const Vector<T, N>& vec1,
return result;
}

template <class T, Dimension auto N>
__host__ __device__ Vector<T, N>& operator*=(Vector<T, N>& vec1,
const Vector<T, N>& vec2) {
vec1 = vec1 * vec2;
return vec1;
}

template <class T, Dimension auto N>
__host__ __device__ inline void DivByElement(Vector<T, N>& result,
const Vector<T, N>& a,
Expand Down Expand Up @@ -529,7 +536,7 @@ __host__ __device__ inline void Normalize(Vector<T, N>& a) {
template <class T, Dimension auto N>
__host__ __device__ inline bool isNearZero(const Vector<T, N>& vec) {
bool result = true;
const auto s = 1e-8;
const auto s = (T)1e-8;

for (Dimension auto i = 0; i < N; i++) {
if (fabs(vec[i]) >= s) {
Expand All @@ -553,7 +560,7 @@ __host__ __device__ inline Vector<T, N> Refract(const Vector<T, N>& v, const Vec

Vector<T, N> r_out_perp = etai_over_etat * (v + cos_theta * n);
Vector<T, N> r_out_parallel =
-(T)sqrt(fabs(1.0 - LengthSquared(r_out_perp))) * n;
-(T)sqrt(fabs((T)1.0 - LengthSquared(r_out_perp))) * n;
return r_out_perp + r_out_parallel;
}

Expand Down
Loading

0 comments on commit 963e61f

Please sign in to comment.