レイトレ合宿メモ
レンダラー名 遍照 Henjou Renderer
仏教用語当たりを使いたい感じがあった
Multi-Scattering的な要素も欲しい
「遍照」あまねく照らすこと
なんかGI的な要素を感じていい感じなので採用、多分HengJohが正しい英語名だけど長いのでHenjouにする(仏教用語ってタブー的なやつありそうだけどそうじゃないことを祈る)
今回の目的
MultiScatteringを実装する
後はTransmission含めたGGX周りをやる
出来たらNon-Height feildをやりたい感じ
""""""実行できるようにする""""""""
論文
前回起動出来なかった理由
→多分ptxファイルの読み込みがsampleのやつそのまま使ってたせいで相対パスになってたんじゃないかなぁって気がする。
開発記
8/17
https://scrapbox.io/files/64df61816d74f70021e0a90d.png
表示までやった。
なんか知らないけどcuファイルがBuildフォルダではなく、cmakeする前のSDK sampleフォルダにあるせいで前のレンダラーのcuファイルがどっか行ってしまっていたので1から書き直すことにした
割と忘れているのでここで書いておく
code:cu
extern "C" {
__constant__ Params params;
}
struct Payload
{
float3 raypos;
float3 raydir;
bool is_hit = false;
float3 position;
float3 normal;
float3 vert_color;
float2 texcoord;
int material_id;
float3 basecolor;
float metallic;
float roughness;
float sheen;
float clearcoat;
float clearcoat_roughness;
float ior;
float transmission;
float3 emission;
bool is_light = false;
int primitive_id;
int instance_id;
};
static __forceinline__ __device__ void* unpack_ptr(unsigned int i0,
unsigned int i1)
{
const unsigned long long uptr =
static_cast<unsigned long long>(i0) << 32 | i1;
void* ptr = reinterpret_cast<void*>(uptr);
return ptr;
}
static __forceinline__ __device__ void pack_ptr(void* ptr, unsigned int& i0,
unsigned int& i1)
{
const unsigned long long uptr = reinterpret_cast<unsigned long long>(ptr);
i0 = uptr >> 32;
i1 = uptr & 0x00000000ffffffff;
}
static __forceinline__ __device__ Payload* get_payload_ptr()
{
const unsigned int u0 = optixGetPayload_0();
const unsigned int u1 = optixGetPayload_1();
return reinterpret_cast<Payload*>(unpack_ptr(u0, u1));
}
static __forceinline__ __device__ void setPayload(float3 p)
{
optixSetPayload_0(__float_as_uint(p.x));
optixSetPayload_1(__float_as_uint(p.y));
optixSetPayload_2(__float_as_uint(p.z));
}
static __forceinline__ __device__ void computeRay(uint3 idx, uint3 dim, float3& origin, float3& direction)
{
const float3 U = params.cam_u;
const float3 V = params.cam_v;
const float3 W = params.cam_w;
const float2 d = 2.0f * make_float2(
static_cast<float>(idx.x) / static_cast<float>(dim.x),
static_cast<float>(idx.y) / static_cast<float>(dim.y)
) - 1.0f;
origin = params.cam_eye;
direction = normalize(d.x * U + d.y * V + W);
}
static __forceinline__ __device__ void TraceOcculution(
OptixTraversableHandle handle,
float3 ray_origin,
float3 ray_direction,
float tmin,
float tmax,
Payload* prd
) {
unsigned int u0, u1;
pack_ptr(prd, u0, u1);
optixTrace(
params.traversal_handle,
ray_origin,
ray_direction,
tmin, // Min intersection distance
tmax, // Max intersection distance
0.0f, // rayTime -- used for motion blur
OptixVisibilityMask(255), // Specify always visible
OPTIX_RAY_FLAG_NONE,
1, // SBT offset -- See SBT discussion
2, // SBT stride -- See SBT discussion
0, // missSBTIndex -- See SBT discussion
u0,
u1
);
}
static __forceinline__ __device__ void RayTrace(
OptixTraversableHandle handle,
float3 ray_origin,
float3 ray_direction,
float tmin,
float tmax,
Payload* prd
) {
unsigned int u0, u1;
pack_ptr(prd, u0, u1);
optixTrace(
params.traversal_handle,
ray_origin,
ray_direction,
tmin, // Min intersection distance
tmax, // Max intersection distance
0.0f, // rayTime -- used for motion blur
OptixVisibilityMask(255), // Specify always visible
OPTIX_RAY_FLAG_NONE,
0, // SBT offset -- See SBT discussion
2, // SBT stride -- See SBT discussion
0, // missSBTIndex -- See SBT discussion
u0,
u1
);
}
__forceinline__ __device__ float randfloat(unsigned int i, unsigned int p)
{
i ^= p;
i ^= i >> 17;
i ^= i >> 10;
i *= 0xb36534e5;
i ^= i >> 12;
i ^= i >> 21;
i *= 0x93fc4795;
i ^= 0xdf6e307f;
i ^= i >> 17;
i *= 1 | p >> 18;
return i * (1.0f / 4294967808.0f);
}
__device__ float3 cosineSampling(float u, float v, float& pdf) {
float phi = 2 * M_PIf * v;
float theta = 0.5 * acos(1 - 2.0 * u);
float cosTheta = cos(theta);
float sinTheta = sin(theta);
pdf = cosTheta / M_PIf;
return make_float3(cos(phi) * sinTheta,cosTheta,sin(phi) * sinTheta);
}
__device__ float3 hemisphereSampling(float u, float v, float& pdf) {
float phi = 2 * M_PIf * v;
float theta = acos(u);
float cosTheta = cos(theta);
float sinTheta = sin(theta);
pdf = 1.0 /(2.0 * M_PIf);
return make_float3(cos(phi) * sinTheta, cosTheta,sin(phi) * sinTheta);
}
__device__ struct Ray {
float3 origin;
float3 direction;
float tmin = 0.001f;
float tmax = 1e16f;
};
static __forceinline__ __device__ void orthonormal_basis(const float3& normal,
float3& tangent,
float3& bitangent)
{
float sign = copysignf(1.0f, normal.z);
const float a = -1.0f / (sign + normal.z);
const float b = normal.x * normal.y * a;
tangent = make_float3(1.0f + sign * normal.x * normal.x * a, sign * b,
-sign * normal.x);
bitangent = make_float3(b, sign + normal.y * normal.y * a, -normal.y);
}
static __forceinline__ __device__ float3 world_to_local(const float3& v,
const float3& t,
const float3& n,
const float3& b)
{
return make_float3(dot(v, t), dot(v, n), dot(v, b));
}
static __forceinline__ __device__ float3 local_to_world(const float3& v,
const float3& t,
const float3& n,
const float3& b)
{
return make_float3(v.x * t.x + v.y * n.x + v.z * b.x,
v.x * t.y + v.y * n.y + v.z * b.y,
v.x * t.z + v.y * n.z + v.z * b.z);
}
__device__ float3 RTAO(float3 firstRayOrigin, float3 firstRayDirection, unsigned int idx) {
Ray ray;
ray.origin = firstRayOrigin;
ray.direction = firstRayDirection;
Payload prd;
RayTrace(
params.traversal_handle,
ray.origin,
ray.direction,
0.001f, // Min intersection distance
1e16f, // Max intersection distance
&prd
);
float sample = 0.0;
if (prd.is_hit) {
float3 n, t, b;
n = prd.normal;
orthonormal_basis(n, t, b);
Payload prd2;
for (int i = 0; i < 100; i++) {
}
}
}
__device__ float3 Pathtracing(float3 firstRayOrigin, float3 firstRayDirection, unsigned int idx) {
float3 LTE = { 0.0,0.0,0.0 };
float3 throughput = { 1.0,1.0,1.0 };
int MaxDepth = 10;
Ray ray;
ray.origin = firstRayOrigin;
ray.direction = firstRayDirection;
for (int i = 0; i < MaxDepth; i++) {
Payload prd;
RayTrace(
params.traversal_handle,
ray.origin,
ray.direction,
0.0001f, // Min intersection distance
1e16f, // Max intersection distance
&prd
);
if (!prd.is_hit) {
LTE += throughput * make_float3(1.0, 1.0, 1.0);
}
float3 t, b, n;
n = prd.normal;
orthonormal_basis(n, t, b);
float pdf = 1.0;
float3 local_wo = world_to_local(-ray.direction, t, n, b);
float3 local_wi = cosineSampling(randfloat(idx + i, i * 10), randfloat(idx, i + 1), pdf);
float3 wi = local_to_world(local_wi, t, n, b);
float3 bsdf = make_float3(0.8, 0.8, 0.8) / M_PIf;
throughput *= bsdf * fabs(dot(wi, n)) / pdf;
ray.origin = prd.position;
ray.direction = wi;
}
return LTE;
}
extern "C" __global__ void __raygen__rg()
{
const uint3 idx = optixGetLaunchIndex();
const uint3 dim = optixGetLaunchDimensions();
float3 ray_origin, ray_direction;
computeRay(idx, dim, ray_origin, ray_direction);
unsigned int p0, p1, p2;
//optixTrace(
// params.traversal_handle,
// ray_origin,
// ray_direction,
// 0.0f, // Min intersection distance
// 1e16f, // Max intersection distance
// 0.0f, // rayTime -- used for motion blur
// OptixVisibilityMask( 255 ), // Specify always visible
// OPTIX_RAY_FLAG_NONE,
// 0, // SBT offset -- See SBT discussion
// 2, // SBT stride -- See SBT discussion
// 0, // missSBTIndex -- See SBT discussion
// p0, p1, p2 );
Payload prd;
RayTrace(
params.traversal_handle,
ray_origin,
ray_direction,
0.001f, // Min intersection distance
1e16f, // Max intersection distance
&prd
);
//float3 lte = { 0.0,0.0,0.0 };
//for (int i = 0; i < 100; i++) {
// lte += Pathtracing(ray_origin, ray_direction, i*(idx.y * params.image_width + idx.x));
//}
//lte /= 100.0;
float3 result;
result = prd.basecolor;
}
extern "C" __global__ void __miss__ms()
{
MissData* miss_data = reinterpret_cast<MissData*>(optixGetSbtDataPointer());
Payload* prd = get_payload_ptr();
prd->is_hit = false;
prd->normal = make_float3(0.0, 0.0, 0.0);
prd->position = make_float3(0.0, 0.0, 0.0);
prd->basecolor = make_float3(0.0f, 0.0f, 0.0f);
}
extern "C" __global__ void __closesthit__ch()
{
const float2 barycentrics = optixGetTriangleBarycentrics();
const unsigned int inst_idx = optixGetInstanceIndex();
const unsigned int prim_idx = optixGetPrimitiveIndex() + params.prim_offsetsinst_idx; const HitGroupData* sbt =
reinterpret_cast<HitGroupData*>(optixGetSbtDataPointer());
const float3 v0 = params.verticesidx0; const float3 v1 = params.verticesidx1; const float3 v2 = params.verticesidx2; const float3 n0 = params.normalsidx0; const float3 n1 = params.normalsidx1; const float3 n2 = params.normalsidx2; const float3 position = v0 * (1.0f - barycentrics.x - barycentrics.y) +
v1 * barycentrics.x +
v2 * barycentrics.y;
float3 normal = n0 * (1.0f - barycentrics.x - barycentrics.y) +
n1 * barycentrics.x +
n2 * barycentrics.y;
Payload* prd = get_payload_ptr();
prd->position = position;
prd->normal = normal;
prd->is_hit = true;
prd->basecolor = sbt->basecolor;
prd->emission = sbt->emmision;
prd->is_light = sbt->is_light;
}
extern "C" __global__ void __anyhit__ch() {
float alpha = 1.0;
if (alpha < 0.5) optixIgnoreIntersection();
}
extern "C" __global__ void __closesthit__shadow() {
const float2 barycentrics = optixGetTriangleBarycentrics();
const unsigned int inst_idx = optixGetInstanceIndex();
const unsigned int prim_idx = optixGetPrimitiveIndex() + params.prim_offsetsinst_idx; const float3 v0 = params.verticesidx0; const float3 v1 = params.verticesidx1; const float3 v2 = params.verticesidx2; const float3 n0 = params.normalsidx0; const float3 n1 = params.normalsidx1; const float3 n2 = params.normalsidx2; const float3 position = v0 * (1.0f - barycentrics.x - barycentrics.y) +
v1 * barycentrics.x +
v2 * barycentrics.y;
float3 normal = n0 * (1.0f - barycentrics.x - barycentrics.y) +
n1 * barycentrics.x +
n2 * barycentrics.y;
Payload* prd = get_payload_ptr();
prd->position = position;
prd->normal = normal;
prd->basecolor = make_float3(1.0f, 1.0f, 1.0f);
}
extern "C" __global__ void __anyhit__shadow() {
float alpha = 1.0;
if (alpha < 0.5) optixIgnoreIntersection();
}
payloadは基本的にアドレス送り付けてポインタとして扱うのが普通らしい。やり方は単純にアドレス64bitを上位、下位32bitに分けてunsigned intの変数として保存する感じ、
anyhitshaderでIntersectionをスルーするときはこれ使う
optixIgnoreIntersection()
slosesthitでprimIDを貰う時、GAS内のprimIDなので全体のPrimIDじゃないことに注意
なので各IASにindexで参照できるようにoffsetを用意しといた方がいいと思う
8/18
マテリアル関係の実装
前日shadow用のprogramがsbtoffsetを設定しても何故か動いてくれないという問題があった。
→これは単にSBTのHitshader配列のstrideが設定ミスしてた。1要素の所を配列全体の幅を渡してしまっていたので多分offset1で配列外にすっ飛んでたんだと思う
修正コード
https://scrapbox.io/files/64df64e3895469001b6f4982.png
GASのoffsetが働かない問題
numSbtRecordsの数が1だったのが原因、いくらoffsetを与えてもnumSbtRecordsが1だと全てのoffsetが(0~0)にクランプされるぽい
以下の様に2とかにして試したらちゃんとGASのoffsetが働いた(fragもその分用意する)
https://scrapbox.io/files/64df651346eb9a001ba5cb34.png
https://scrapbox.io/files/64df65a78e4f7f001b852351.png
fragsの設定はOptiX memo参照。別段何もしなくていいのでOPTIX_INSTACE_FLAG_NONEでいい(ちゃんとanyhitは働く) https://scrapbox.io/files/64df6826b40ddb001c61b355.png
乱数生成
折角なのでCMJをやってみる
マテリアル実装
取りあえずLamberのみ 1000spp
https://scrapbox.io/files/64dfc364c4ad99001b44d5ef.png
device側のOptixで提供される関数に補完が効かない問題
→deviceでのOptix関数はoptix_device.hにある。optix.hはこれをincludeしてるけどマクロで切り分ける形になっておりVS上では__CUDACC__っていう定数は定義してないのでoptix_host.h側を読み込んでる
なので一応直接optix_deviceをincludeすれば補完が効く(ただ正規の方法ではないと思うのでbuildするときは消しておくといいかも)
https://scrapbox.io/files/64dfc5435d8cf1001cd48b6a.png
8/19
大セグフォ時代
何故かLambertのクラスを作ってやってみたところ謎の実行エラーが出た
https://scrapbox.io/files/64e090fc08082b001be27025.png
OutputBufferのcudaFreeでエラーが起きた 調べてみると
sbtIndexOffsetBUfferの設定ミス これは確かめたけど問題なし
調べてみる限り文字に出ているようなcudaFreeのエラーではなく、恐らく実行中にセグフォ等が起きたことによるエラーだと思われる。
GPU計算中→実行時エラー→即終了→bufferのアドレスがバグってるままFree
(長いサンプル数で実行した際直ぐにエラーが出たため、エラー時即終了という仕組みになっているっぽい?)
8/22追記、なんか知らないけど治った。同じ現象が再び起きたが標準のreflect,refract関数を使わないことで何故か防げた
8/22 治ってない
8/25追記、恐らくこれはスタックオーバーフローが原因だという事になった。解決法はCudaのHeapSIzeを増やす関数を呼び出しておくことで治った。
8/20
アニメーションレンダリング追加
前回のやつをそのまま流用
Texture 、render設定json等
8/22
BRDFの日
なんか知らないけどOptixが標準で渡しているReflect関数を使うとエラーが起きるという現象があった。(エラー文は上記)
治ったと思ったら再び発症
調べてみる限りGGXのサンプリング部分で起きている(だけどこれは前回では問題なく使えたし、単なる関数)
Nsightを使ってみると各数学関数に~fと付けてないから起きてるっぽい?→全部つけてもエラーが出た
Nsightの使い方がよくわかっていないのでどうにかしたい
乱数が原因かと思ったら、固定の値を入れてもおかしくなる?
というかそもそも他のやつで同じの使ってるから何でそれでおかしくなるのかが不明
乱数部分の異常?
乱数生成するだけではエラーが起きなかった
と思ったら乱数を生成するとエラーが起きることがあった
そもそも他の場所で使ってるしシード値の影響?
特にシード値にも特筆するような値とかは見当たらなかった
他の関数に異常?
今のところサンプリング部分と乱数生成部分を
あまり思いつくことがなかったのでOptixの不調またはドライバの不調を疑った
再インストールしても同様のエラー
他バージョン(8.0)に移植しても同様のエラー
再起動も同様
マジでなんだよ
仕方がないのでどうにか書き直して見る、乱数あたりが怪しい感じがあるので一旦rnd関数で実行するようにしてみる。
これでダメそうならちょっと相談してみる
だめそう
Nsightで出たエラーはこんな感じ
https://scrapbox.io/files/64e521ef5532cf001bc4d505.png
https://scrapbox.io/files/64e521e58d217e001b76a34f.png
最悪の場合を考え、お茶を濁すシーンを考える
8/25
調べてみた結果、等価なプログラムのはずではあるのに何故か関数を挟む挟まないでエラーが出ることが分かった。
それをツイートしてみたところYoshi's Dreannさんからスタックオーバーフローではないかとのご指摘を頂いた
そしてスタックサイズを増やす方法を教えてくれた
これをOptixのContextの設定部分に追加してみたところ、なんとエラーが出ずにそのまま通った。ありがとうYoshiさん
前回ではこのようなことは起きておらず、また知り合いの人からもそんなエラーを聞いたことがないのでもしかしたら単に設定ミスがあったのかと思うが何処だかよくわからない、取りあえずこれで行ってみる
ちょっと外法感があるので実機で動くかどうかが心配
メタマテリアル実装
多分正常
https://scrapbox.io/files/64e793c775585c001c793e00.png
8/26
Disney BRDFを実装するなど
https://scrapbox.io/files/64e923f1605d75001b5b13a2.png
Reference PT spp 1000
https://scrapbox.io/files/64e923a83039a4001ca140ce.png
NEE
https://scrapbox.io/files/64e924f8dba05c001c871f31.png
MIS
比較
PT 50 spp 5.7 s
https://scrapbox.io/files/64e92aa2861f94001b287b8e.png
NEE 15 spp 6.8s
https://scrapbox.io/files/64e92a2e2b6dce001b2bf150.png
MIS 10 spp 5.7 spp
https://scrapbox.io/files/64e92ac7d83007001d1501a7.png
Uchimura Tone Mapping
実装はShadertoyの人から引っ張ってきた
仕事をする
https://scrapbox.io/files/64e9c643d83007001d1976f3.pnghttps://scrapbox.io/files/64ea68a537ecca001bb9b10e.png
ACES ToneMapping
8/27
NEEのバグ取り
なんか知らないけどNEEでDisneyBRDFをやるとやたらと明るくなるバグに遭遇
https://scrapbox.io/files/64ea68d01d66db001ced3fa2.png
NEEのPDF、Samplingに異常がある等を探ったが、特に問題が見つからず途方に暮れる
なんかハイライトがやたらと滲んでいることがわかる
Lambertでやると問題がなかったのは昨日確認していたため、ハイライトのGGXがバグってるのでは?という疑問が浮かび上がる
GGXがバグってました
実際、昔のレンダラで使ったGGXをやってみると一致した。つまり、GGXの評価にバグがあったということである。
多分PTだとわからなかったのはサンプリングの方が正常だったせいでバグが目立ちにくかったんだと思われる。(全体的に高い値にBSDFが鳴るバグだったので)
調べたらGGXのLamda関数におけるsqrt(1.0 + 1/a^2)をsqrt(1/a^2)と計算していた。なんたる様https://scrapbox.io/files/64ea6fbec8c46d001c239e14.png
https://scrapbox.io/files/64ea70c14bad9f001ca6b76e.png
治った
Oh Multiple Importance Sampling
残り実装したいこと
Denoise
Thin-Film Interference
薄レンズ
Denoiser
以前のコードを再利用
ちょっとバージョン違いでAlphaModeの指定値がOptixDenoiserAlphaModeという列挙型になっていた
https://scrapbox.io/files/64ea771ca458ee001c55f966.png
こんな感じの定義
Upscale2x
解像度を2倍にしてくれる機能
やり方は簡単そう inputとoutputそれぞれ解像度を別々に用意しておきメモリ計算ではOutputの解像度を、
https://scrapbox.io/files/64eb5cf6daafe4001b146413.png
SetUpの方ではinputの解像度を渡す
https://scrapbox.io/files/64eb5d0b1ac8c7001bb76854.png
それでinputに渡すレイヤーの解像度はみんなinput_width,heightの解像度を使い、outputではoutput_width,outputheightを使う。
0.5解像度でレンダリングしてデノイズで大きくしてもらって高速化を目指す。
通常解像度 10second 10spp
https://scrapbox.io/files/64eb6254878473001b70f4b1.png
半分解像度 12second 100spp
https://scrapbox.io/files/64eb628a0cc709001b02a334.png
解像度が低いとあんまり高速化にはならない気がした。(当然と言えば当然)
どちらかというとそこそこ解像度を更に高い解像度にする時に使うのが良い気がする
なのでせめて1000以上の解像度あって、ノイズがかなり少ないことが条件になりそう。
なんか自由な解像度でも良さそうな雰囲気を感じたので0.75倍レンダリング出来ないかやってみる
出来なかった 変な場所が映ってる
https://scrapbox.io/files/64eb63257c6a9000210d3d6c.png
やっぱ2倍じゃないと行けなさそう
多分本番ではやらないと思う
8/27 ~8/28
大体終わったのでテストシーンを提出してみる。
8/27の時にはPTXのコンパイルエラーで実行できない連絡を頂いた。相談したところCUDAのバージョンが違うことが分かりCUDA12.2からCUDA12.1にしてビルドした
EC2インスタンスを自分で建てて実行できるかの確認をした。インスタンスの立て方は実行大臣の資料に従って行っただけ
(去年vCPUの申請をしたリージョンが間違えて大阪だったので一応大阪リージョンでやった)
CUDA12.1を入れて見たら実行出来たので多分本番環境でも大丈夫だろうと思われる
8/28の提出、無事に通った。やった~!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
後はシーン制作、メインはお寺なのでブッダの例の銅像で、あと負屈折物質を見せたいので水がある場所。
なんとなくでけえ建物が作りたかったのででっかいお堂みたいな感じにしてみた。段々、これ千と千尋の油屋じゃんとなった
https://scrapbox.io/files/64ecda8fa95212001c7ddb5a.png