From b0dd97ee6bf8d5daa587da40ad941efac68152df Mon Sep 17 00:00:00 2001 From: brett weiland Date: Sat, 1 Jun 2024 01:36:18 -0500 Subject: [PATCH] got raycasting laid out --- .gitignore | 2 ++ 3d_fractals_cuda.prf | 1 + build/main.o | Bin 32912 -> 0 bytes entity.cuh | 31 ------------------------- makefile | 15 ------------ render_object.cuh | 12 ---------- scene.h | 12 ---------- sphere.cuh | 17 -------------- src/.vscode/launch.json | 19 +++++++++++++++ camera.cuh => src/camera.cuh | 26 ++++++++------------- common.cuh => src/common.cuh | 26 ++++++++++++++------- src/entity.cuh | 30 ++++++++++++++++++++++++ {include => src/include}/helper_math.h | 0 kernel.cu => src/kernel.cu | 3 ++- kernel.cuh => src/kernel.cuh | 0 main.cu => src/main.cu | 8 +++++-- src/makefile | 25 ++++++++++++++++++++ src/render_object.cuh | 14 +++++++++++ src/scene.cu | 14 +++++++++++ scene.cuh => src/scene.cuh | 23 +++++++++--------- src/sphere.cu | 9 +++++++ src/sphere.cuh | 17 ++++++++++++++ 22 files changed, 178 insertions(+), 126 deletions(-) create mode 100644 .gitignore create mode 120000 3d_fractals_cuda.prf delete mode 100644 build/main.o delete mode 100644 entity.cuh delete mode 100644 makefile delete mode 100644 render_object.cuh delete mode 100644 scene.h delete mode 100644 sphere.cuh create mode 100644 src/.vscode/launch.json rename camera.cuh => src/camera.cuh (64%) rename common.cuh => src/common.cuh (71%) create mode 100644 src/entity.cuh rename {include => src/include}/helper_math.h (100%) rename kernel.cu => src/kernel.cu (82%) rename kernel.cuh => src/kernel.cuh (100%) rename main.cu => src/main.cu (87%) create mode 100644 src/makefile create mode 100644 src/render_object.cuh create mode 100644 src/scene.cu rename scene.cuh => src/scene.cuh (50%) create mode 100644 src/sphere.cu create mode 100644 src/sphere.cuh diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..8765681 --- /dev/null +++ b/.gitignore @@ -0,0 +1,2 @@ +old_cuda/* +src/build/* diff --git a/3d_fractals_cuda.prf b/3d_fractals_cuda.prf new file mode 120000 index 0000000..8e05f78 --- /dev/null +++ b/3d_fractals_cuda.prf @@ -0,0 +1 @@ +/home/indigo/.unison/3d_fractals_cuda.prf \ No newline at end of file diff --git a/build/main.o b/build/main.o deleted file mode 100644 index 737de980dc6f5924fa801648135a1740af6cf608..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 32912 zcmcJ&2{={V8~=SILmrX2B10uZA(0_UDMOUXkdWCiWy(wgB^3%KDm19flp#}PN`;g{ ziZmD^3Kd0EynCN>uO83()%E_b|9idby3W~qefD?VYr5BPwtbEetu5NK7#JuhFi_S} z{*56@p==NO_k2C=bUkGeMTlae8_}y9ap1ddl-76M2qyAMtYr;JjdTwUo}(U2Gm(bLPpZI+iOPrBKsSQ6voQcxT&iy`1 zhO3oq#D#w1MbSMOUPVCNh+0}SB78p)3o9qzDV*_DN2+{fqi75Sx)Bd_Ptrygfbmb z17(V6zeOgYB}V1bJy~>_EN<8!EAg7I8_}g3MXx|qd>ySG47w3JSPA%)7TYjSH)5P# zny`DK1Ik1J3G1Q6wF(N+%)iY>(V8pjB-uc^5f62v%E`u!Dk7V6Lm5#l$_83xbt9T| zBks}+hFAG<68Vs}kQsVLQAIR%GQ5+xR7TW*a%*N@bkWHo^6G!A3otLDwfh^!opGlCY8*+vi}FUo{?q3oxd z|2?PI<0R7-nv$M9VVUA+>*7Xr5cTyGb)x!;I#B&x?5SdGTHfBC-fKj~eb^}Weh#*x zzEmHoFV)-A!-dS34A#~LJFT~F+hnpu%X;%BVqO?Bxy3-4N70>4zdnn?z@STs%er_tdeTnpZQb0QyluUlX)!NvPhU^p zATKH{=1p~zwYB$k@$~rr**LvIe|cX;a7GF=_6&^yjnh0}2Bn`8T|Z5p5w1=<{(>Xn zWYa$Xs=34|OiKSo|L@Mo?!da9W@tT_vxFl$K=+|2{& z*@%G55Ww`{eO0$?{i8Zcbl8&XRe8%B8TNANIBt4g66Hp!MrdQ>OEH9AM9|G$*7p+8%1o5`L|N)^{dC4+|IKtsjH2U6+I=vnj2HY+UF2?>PviT zkcP%0>Vc`W@P?cHO1|208+sP-3{6cL4WzGdsI%T$u5F#NV!p`EW8bQRiwtJ-tF)Ev zesS|Vn~ay_W@(uz#s-mF7Zy&s4=xPILe>#r*(_J?w_Rpi9Ic5AnMKHaUK&?J?0;^@KKPd~{t2*>^U zkeSOJ7pV0~>srDS$tk{DI_VZGZJZ7FSN!F;CHU<0ZBvef8%ITys!W*^ zZiLTiuvwC~w{5oD=^FxWY{m16S(WA{9U7W#!QQc;_#d}`Htx2IdF3|xw)<*!-W>}c zVzFSdVBE(uNNMBBTVktI9es2WqZ*ID(x6aVS>AZf^}`2Pisz0=4$8L$*v8aUux62`K0oEulnOJ z<3iI_`4@Y9q60;8<>dP<+|7DqtkeBMz9clC;LqNs@#@?InXZR>Cs zIISGd|4!NLCF|U3rG2d*5=12wpH@4wBS6rBw6`H$% zr!7~}ey`t)u1EYSC&bJX#c!rp9E*8TN&_+EJ^g{nQ8=)6H_ z415H>Z|mFqnL0YFv|)>{*unVZh#>3DviHBVzSI7^t$uB1N&VsL_x98;QK`6i$3cGH z^REu^d^RT)M4q)J8k*nI{uo&Kt?6pqlEndEuhrRQckycRXV<+`)arK~xXh}$cokQO z+~m7|%tACT2yrbqd#z44JLSycB_tEHHN-HlI^+fOFnKHT@!qkCWN^?CDH1D4&JOgd6mKB8`=FgV(IiaAJF;$6_^ zI=f7l0GB@ftyR9R!f}b*EtzKGtBYhGD;T!e?C}oIH7vc$AK7}bsDFb@flii1?{W2! zNZ#QyOSFaE%rky%I-9uBFsv@v(zQ|YalwP|mpf9+vJ`n;j1rCOf*e2TZZ)Y3GGs}0 zyz{Ix+v=V8>ZdEDx>E$6=`Od-QqU`EN}bfr^w;N0DzwYzSC!GKy>?t@U2J0E^;3FZ z&&aI0J9g<~`0%PXS}`x1NAiA+Ka6;r;kd@?Lw?CfU{T_*dgj{)a<{d{_t?k(Y|8to zw%y?x_c}rU#)N(e@A<~{rC#1mHhtw|SokEhOmek^W#Xuz&nZ=> z^z8QYTyL{4$uj9R<*ktyK zzsk}4xO+}Wj;350cSV+FrH}Z}BDMpAdf%T&e?7UpYX_h7aaEq;%d@tKSA(vgc;7Odu(Ah7Mffs)= zNn42!v9#{36C*D&z9;UV_>uX2Z2!b^<2y3eQxVb!Wv!?7N_#9Dc_$L@GZmpe@%~+M z$l`5N50XcpoSl9{dvCb>idwuZIa7|gmwD9jZsy1W!*412sS6DK+Eb+svjUql^`uRx z&6(1sS^2UinTD%x85ik&t8uvKciU|9rEP1Ti?4DT*zJDvbfqkdmP09b#ixZnx48XE z9fOk$!~aRn3eq##*=hWEXBL-iYmJ^?jDFGXOkwZL=Tfal536U7ED?Wg(YZUm?4N~2 zPG^4J+$j6}*yPPD-WQ9cganK9o~?~{5`L1k)Gt$ga*->q)L z*IAmGH;q;sJmo&Ggt%f&Cs=SOztG^%*qPp-Y}v?b{5R-dKyLpMzxJM4I%8X=<^6}vZjI?J%U z-|&}tCP%69C&|_Ys(W}B2L8N!+OX@=lhqPaDa~hn3%;p7@k>8*u=LscdmZsp!v>9J zmrR;_FC-UMq#4V|wwk2AlroWw++=Zcp8VH{V@>SQ>o-?A&aaiRU_8*Tz?rrBS&F6q zdLw6#6Jz?)@dGT1?JtiwEHM4}lU2qr-mQ0u$Y9e(lm2DZ6sm>8k4`i$h7GzHNdj+?U_v{F<0}eO%adEtfRMNXv)AX(>NH z-BoxNwAgsXGt=zVmeJorvI^G<2PJM;en{u(QU%~s`&37{1FPRLHTv^>+hWW-9eUvGR~bd-143U;FP0q%^pua89`nDtZeQ|& z@0zkwDU1svFNs(uMv2dk-2Gw1vSGZEZ}#Mu=w*GwHHFjeyqBhif?k~w2xp+@>tU+|qrPS_=_!k$CT@Rnd!FfQZS zI6v!EL<+mf@TNr@{o1?P!+18Juv+2$op3(zu+5 zXF=BMh)?pyw|U~`6KD18t;6?tkGclfNpCDt+p9*r$)DU&q?WO-|LbRQ$qCI939&yc zdVE_Otaqq(PegbBJZr-#8!H-OwJ7u&W6ceVCtHgF-Hh3~ISecXQimURijH0DQ|;lBeC=wH!zG39nXFmT^T(Ogs?_GM z{mR`?HQ;unb8Y8^L9yz&o$SS?pB3AxuK&~Aaz(;VAe-k7=bfZn?pM-n%K~@A3g1@G z+_x*+x zB_B~xj8a?F$gfgUP+HRU^zkCeoCdpq&=z*5P5eJtk{7C6Eoh7CRA+PY<-g69T<-P2 zAz#000h5y(e-cl!T-S*PyCp|_*c-pE@ZX@qJo(HipY!elPA5(N4%K9T6`sPj%Uwc; zefIEMtV}lZsyKW9*sSCiom(E-B^_&FbModFR85vvF)H|;5o$EcNqRajw1x7%c*Q!7 z&fVcY`}hUd3BCOCh0E-W)c4G^FXN~85BT@!dQOU4>Q8TTW2oSMeIr%wZflQ1(O2=( zi-A`YhC@I0nl)PoS_DU}IREw1iRM*bdiGCth#ok}$MWJcH$_hE+Kryj6xpgu#`|3# z*Gw;tvL0#?e*I&E=NDHa>7f@Vb#k6*Z7nSFUC)`Dw)6Pmc@v$2pM$EapS6mLtvjwO z^t^mEwfl9bukp6+Z!Eq?@eHiuv|?4W9Geai*}ODoEKqsZhwBYFVV*0}B@XO~c5N!8 zc2ieBxp8a9?O-!^QS+YQM-v8IBxRD_aBcq{ltN2MiC9k0lo6?la9gX z9;;Z|XZ-S$Y-va;>TVKG7|_ejJ#g>QbJr#vnUHzkJ{Y%zz1-8qIXJ;eH6E>2v+@)w zUg(!7IqmUb;;ezMS6WPwXwMp^nAgadCdM#$~^^0=b(mP)JZYL68qmu@`U zxU~GhgLLCpm#@CCS<^pQQ3FLUJ{uYy8pgvmX1$N({ct>8bya z&(`E~{9K3N6z?LBD`PuX_i`{A9ZbC|d;Rmhwfhf+t&-oW9o6>YzUAS`!iza+euYVe zhfDj)PgQVbSGTO%)^hZ8tsG+*Z`}hov-6LXl$S@}8C>?kJ%T!tw?s8Nwc>+b>bOQO zvwc4GXlqlxM9o{VqG`T{9ac9d*##__kK5kla@PG2?mzveU|RQxQeC0k87}M8b2pqe z7N3{s9{+K4;ft@@8*E0-Ub|O$RaoswxZaQ4o%{#e4_Ee9UU#>;Z?X7FZ;o3;LFh5_ zic>Q3O(Wm^Zi{nWX?XFrS?tHsFvZBq-qJy%$cDzZ1r4qPg>H>(xg`^<3!hue^?se+ zcm4Lo>IJNilT-DJ47hIj|CAYezW#Ztj_J_qJ;_Zsn|9wQ98QyeI^Z+7U(5U310|W8 zvvS9**Z!!s`r zmlXAB1okoRuK4Jb6xR@&P~z}DxqQ5AlkUxh{k^-Mimz5II-A4&ag`B&?A+H`JLU9k zyQZX8`RRuGvdm%Mo$%GH#(rL|q|y2c##O4L>-sY!Q-jY*Z;ExjSGRLKt!jxtyr+Z! zgVOypjnx6EgM$)^RUU_Qn?C&W@u2b2Pwa!K!b{F`tY|J&H<i}^*b*|_@5M4&&g|7>CBv$+jquk zskWUb>!UTR2I?=ej5}62#(1B)%UAv*{<0;G69s+G=kGn19DOgM`OISu$K*0K2CMe+{0}#3mi1O8{m2fq zA9#C9iTl(ZyNr?7LhDm=S6ZkE9xJ}%xk-r4TH}QKvs#bB)4zV4as9=Ys`F8M(}^4J zyB0q98g%MslH7q)!6Q7C8F5qNiJh;G2qZ$p)!`-1iWB#%c9iSYZ0B9~{?-10+|5gr7T+}fXHCRn zR*oqXf%Gw^g|!@$vt4<5u~^^-5RcYMS-a`KNcerSFpa&hGYIINqwO_G$6D z+PjR&tTG3}R;zxl_P)vbp10vvn5BY*nxC(anB@wp*Wo*pbFC<*90!~ieD_+V*s3M+ z^2B(&Kp3NbUar;Vw?}QhY1u0VJ5xIc=Rf8ObdVa%jLDW6lqfy1_wgL1HErzaf=TZk ziY~_K$>$mEm=HL~_>fIbaz^`Ah|*j3-}XCT*5Th#%zZ!odwI6ysa}=(jUv2e{s*KR zvZ$*+ehvR;F~giWb5>F$^?yw&g#9z;^hkXeW2EvU_fKm#S{;vcn06bSd&9W;oW-_* zM&6~KY?D_OnyFXM8UK-6?@|`G`=;sUl1YbqR>MbRj&n^HJQ_RrZg|5#x$zg5n<#9} zRVfg6cQ_bm;=1`;#lDXb4&$j;#>V=11H_lU_xP1>{G*;FAtz49lefgvxUkYiip}=P za;m_CW=p9@hxlIm#f%&`eU*Q;u6^xbHJkYDmTh;d2Fe}o=SV6lIQqU%YfA3R8C@{g zSgLk?gt2vLL-N};tN9II*A`iaUz)t5$OSL6W zWR~ls4;m{NCWNcp>XAErVMXufg$e5RxqJfRErr}~jFOcTT=!KoZp(T2arDB|N4F`9 zhQf9ayO``q5DtylX<-sA((=}I-6Byd%pcz z-I4ugrVN~T+fMk6bGAr3+~IF7aJX=1*`3aa=c_jcNXj>M_!Jq5U))$(qY?BdGA*ld z!C?5#^Nz`bdhyfeYC6I!%MN#oQ5k}C5}t6jm$Wy0*!yk0@Xn6x3AxHG>nyjXRa#Cf zmED^1IxyPB937JD;n^AEzVyd|CreuvAKREUTJtIWNl9gQ{Q=Xe)xlR8jMjxo2%p=Z z(|_@?N#@r&&Krr(z4G&CjbBWwV;>w?#X9$$M`1wJjSdzk&-ee-o!e^k!gUu{TG#G( zvkXTK7*z(YnWvX>Ma2?pu~O~}bDZZtqWH6}e763c{qvh1qoyggITqfBDT{25 zz6jdPU_DVb|JPmZ%B8lScYk};<}`KpgsX2OyS$TbMV62Jr!PXhrCZ~M%#KQyeH0tN z5hwnMDZV9b(>%sEOJ4i$c4V6VuEub6#G!uk5+mV-gTnTA>JJ@UbZEoH(1v9%F5T$f zI<=u+rr7bG@Vss6+*Mnh?mV#1^vPN&zJ1A$_NUj1=W$dsr#@jgX?N)8u@HV{o3bmt zow46Id6Yd;WW%Pb_6hD-xnQLt=aVZ?n@Ae(rApgiVWoYMNf2oyJMx7fK4((P-W(U`7 z?K*pZKrz_-gw)iq&+L=x&f**NCNkDb3!N694B-?mmNJ;Ic)YWD37`1sL^-XqI$E9H zC;1wi84Z$SZ}pF+W>7YMK6#HpVWNL6N146x!L^Sy<(D;wF$lfSn;SNw-z^&%KQ
    #v;S+h@PVMj_pwr888e_58U+K_105^)(F1wy#f@G>48h zNR zTiZqV8?i2IS{;@!@{1}B0SINwgpb69l3W3BwRXs^h9 z!3Vtgg$;|htp6C-GL`DkE6VUtwEg)>RhzoW$UI@m_>*BbI(NUp-?TSEpQp=bq#8HJ!DK#{$F_GDoQ71curL)*dwzte;!8!q`5`_)cQd z`6A8ur!``>OzH#H@9oB0bmboFo3{##E(}=9SjNby8k7I{;?YN6Hf=jl>>>V)0A5~AM>Hc){oP40wms*bfj$2>DN1_Vn1xHVC3F>Q_brA$0I800cssQhTrEG&U5w*9=a@i zxuIu>;T^SpmDg4_y9{;agPK7Dx;0HFwkB;BOswCVP?Mx`WuETod0cL5N=xs_Wcm#a zJX59c%K2SNPUg)wwIn-X!82!JPjLbf9A3N zQ1SXv^(PfBU*&JH7J7ZVV}4>yt=VuOzlzrCj71NeKjifpoMXtW9d6F-zP#h0Yto7X z%44OsLWCKa&b>WsHu)@ml4qgF&NZ+1inKoN&*cc$@(A^g-5a6*cD>H{vqr(X4vpP! zxX)Uyt>kQ@ZePEBI>^&&_GTvC{yBA3<-ZC?uMB@nIB;u$^N#*qCc^_ww*?pT%)hL| z(toe@l5G6QRnv^9B?Bn}xBS11jQl*k=_bF}w@O3ZwvtyiEiT8;7OhK9K2xTs7_3~C z|7Fhbew}9~D!aHME;&72M%l7{r=!I9cyg7Zb(DwNhOWugoSHH*;_nlrO;?i=!+1@qLstACN%mgKVaNlcp7<;%}>W|hco6K%hC={!@tvXI>&>#&`R)l%=q z?m4|B^$lM@^$W$t9cg<*1)h5hpYy(Bs*xL$)~Ka1*TH6Tf4_d>n!_oKvDwSxOd`6y59>DcH+y zX3aNxS~NDOLR9v4ocd_q&qrc@>`PPJYhSfR+yB1vg0?H0V}5P6)?UHVF&i8-%3_SeC& zH5OKbbA)`WdrJ}-w)KBIWPIhiv%8n}B(p%u=Gd3~AKsk3$D5U0wET_1CI8o=uCPci#@JtM`d$V-#(K?+I&WD3y-`>+LW;%>RG{tjY@kLeq{!Ot^xzPtm~%Ke|Nt=T;{UA$X5 zCT4N0?ha1o;aM&zA<-^MoIb-(nZKKR@3`J-Ayg4{{pHc&tz2TOmKIj<&CcFXA;vE8 z{rIeJ&EMx}CC}c{7`krlrG*+|4lv zNOoSl{oy10{QF-WeCKze@0}9sN!g+=Nm+t1-})C0&U@)T<`t7`q&3T_M0NF} z5l$maGD6qFq&fx@?9z8+1SV=~hAhiM*l>IKucoTrPo#E z7Yco6M@Dj`@_)JXB)(!qvfX{%gkE?DM|+!9Ww}4wd&An8>H0$vzr;>ZGo~r+xvF1A zi`#UjIS)(=zWV*SOb0%={rkBhen*r?jVSx0=}mJQ%uRj>ZaT5KjeEUK*4QVf)9Xt7 zKg6>34XUXfoB!*?-8Jqka++T+v7Wi7t1|LOvu) z28Lim3Pa#L20?Zf)>z`A5W&b78TiS!4IJ~1&e}Mac?TgNqN2o0op_PqG5N_IUDpj` zWUQJ+6Orj{h!+|7$Ty_fPPai`4<$Jf#K|{w-R9rxmq|{9h2$H$Zu0N-Dv}dn8Tp2; z+tRO#5~!1Sk-?mNL)T}40h0YCknJ;@@`d(%7cOBWTJcZePtTFb5JCo8uh3JVl$(L0 z+a2iWNpPai=wjlX*7gu%r1k3WOSBX6IWE#ZdJYQNsUbLNr-y!5h2H*2kUj*Zy)^A{ok%FE+1V9|jX;uASPAV9G zKTYIy_zQ3N3vd4m@BIrO`U@w=!=L(vjqtNS zalyav)qmmIf8kbt;X!}l6km6*Ku2Hel|)dqlb5HRkP#UrYj;~04_WegWk;&Kti6{P z#oBz!s#Vt39{$!Y9xlFA2SVXt>qH#AeaUMI^lOf`zIHAi&bA&7Zd7YWPj6hMOtau) z>rY($yL9_@f5Ij`Ls|j{KX>;aYcE@GTX$=3s^fNlt<~1#2`bf}MdVd^Lg?=4;71h4 z#lf0*YMGW~rL{F#H*H(rjebOVb!p|LT~VT6G4b%GI=T1|&uMGAc-VRenfSW6k!3>l zB->y{-K+y#e4Qa5vheh}ZL;;XMP(q!CfWnx#?(Nv!3az8v{IAh(DLwg@eMMB95)gV zOxxMo?;#A$sOJBPoK^!xYoeMi)$j3PnRP z+7s(WzTPe#P82WV-Pe(p$lKPQy2aCJ3)P?MM%hd>fT@dzgJ%Fmhw8D}op6PXo^GDr zlr5gN4rCquys6sWp6;|OG{akoYI)id54)4a+V1D$OxAOwr-!49lb^RO?HO=ub0u%8 zhXd7nyE82by~T`zJnWsx=jDT`bVF2kd#@mhCed6To4joU$YPp!IoSGAAzunfXP%Qa z)$(vaVtToaoIU;A95%Xn`cR2N`S=noZf)Wb_%0w?W@%;4QKxabT%036vN^CrU&`hm8y5hKo#eI;VOA_HlEG4Vp&YX)C3gP)wi zm(JkTGdQ`oKz5$Z;N)74c>4@4F@sOc;F2@=)C?{)gDa60B?4J~Y2t}5drBD=Lb&O zUrD@>efb$&eg@aW?2x`l2J}4?(muKCL3X@m*dd?!M?3(tgUWOAFZ}#pc*+d>3N!5I z&ESeNINJY_-HgT>ar!%uoU8-d{~*eTNe0iqNgI6->h zIdrfC=dcN}0*>sU`g>ve26FNs!hZVUcX`f}lMNBV=!f6D0Jgl582RrAVY~qb z#wf-^!7nEpZ3dGaxi?37^qCpgv@e@4`gVgq3u62{?S=z|B8G86 z_~-R{w^5TgUQkp<9oqA`MHe@sQq_=oqZS|gnl`M@d2psF^pTlqM?$rEC3OnU|bB^trg=RA>M)UYM7t9Fun{HnfDm4g#INz*OCFX!x`w8QH+N}c_uK< zBT7q7LHiSAhX?$b3H%Xp8SuA%Fy0R94j0BMR=>9meV=YQIv=!6K!~AT7@#|p6 z4CAA)!Lr79ImD?LKMC{F9*mEJpYO%EAB=0XuS4a(N^XvbZ~)WCz_<&;xG=PH1jboN zgG7kI_)dt&WBf9dCkf-^F9(nT?Hf@!bD_TJn0^SxMJ~qc!A}Y?&J6uhhH(L~Uy1QI zP@el3kAr!m0prSGry1ivV4s8LbyP0_;GLMh9LB|ajNbx38OHd0;NLL53EJl;#?>HS zMwrh~dDy@o*f2f}_2R;KHMGNgj0?g%AcFA#u)hrB(y$+s!FV*3UlHS*z~9s`?g#Ug zCdMsc9@vcWozQOEF@6X3X{H$80QI%Pcs1CkVq6;5H8+eOfcpAioEhflV2qc;xITpO zWT@|Pj30sdIR@hn^x;KG!1#G6&jpP4Li`%WonU-rV|)nKodS$+0e>jN_%fJpt1#XG z^FuAhk3qhT7#D^8TPwybpkA*qeg*9GV0<0e`GD~oVCM_QE1{jgV>|%rHH~o>=wD_S ze`wsHKRL|BxIC|aKwJys zr@;TWVEj6?n-RwAq5Ul|J|FV+z_xIWCE8W=By`BN9;$c{b6=R&@o825&D%Yyzv<&1>(zlU*U znD^db+#mGcF&+w>8_J99!!%4;f$?KtUmN2uz>W>Zi=h1@G2R4zdkW*9VE#FW@!9Be z6UHOJej3Kl!TeB$aip)p_*KxiU_1}(zr^@EvZ07DgmI+*it$gNXMxW#sGY~aUpX-T z5d1(4<47-w@mA2QVSG8vdzu*k0`reK#*y9@<9*ORp%^cR@tBD5THrS^{t3#{i17*F z{TMffaXJgy54GEJXkR{zzkz*%5XRBIO%CHouY~clXdi`fF^KQL_$4T}1ICfw4dYp$ z55YJ;^y4v%uZMXc5#vaI1>^dlzk~4%SeKt;yb|^a0~j}fas3734baYWpx?;L_jNV zIMP>PypS%Zv|zjw`nLn)bHJZRF^=>)5r(hiEGce8y^}3DmcJRZy7~cw?D_b#+^sg{} z4f^F1#!rF$ag1LFe`bUI3#u>Db76cNtZPd#UI_iZ0^_hOP&Q#4>9=4U-G6t$_(fRn z-7r25^$o!|(jUWkFX$67j`UYBJ_qtG!uWAm$ICH3ALbd_Zvhfx2el{C(|%xa2A>1^ zkC+{#AH}!`?9-T`Kau@&P)-hvSC9fCh+!P*r7+$B^SL_4ksjSINBKSk`=*%w2E?r~ zt_~YAZ;T^-FveGc{v^hcJ|5$`pwGbgaZy?`Q1USD2Yy?FaimA<1FEkh=+SzMIMVlF zc8-940^_YPFc@I|Lw4F=Jn~{3>4h-9fnF_&9LAAe8RI)akN*CE@-2dXH^lTPKUa(+ zy*I`kpq`N!NBYwk-vsMj7RHAmUVw2Pn13E(9O;`e&I9B01ICel1moT?KeIu(Q2T6! z`JWr(_rbn2#*uy%#^qr@riXE)H^g`s)Ylc`tHH0lG2RUGLnOwL{xrth;d4P2#*w}N zML5-{KDVI1iV zF>Vk0e^-nny*I|~(LN31EzoYKF@6r_rBsX~eJ;kYfxZgkCeV&`7!Lw}>%chD_h396 z^kW#G3;TvCjEBN}#RdB&)L%%y0OKb>zXIbez!foW5Bcd}9O<`VJPF3L9mW^KcyYz} zVY0wPIE-p)it*R5Zl1$9 z(qG2-dT5_YjK78Xp$FrW;1B50aa2xG7#FPIr-+ZixDdrS(o14|zbL(=>oGnD%A<>M zebC!r9O<1go&*gMit)LyjvUAMcGwT6VjSsnF`fi|^$_Do-;8l{nD;(m9O=g~ZVLU& z3G*&$XQUUv_yFkTFwO{mtBi3+Sl{(9j`W5Y-v)Xr#-+g@+%e7p{TPaIq(6>v^gK-x z##N!;uVTCp_94X>NBRnkkAl7lf$!@4GOL^fnmZ2YMfje*pXYFg^k6do;$8J^|yi!A~+V zJ_hSiKE_kwe%C#WBmEDtV;NRI8NBTmH zAB6GVfN?odTDFuHjORkXe8f1?k7E2dj9;|xM*Z>u#wicXuZW}j?<+Bm^ePzt4D-1O z#>K&JEiv8#<@drk(g$Mv4d|mV{uuT#aTxyv`?oZVBmH%Z&w+hZHO6;BKh|Sh9MH$ z(Q_QA{8doT5KMmx_S>g1j`WEbzX1ANjF*Dn7Gb;s+N~Mm_hB5p{0kq#coU3YM(|Tq zuV*mc7ht>-^l}(ahk4K+<2^9{pTxK+%!ip67l-xkKE_jF|1gAcQ}B~E-fgL%FOG3U57(Wi<-3R05Fn$c7-pIZcP>{=R|igh2lVVtRjwhhSU<;zu#g2>x>tGm=Yx8QV;q%d6~;Xv->n!&>Zo<0#)KjGttt*YgU-`@m1K zF&+;3e2hCoJCtDjIjjri7bbsLq#?M2)v|@Z1+W8g6O<{aZVZ0Xh zA^%9w+k>2@z@=z_7dY~V7AXJX8G3RG1GlBfVfs0+4ys^02gZpO#v8zpA;ytETVwnQ z%tziBe+lhy4&#MTxoa5b0DS>)vK;`?+Tk{)C%@-Th6ar9fO)=&Lv=56U z-5zRZVaS{p<6mJu6vX&8CVF}WjAydYxe3PeMCsfe<7b)a+!y1^p!|s#kAQVB9pi7I zf3IV_9Q^h!#y7#Z>&N&2)OR*)Xi@vPLAwcJJQBu{I>zEWq^iROwcF)k0QZPgD#`KG!-w)2vll4X4_lU>zo1op& zFzx~K!*$FK`W=&MOg{$Yufz1{cU9hEdT*G|hcS-o`(uV3(ljjl6sF}wga~ZJi>zu6 zw9h=?sQsZ@6bVf42=^;hG42LjALDbue=IT1&q|hpLh;0S56qK?FuwC2x;_r$)xe8^ zqw+`5<&;`X{|hQMjqyXknOD#gB0FG)0`5y|H`F%dHPo)Sela8@+nF4?yC4qB8qE&L zLx4kf{pLr2%fb=aNrd<+jORdH3FEIIz6RsN5MPJ!9}w5XI1gDtB52bOzstE2xE{t0 zfE!~x1h^T-vw>S-yb!oO#_s}m#yI)=He_(8AAXk;eLvqD@^Dit!ZS&GHqpvpAGS)$P&MFVoQ;=^YNj`dVA6$v>%fB{VO-*f7XOw*C7;SIuC#P z&ypBXo!qG&gb5m@{f3RLx3_H&00*j_pA+%(KIA{{q@UZ`*?Cj_|8tr64IU3CDw#KF znJ#m4@gV-6C;goKiJcjk_LDoZ_MYzaj6FPksj|dx(#g<%Z-po+8F%vVlO_Hys2A1S z7Zru-Z0$(=*bfz2>i^q8lt@nc1X?GMwe@gP0{(+IOj9UIA-9jSiP!qygnlBaNI%*@ zdUDMm)57xk`|}LBu8mmIvI3fdhdUAax)8q9=-(P_397_H_ z`5O^e{?vcRmFY=Pdb0n>yixv$dk~C+c;Wfyfe`5l2`!*FI=T?_Px)W|L;mF2fzqM; z$$bb}|KAtLb0VPczZ^pw9XiHklwgAYDgRQ)zf_H$kgPvD@xt>b*T+BQzaIAE2rKD_-|x8gJaOeu<@fTTCqe1S@{_)Z@<;7gK(IfR-x=~> zi)Dz?)2&k|xx~qz^1lrEqx-L9{#?WhrAPT!5RA-XCAvVIlOY4{vzn5PJ|iF*Dpoyl z?N8|`@Hx&B3Bd-Qz8$5Xkq3EdDNP4ZjNez0IdtqMuAuYZBH|6wV@f9k(j*jQ#jMr8kyeTLGb{`*tj zWRg7ioPxfuNxBpH!qd+o7(RCph>`&n55+EG)uH(mCLxnkR03%S3NMxbi&Fj{JLe%T diff --git a/entity.cuh b/entity.cuh deleted file mode 100644 index 737ef8b..0000000 --- a/entity.cuh +++ /dev/null @@ -1,31 +0,0 @@ -#ifndef ENTITY_H -#define ENTITY_H -#include "common.cuh" - -//we could make a template to allow double percision, but start with float -//idk how nessesary it is yet so I'll go ahead. -//I know I needed it for zoomin far into the mandelbrot ig, so it's not -//out of the question -template class entity { - using T3 = typename vect_t3::vect_t; - public: - __device__ entity() : pos_(vect_create(0)), rot_(vect_create(0)), scale_(vect_create(0)) {}; - __device__ entity(const T3 pos, const T3 rot, const T3 scale) : pos_(pos), rot_(rot), scale_(scale) {}; - __device__ entity(const float3 pos) : pos_(pos), rot_(vect_create(0)), scale_(vect_create(0)) {}; - - - T3 get_pos() const { return pos_; } - T3 get_rot() const { return rot_; } - T3 get_scale() const { return scale_; } - - __device__ void set_pos(const T3 pos) { pos_ = pos; } - __device__ void set_rot(const T3 rot) { rot_ = rot; } - __device__ void set_scale(const T3 scale) { scale_ = scale; } - - protected: - T3 pos_; - T3 rot_; - T3 scale_; - -}; -#endif diff --git a/makefile b/makefile deleted file mode 100644 index 7f6bddb..0000000 --- a/makefile +++ /dev/null @@ -1,15 +0,0 @@ -LIBS = -lraylib -lGL -lm -lpthread -ldl -lrt -lX11 -$CC = gcc -INC = -I /opt/cuda/include -make: - nvcc $(LIBS) $(INC) -O0 --debug -c main.cu -o build/main.o - nvcc --device-debug --compile kernel.cu -o build/kernel.o - nvcc $(LIBS) -O0 -o build/indigo_worlds build/main.o build/kernel.o - -run: - build/indigo_worlds - -clean: - rm -rf build/* - - diff --git a/render_object.cuh b/render_object.cuh deleted file mode 100644 index 0063d94..0000000 --- a/render_object.cuh +++ /dev/null @@ -1,12 +0,0 @@ -#ifndef RENDER_OBJECT_H -#define RENDER_OBJECT_H -#include "entity.cuh" - -template class render_object : public entity { - using T3 = typename vect_t3::vect_t; - using entity::entity; - public: - virtual __device__ T distance_estimator(T3 point) const = 0; -}; - -#endif diff --git a/scene.h b/scene.h deleted file mode 100644 index 05ce2af..0000000 --- a/scene.h +++ /dev/null @@ -1,12 +0,0 @@ -#ifndef SCENE_H -#include SCENE_H -#include "camera.cuh" -#include "render_object.cuh" - -//for now we only neeed once scene, later we'll expand this to just be a virtual template -template class scene { - - -} - -#endif \ No newline at end of file diff --git a/sphere.cuh b/sphere.cuh deleted file mode 100644 index 33c77f2..0000000 --- a/sphere.cuh +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef SPHERE_H -#define SPHERE_H -#include "render_object.cuh" -template class sphere : public render_object { - using render_object::render_object; - using T3 = typename vect_t3::vect_t; - public: - __device__ T distance_estimator(T3 point) const; - private: - T r_ = 1; -}; - -template __device__ T sphere::distance_estimator(T3 point) const { - return length(point) - r_; -} - -#endif diff --git a/src/.vscode/launch.json b/src/.vscode/launch.json new file mode 100644 index 0000000..33752cd --- /dev/null +++ b/src/.vscode/launch.json @@ -0,0 +1,19 @@ +{ + // Use IntelliSense to learn about possible attributes. + // Hover to view descriptions of existing attributes. + // For more information, visit: https://go.microsoft.com/fwlink/?linkid=830387 + "version": "0.2.0", + "configurations": [ + { + "name": "CUDA C++: Launch", + "type": "cuda-gdb", + "request": "launch", + "program": "${workspaceFolder}/build/indigo_worlds" + }, + { + "name": "CUDA C++: Attach", + "type": "cuda-gdb", + "request": "attach" + } + ] +} \ No newline at end of file diff --git a/camera.cuh b/src/camera.cuh similarity index 64% rename from camera.cuh rename to src/camera.cuh index 61944fe..4d04fcd 100644 --- a/camera.cuh +++ b/src/camera.cuh @@ -9,38 +9,32 @@ //template class scene; //I am soooo high lol -template -class camera : public entity { - using T3 = typename vect_t3::vect_t; - using T2 = typename vect_t2::vect_t; +class camera : public entity { public: - __device__ void render(); - __device__ camera(scene *pscene, const T fov = 1, const T3 pos = vect_create(0), const T3 rot = vect_create(0)) - : pscene(pscene), fov(fov), entity(pos, rot, vect_create(0)) {}; - - //__device__ ~camera(); + __device__ camera(scene *pscene, const T fov = 1, const vect3 pos = make_vect3(0), const vect3 rot = make_vect3(0)) + : pscene(pscene), fov(fov), entity(pos, rot, make_vect3(0)) {}; private: T fov; - T2 size; + float2 size; int steps = 100; T clip_min = .1; T clip_max = 100; - scene *pscene; + //scene *pscene; }; /** //later we'll make scenes objects, rn im lazy (TODO) -template __device__ void camera::render() { +template __device__ void camera::render() { //TODO *really* need to clean this up once you get further //extra dimentions is extra math //either generisize float3 or stop using this fucking template nonsense const uint3 unnormalized_uv = ((blockDim * blockIdx) + threadIdx); const unsigned int img_index = (unnormalized_uv.x + (unnormalized_uv.y * blockDim.x * gridDim.x)) * 4; - const T3 uv = ((2 * vect_create(unnormalized_uv)) / vect_create(gridDim * blockDim)) - 1; - const T3 ray_direction = normalize(vect_create(uv.x, uv.y, 1)); + const vect3 uv = ((2 * make_vect3(unnormalized_uv)) / make_vect3(gridDim * blockDim)) - 1; + const vect3 ray_direction = normalize(make_vect3(uv.x, uv.y, 1)); T dist; T total_dist = 0; - T3 ray; + vect3 ray; int i; @@ -48,7 +42,7 @@ template __device__ void camera::render() { T min_dist = clip_max; - render_object **objs = pscene->get_objs(); + render_object **objs = pscene->get_objs(); for(i = 0; i < steps; i++) { ray = this->pos_ + (total_dist * ray_direction); //gyagh memory lookups diff --git a/common.cuh b/src/common.cuh similarity index 71% rename from common.cuh rename to src/common.cuh index d5a9cdf..9e026f4 100644 --- a/common.cuh +++ b/src/common.cuh @@ -6,16 +6,16 @@ /** template class vect_t2; -template class vect_t3; +template class vect3; template class vect_t4; //this feels so hacky... idk why people are so scared of metaprogramming template <> class vect_t2 { public: using vect_t = double2; }; -template <> class vect_t3 { public: using vect_t = double3; }; +template <> class vect3 { public: using vect_t = double3; }; template <> class vect_t4 { public: using vect_t = double4; }; template <> class vect_t2 { public: using vect_t = float2; }; -template <> class vect_t3 { public: using vect_t = float3; }; +template <> class vect3 { public: using vect_t = float3; }; template <> class vect_t4 { public: using vect_t = float4; }; @@ -32,12 +32,22 @@ template __device__ inline float3 vect templates, but this changes the structure of my entire project in unwanted ways, so I'm switching over to typedefs. **/ -typedef float2 vect_t2; -typedef float3 vect_t3; -typedef float4 vect_t4; +typedef float2 vect2; +typedef float3 vect3; typedef float T; -#define vect1to3(x) (make_float3(x)) -#define make_vect(x, y, z) (make_float3(x, y, z)) +#define make_vect3(...) (make_float3(__VA_ARGS__)) + +//TODO move to color.cuh +typedef float3 Color; +#define make_color(...) (make_float3(__VA_ARGS__)) + +//TODO move to ray.cuh +struct Ray { + Color color; + vect3 start; + vect3 direction; //MUST BE A UNIT VECTOR + int bounces; +}; #endif diff --git a/src/entity.cuh b/src/entity.cuh new file mode 100644 index 0000000..b086644 --- /dev/null +++ b/src/entity.cuh @@ -0,0 +1,30 @@ +#ifndef ENTITY_H +#define ENTITY_H +#include "common.cuh" + +//we could make a template to allow double percision, but start with float +//idk how nessesary it is yet so I'll go ahead. +//I know I needed it for zoomin far into the mandelbrot ig, so it's not +//out of the question +class Entity { + public: + __device__ Entity() : pos_(make_vect3(0)), rot_(make_vect3(0)), scale_(make_vect3(0)) {}; + __device__ Entity(const vect3 pos, const vect3 rot, const vect3 scale) : pos_(pos), rot_(rot), scale_(scale) {}; + __device__ Entity(const float3 pos) : pos_(pos), rot_(make_vect3(0)), scale_(make_vect3(0)) {}; + + + vect3 get_pos() const { return pos_; } + vect3 get_rot() const { return rot_; } + vect3 get_scale() const { return scale_; } + + __device__ void set_pos(const vect3 pos) { pos_ = pos; } + __device__ void set_rot(const vect3 rot) { rot_ = rot; } + __device__ void set_scale(const vect3 scale) { scale_ = scale; } + + protected: + vect3 pos_; + vect3 rot_; + vect3 scale_; + +}; +#endif diff --git a/include/helper_math.h b/src/include/helper_math.h similarity index 100% rename from include/helper_math.h rename to src/include/helper_math.h diff --git a/kernel.cu b/src/kernel.cu similarity index 82% rename from kernel.cu rename to src/kernel.cu index c6895dc..40944c0 100644 --- a/kernel.cu +++ b/src/kernel.cu @@ -4,6 +4,7 @@ #include "scene.cuh" __global__ void render(uint8_t *image) { - scene scene; + Scene scene; + scene.debug(); //scene.render(image); } diff --git a/kernel.cuh b/src/kernel.cuh similarity index 100% rename from kernel.cuh rename to src/kernel.cuh diff --git a/main.cu b/src/main.cu similarity index 87% rename from main.cu rename to src/main.cu index 4aa6cbb..aa74a6c 100644 --- a/main.cu +++ b/src/main.cu @@ -1,10 +1,11 @@ +#include #include #include #include #include +#include #include "kernel.cuh" -#include "raylib.h" int main() { //bluuuuugh i'll figure out occupancy later, this res are easy @@ -37,9 +38,13 @@ int main() { Texture tex = LoadTextureFromImage(image); while(!WindowShouldClose()) { + cudaError_t err; //cuda stuff cudaMalloc((void **)&image_d, res_x * res_y * sizeof(Color)); render<<>>(image_d); + if((err = cudaGetLastError()) != cudaSuccess) { + printf("kernel did not launch! Error: %s\n", cudaGetErrorString(err)); + } cudaDeviceSynchronize(); cudaMemcpy(texture_data, (void **)image_d, res_x * res_y * sizeof(Color), cudaMemcpyDeviceToHost); @@ -49,6 +54,5 @@ int main() { DrawFPS(0, 0); EndDrawing(); } - return 0; } diff --git a/src/makefile b/src/makefile new file mode 100644 index 0000000..4d86411 --- /dev/null +++ b/src/makefile @@ -0,0 +1,25 @@ +LIBS = -lraylib -lGL -lm -lpthread -ldl -lrt -lX11 +$CC = gcc +INC = -I /opt/cuda/include +COMPILED_BIN = build/indigo_worlds + +CU_SRCFILES := $(wildcard *.cu) +CU_OBJFILES := $(patsubst %.cu, %.o, $(CU_SRCFILES)) + +all: $(CU_OBJFILES) + nvcc $(LIBS) -o build/indigo_worlds build/*.o + +%.o: %.cu + nvcc --device-debug -dc -c $< -o build/$@ + +run: all + build/indigo_worlds + +clean: + rm -rf build/* + + +#make: +# nvcc $(LIBS) $(INC) -O0 --debug -c main.cu -o build/main.o +# nvcc --device-debug --compile kernel.cu -o build/kernel.o +# nvcc $(LIBS) -O0 -o build/indigo_worlds build/main.o build/kernel.o diff --git a/src/render_object.cuh b/src/render_object.cuh new file mode 100644 index 0000000..f25f835 --- /dev/null +++ b/src/render_object.cuh @@ -0,0 +1,14 @@ +#ifndef RENDER_OBJECT_H +#define RENDER_OBJECT_H +#include "entity.cuh" +#include "common.cuh" //TODO color + +class Render_object : public Entity { + using Entity::Entity; + public: + virtual __device__ T distance_estimator(vect3 point) const = 0; + virtual __device__ Color get_color(struct Ray ray) const = 0; + +}; + +#endif diff --git a/src/scene.cu b/src/scene.cu new file mode 100644 index 0000000..100af14 --- /dev/null +++ b/src/scene.cu @@ -0,0 +1,14 @@ +#include +#include "scene.cuh" + +__device__ void Scene::debug() { + //const uint3 unnormalized_uv = ((blockDim * blockIdx) + threadIdx); + //const unsigned int img_index = (unnormalized_uv.x + (unnormalized_uv.y * blockDim.x * gridDim.x)) * 4; +} + +__device__ Color Scene::raycast(struct Ray ray) { + if(ray.bounces++ > max_bounces) return make_color(0); + const size_t obj_count = sizeof(objs) / sizeof(objs[0]); + for(size_t obj_i = 0; obj_i < obj_count; obj_i++); + return make_color(0); +} diff --git a/scene.cuh b/src/scene.cuh similarity index 50% rename from scene.cuh rename to src/scene.cuh index 31bb99f..e3d5896 100644 --- a/scene.cuh +++ b/src/scene.cuh @@ -5,28 +5,27 @@ #include "sphere.cuh" #include "render_object.cuh" #include "include/helper_math.h" +#include -template class camera; +//template class camera; //when we get animations with multiple scenes, we'll make this a virtual function //with array of DE objects and cam -template -class scene { - using T3 = typename vect_t3::vect_t; +class Scene { + const int max_bounces = 10; public: //__device__ void render(uint8_t *image) { cam.render(); }; - __device__ render_object **get_objs() { return objs; } - __device__ render_object **get_image() { return image; } + __device__ Render_object **get_objs() { return objs; } + __device__ void debug(); + __device__ Color raycast(struct Ray ray); private: - camera cam = camera(); - sphere sp1 = sphere(vect_create(0,0.4,-5)); - sphere sp2 = sphere(vect_create(0,-0.4,-5)); + //camera cam = camera(); + Sphere sp1 = Sphere(make_vect3(0, .4, -5)); + Sphere sp2 = Sphere(make_vect3(0, -0.4,-5)); protected: //idk why I need to specify the size... why can't the compiler figure that out? - render_object *objs[3] = {&sp1, &sp2, NULL}; + Render_object *objs[3] = {&sp1, &sp2, NULL}; uint8_t *image; }; -#include "camera.cuh" - #endif diff --git a/src/sphere.cu b/src/sphere.cu new file mode 100644 index 0000000..849a37d --- /dev/null +++ b/src/sphere.cu @@ -0,0 +1,9 @@ +#include "sphere.cuh" + +__device__ T Sphere::distance_estimator(vect3 point) const { + return length(point) - r_; +} + +__device__ Color Sphere::get_color(struct Ray ray) const { + return make_color(0); +} diff --git a/src/sphere.cuh b/src/sphere.cuh new file mode 100644 index 0000000..0d228ff --- /dev/null +++ b/src/sphere.cuh @@ -0,0 +1,17 @@ +#ifndef SPHERE_H +#define SPHERE_H + +#include "render_object.cuh" +#include "common.cuh" + +class Sphere : public Render_object { + using Render_object::Render_object; + public: + __device__ T distance_estimator(vect3 point) const; + __device__ Color get_color(struct Ray ray) const; + private: + T r_ = 1; +}; + + +#endif