From d2d1ef622e36adfbd43b0113921a7f76696d79ad Mon Sep 17 00:00:00 2001 From: Aaryaman Vasishta Date: Wed, 7 Feb 2024 16:57:44 +0900 Subject: [PATCH] [HIP] Add PBRT_CPU_GPU in some signatures --- src/pbrt/bssrdf.h | 2 +- src/pbrt/bxdfs.cpp | 29 ++++++----- src/pbrt/bxdfs.h | 10 ++-- src/pbrt/cameras.cpp | 34 ++++++------ src/pbrt/cameras.h | 10 ++-- src/pbrt/film.cpp | 16 +++--- src/pbrt/film.h | 10 ++-- src/pbrt/filters.cpp | 2 +- src/pbrt/filters.h | 8 +-- src/pbrt/interaction.cpp | 8 +-- src/pbrt/lights.cpp | 78 ++++++++++++++-------------- src/pbrt/lights.h | 11 ++-- src/pbrt/lightsamplers.cpp | 4 +- src/pbrt/lightsamplers.h | 8 +-- src/pbrt/materials.cpp | 12 ++--- src/pbrt/materials.h | 8 +-- src/pbrt/media.h | 12 ++--- src/pbrt/samplers.cpp | 16 +++--- src/pbrt/samplers.h | 10 ++-- src/pbrt/shapes.cpp | 61 +++++++++++----------- src/pbrt/shapes.h | 22 ++++---- src/pbrt/util/color.cpp | 19 ++++--- src/pbrt/util/image.cpp | 5 +- src/pbrt/util/math.cpp | 20 ++++---- src/pbrt/util/math.h | 7 ++- src/pbrt/util/rng.h | 20 ++++---- src/pbrt/util/sampling.cpp | 17 +++--- src/pbrt/util/sampling.h | 2 +- src/pbrt/util/spectrum.cpp | 24 ++++----- src/pbrt/util/spectrum.h | 14 ++--- src/pbrt/util/taggedptr.h | 94 +++++++++++++++++++++++++--------- src/pbrt/util/transform.cpp | 62 +++++++++++----------- src/pbrt/util/transform.h | 26 +++++----- src/pbrt/util/vecmath.cpp | 2 +- src/pbrt/util/vecmath.h | 8 +-- src/pbrt/wavefront/workitems.h | 4 +- 36 files changed, 371 insertions(+), 324 deletions(-) diff --git a/src/pbrt/bssrdf.h b/src/pbrt/bssrdf.h index b024f15a6..637fdbe36 100644 --- a/src/pbrt/bssrdf.h +++ b/src/pbrt/bssrdf.h @@ -288,7 +288,7 @@ PBRT_CPU_GPU inline void SubsurfaceFromDiffuse(const BSSRDFTable &t, } } -inline pstd::optional BSSRDF::SampleSp(Float u1, Point2f u2) const { +PBRT_CPU_GPU inline pstd::optional BSSRDF::SampleSp(Float u1, Point2f u2) const { auto sample = [&](auto ptr) { return ptr->SampleSp(u1, u2); }; return Dispatch(sample); } diff --git a/src/pbrt/bxdfs.cpp b/src/pbrt/bxdfs.cpp index 0763bb4fd..ce74b51c3 100644 --- a/src/pbrt/bxdfs.cpp +++ b/src/pbrt/bxdfs.cpp @@ -74,7 +74,7 @@ std::string LayeredBxDF::ToString() const { } // DielectricBxDF Method Definitions -pstd::optional DielectricBxDF::Sample_f( +PBRT_CPU_GPU pstd::optional DielectricBxDF::Sample_f( Vector3f wo, Float uc, Point2f u, TransportMode mode, BxDFReflTransFlags sampleFlags) const { if (eta == 1 || mfDistrib.EffectivelySmooth()) { @@ -169,7 +169,7 @@ pstd::optional DielectricBxDF::Sample_f( } } -SampledSpectrum DielectricBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { +PBRT_CPU_GPU SampledSpectrum DielectricBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { if (eta == 1 || mfDistrib.EffectivelySmooth()) return SampledSpectrum(0.f); // Evaluate rough dielectric BSDF @@ -208,7 +208,7 @@ SampledSpectrum DielectricBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) } } -Float DielectricBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, +PBRT_CPU_GPU Float DielectricBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, BxDFReflTransFlags sampleFlags) const { if (eta == 1 || mfDistrib.EffectivelySmooth()) return 0; @@ -272,7 +272,7 @@ std::string ConductorBxDF::ToString() const { } // HairBxDF Method Definitions -HairBxDF::HairBxDF(Float h, Float eta, const SampledSpectrum &sigma_a, Float beta_m, +PBRT_CPU_GPU HairBxDF::HairBxDF(Float h, Float eta, const SampledSpectrum &sigma_a, Float beta_m, Float beta_n, Float alpha) : h(h), eta(eta), sigma_a(sigma_a), beta_m(beta_m), beta_n(beta_n) { CHECK(h >= -1 && h <= 1); @@ -300,7 +300,7 @@ HairBxDF::HairBxDF(Float h, Float eta, const SampledSpectrum &sigma_a, Float bet } } -SampledSpectrum HairBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { +PBRT_CPU_GPU SampledSpectrum HairBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { // Compute hair coordinate system terms related to _wo_ Float sinTheta_o = wo.x; Float cosTheta_o = SafeSqrt(1 - Sqr(sinTheta_o)); @@ -365,7 +365,7 @@ SampledSpectrum HairBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const return fsum; } -pstd::array HairBxDF::ApPDF(Float cosTheta_o) const { +PBRT_CPU_GPU pstd::array HairBxDF::ApPDF(Float cosTheta_o) const { // Initialize array of $A_p$ values for _cosTheta_o_ Float sinTheta_o = SafeSqrt(1 - Sqr(cosTheta_o)); // Compute $\cos\,\thetat$ for refracted ray @@ -394,7 +394,7 @@ pstd::array HairBxDF::ApPDF(Float cosTheta_o) const { return apPDF; } -pstd::optional HairBxDF::Sample_f(Vector3f wo, Float uc, Point2f u, +PBRT_CPU_GPU pstd::optional HairBxDF::Sample_f(Vector3f wo, Float uc, Point2f u, TransportMode mode, BxDFReflTransFlags sampleFlags) const { // Compute hair coordinate system terms related to _wo_ @@ -489,7 +489,7 @@ pstd::optional HairBxDF::Sample_f(Vector3f wo, Float uc, Point2f u, return BSDFSample(f(wo, wi, mode), wi, pdf, Flags()); } -Float HairBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, +PBRT_CPU_GPU Float HairBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, BxDFReflTransFlags sampleFlags) const { // TODO? flags... @@ -545,7 +545,7 @@ Float HairBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, return pdf; } -RGBUnboundedSpectrum HairBxDF::SigmaAFromConcentration(Float ce, Float cp) { +PBRT_CPU_GPU RGBUnboundedSpectrum HairBxDF::SigmaAFromConcentration(Float ce, Float cp) { RGB eumelaninSigma_a(0.419f, 0.697f, 1.37f); RGB pheomelaninSigma_a(0.187f, 0.4f, 1.05f); RGB sigma_a = ce * eumelaninSigma_a + cp * pheomelaninSigma_a; @@ -556,7 +556,7 @@ RGBUnboundedSpectrum HairBxDF::SigmaAFromConcentration(Float ce, Float cp) { #endif } -SampledSpectrum HairBxDF::SigmaAFromReflectance(const SampledSpectrum &c, Float beta_n, +PBRT_CPU_GPU SampledSpectrum HairBxDF::SigmaAFromReflectance(const SampledSpectrum &c, Float beta_n, const SampledWavelengths &lambda) { SampledSpectrum sigma_a; for (int i = 0; i < NSpectrumSamples; ++i) @@ -996,7 +996,8 @@ MeasuredBxDFData *MeasuredBxDF::BRDFDataFromFile(const std::string &filename, } // MeasuredBxDF Method Definitions -SampledSpectrum MeasuredBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { +PBRT_CPU_GPU SampledSpectrum MeasuredBxDF::f(Vector3f wo, Vector3f wi, + TransportMode mode) const { // Check for valid reflection configurations if (!SameHemisphere(wo, wi)) return SampledSpectrum(0); @@ -1032,7 +1033,7 @@ SampledSpectrum MeasuredBxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) co (4 * brdf->sigma.Evaluate(u_wo) * CosTheta(wi)); } -pstd::optional MeasuredBxDF::Sample_f(Vector3f wo, Float uc, Point2f u, +PBRT_CPU_GPU pstd::optional MeasuredBxDF::Sample_f(Vector3f wo, Float uc, Point2f u, TransportMode mode, BxDFReflTransFlags sampleFlags) const { // Check flags and detect interactions in lower hemisphere @@ -1083,7 +1084,7 @@ pstd::optional MeasuredBxDF::Sample_f(Vector3f wo, Float uc, Point2f return BSDFSample(fr, wi, pdf * lum_pdf, BxDFFlags::GlossyReflection); } -Float MeasuredBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, +PBRT_CPU_GPU Float MeasuredBxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, BxDFReflTransFlags sampleFlags) const { if (!(sampleFlags & BxDFReflTransFlags::Reflection)) return 0; @@ -1127,7 +1128,7 @@ std::string NormalizedFresnelBxDF::ToString() const { } // BxDF Method Definitions -SampledSpectrum BxDF::rho(Vector3f wo, pstd::span uc, +PBRT_CPU_GPU SampledSpectrum BxDF::rho(Vector3f wo, pstd::span uc, pstd::span u2) const { if (wo.z == 0) return {}; diff --git a/src/pbrt/bxdfs.h b/src/pbrt/bxdfs.h index 0c72c9ddc..1871f1718 100644 --- a/src/pbrt/bxdfs.h +++ b/src/pbrt/bxdfs.h @@ -1131,12 +1131,12 @@ class NormalizedFresnelBxDF { Float eta; }; -inline SampledSpectrum BxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { +PBRT_CPU_GPU inline SampledSpectrum BxDF::f(Vector3f wo, Vector3f wi, TransportMode mode) const { auto f = [&](auto ptr) -> SampledSpectrum { return ptr->f(wo, wi, mode); }; return Dispatch(f); } -inline pstd::optional BxDF::Sample_f(Vector3f wo, Float uc, Point2f u, +PBRT_CPU_GPU inline pstd::optional BxDF::Sample_f(Vector3f wo, Float uc, Point2f u, TransportMode mode, BxDFReflTransFlags sampleFlags) const { auto sample_f = [&](auto ptr) -> pstd::optional { @@ -1145,18 +1145,18 @@ inline pstd::optional BxDF::Sample_f(Vector3f wo, Float uc, Point2f return Dispatch(sample_f); } -inline Float BxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, +PBRT_CPU_GPU inline Float BxDF::PDF(Vector3f wo, Vector3f wi, TransportMode mode, BxDFReflTransFlags sampleFlags) const { auto pdf = [&](auto ptr) { return ptr->PDF(wo, wi, mode, sampleFlags); }; return Dispatch(pdf); } -inline BxDFFlags BxDF::Flags() const { +PBRT_CPU_GPU inline BxDFFlags BxDF::Flags() const { auto flags = [&](auto ptr) { return ptr->Flags(); }; return Dispatch(flags); } -inline void BxDF::Regularize() { +PBRT_CPU_GPU inline void BxDF::Regularize() { auto regularize = [&](auto ptr) { ptr->Regularize(); }; return Dispatch(regularize); } diff --git a/src/pbrt/cameras.cpp b/src/pbrt/cameras.cpp index 20428ac88..6b8fb3c92 100644 --- a/src/pbrt/cameras.cpp +++ b/src/pbrt/cameras.cpp @@ -62,24 +62,24 @@ std::string CameraTransform::ToString() const { } // Camera Method Definitions -pstd::optional Camera::GenerateRayDifferential( +PBRT_CPU_GPU pstd::optional Camera::GenerateRayDifferential( CameraSample sample, SampledWavelengths &lambda) const { auto gen = [&](auto ptr) { return ptr->GenerateRayDifferential(sample, lambda); }; return Dispatch(gen); } -SampledSpectrum Camera::We(const Ray &ray, SampledWavelengths &lambda, +PBRT_CPU_GPU SampledSpectrum Camera::We(const Ray &ray, SampledWavelengths &lambda, Point2f *pRaster2) const { auto we = [&](auto ptr) { return ptr->We(ray, lambda, pRaster2); }; return Dispatch(we); } -void Camera::PDF_We(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void Camera::PDF_We(const Ray &ray, Float *pdfPos, Float *pdfDir) const { auto pdf = [&](auto ptr) { return ptr->PDF_We(ray, pdfPos, pdfDir); }; return Dispatch(pdf); } -pstd::optional Camera::SampleWi(const Interaction &ref, Point2f u, +PBRT_CPU_GPU pstd::optional Camera::SampleWi(const Interaction &ref, Point2f u, SampledWavelengths &lambda) const { auto sample = [&](auto ptr) { return ptr->SampleWi(ref, u, lambda); }; return Dispatch(sample); @@ -113,7 +113,7 @@ CameraBase::CameraBase(CameraBaseParameters p) "the system may crash as a result of this."); } -pstd::optional CameraBase::GenerateRayDifferential( +PBRT_CPU_GPU pstd::optional CameraBase::GenerateRayDifferential( Camera camera, CameraSample sample, SampledWavelengths &lambda) { // Generate regular camera ray _cr_ for ray differential pstd::optional cr = camera.GenerateRay(sample, lambda); @@ -281,7 +281,7 @@ CameraBaseParameters::CameraBaseParameters(const CameraTransform &cameraTransfor } // OrthographicCamera Method Definitions -pstd::optional OrthographicCamera::GenerateRay( +PBRT_CPU_GPU pstd::optional OrthographicCamera::GenerateRay( CameraSample sample, SampledWavelengths &lambda) const { // Compute raster and camera sample positions Point3f pFilm = Point3f(sample.pFilm.x, sample.pFilm.y, 0); @@ -305,7 +305,7 @@ pstd::optional OrthographicCamera::GenerateRay( return CameraRay{RenderFromCamera(ray)}; } -pstd::optional OrthographicCamera::GenerateRayDifferential( +PBRT_CPU_GPU pstd::optional OrthographicCamera::GenerateRayDifferential( CameraSample sample, SampledWavelengths &lambda) const { // Compute main orthographic viewing ray // Compute raster and camera sample positions @@ -401,7 +401,7 @@ OrthographicCamera *OrthographicCamera::Create(const ParameterDictionary ¶me } // PerspectiveCamera Method Definitions -pstd::optional PerspectiveCamera::GenerateRay( +PBRT_CPU_GPU pstd::optional PerspectiveCamera::GenerateRay( CameraSample sample, SampledWavelengths &lambda) const { // Compute raster and camera sample positions Point3f pFilm = Point3f(sample.pFilm.x, sample.pFilm.y, 0); @@ -426,7 +426,7 @@ pstd::optional PerspectiveCamera::GenerateRay( return CameraRay{RenderFromCamera(ray)}; } -pstd::optional PerspectiveCamera::GenerateRayDifferential( +PBRT_CPU_GPU pstd::optional PerspectiveCamera::GenerateRayDifferential( CameraSample sample, SampledWavelengths &lambda) const { // Compute raster and camera sample positions Point3f pFilm = Point3f(sample.pFilm.x, sample.pFilm.y, 0); @@ -527,7 +527,7 @@ PerspectiveCamera *PerspectiveCamera::Create(const ParameterDictionary ¶mete lensradius, focaldistance); } -SampledSpectrum PerspectiveCamera::We(const Ray &ray, SampledWavelengths &lambda, +PBRT_CPU_GPU SampledSpectrum PerspectiveCamera::We(const Ray &ray, SampledWavelengths &lambda, Point2f *pRasterOut) const { // Check if ray is forward-facing with respect to the camera Float cosTheta = Dot(ray.d, RenderFromCamera(Vector3f(0, 0, 1), ray.time)); @@ -555,7 +555,7 @@ SampledSpectrum PerspectiveCamera::We(const Ray &ray, SampledWavelengths &lambda return SampledSpectrum(1 / (A * lensArea * Pow<4>(cosTheta))); } -void PerspectiveCamera::PDF_We(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void PerspectiveCamera::PDF_We(const Ray &ray, Float *pdfPos, Float *pdfDir) const { // Return zero PDF values if ray direction is not front-facing Float cosTheta = Dot(ray.d, RenderFromCamera(Vector3f(0, 0, 1), ray.time)); if (cosTheta <= cosTotalWidth) { @@ -581,7 +581,7 @@ void PerspectiveCamera::PDF_We(const Ray &ray, Float *pdfPos, Float *pdfDir) con *pdfDir = 1 / (A * Pow<3>(cosTheta)); } -pstd::optional PerspectiveCamera::SampleWi( +PBRT_CPU_GPU pstd::optional PerspectiveCamera::SampleWi( const Interaction &ref, Point2f u, SampledWavelengths &lambda) const { // Uniformly sample a lens interaction _lensIntr_ Point2f pLens = lensRadius * SampleUniformDiskConcentric(u); @@ -607,7 +607,7 @@ pstd::optional PerspectiveCamera::SampleWi( } // SphericalCamera Method Definitions -pstd::optional SphericalCamera::GenerateRay(CameraSample sample, +PBRT_CPU_GPU pstd::optional SphericalCamera::GenerateRay(CameraSample sample, SampledWavelengths &lambda) const { // Compute spherical camera ray direction Point2f uv(sample.pFilm.x / film.FullResolution().x, @@ -746,7 +746,7 @@ RealisticCamera::RealisticCamera(CameraBaseParameters baseParameters, FindMinimumDifferentials(this); } -Float RealisticCamera::TraceLensesFromFilm(const Ray &rCamera, Ray *rOut) const { +PBRT_CPU_GPU Float RealisticCamera::TraceLensesFromFilm(const Ray &rCamera, Ray *rOut) const { Float elementZ = 0, weight = 1; // Transform _rCamera_ from camera to lens system space Ray rLens(Point3f(rCamera.o.x, rCamera.o.y, -rCamera.o.z), @@ -894,7 +894,7 @@ Bounds2f RealisticCamera::BoundExitPupil(Float filmX0, Float filmX1) const { return pupilBounds; } -pstd::optional RealisticCamera::SampleExitPupil(Point2f pFilm, +PBRT_CPU_GPU pstd::optional RealisticCamera::SampleExitPupil(Point2f pFilm, Point2f uLens) const { // Find exit pupil bound for sample distance from film center Float rFilm = std::sqrt(Sqr(pFilm.x) + Sqr(pFilm.y)); @@ -916,7 +916,7 @@ pstd::optional RealisticCamera::SampleExitPupil(Point2f pFilm, return ExitPupilSample{pPupil, pdf}; } -pstd::optional RealisticCamera::GenerateRay(CameraSample sample, +PBRT_CPU_GPU pstd::optional RealisticCamera::GenerateRay(CameraSample sample, SampledWavelengths &lambda) const { // Find point on film, _pFilm_, corresponding to _sample.pFilm_ Point2f s(sample.pFilm.x / film.FullResolution().x, @@ -956,7 +956,7 @@ std::string RealisticCamera::LensElementInterface::ToString() const { curvatureRadius, thickness, eta, apertureRadius); } -Float RealisticCamera::TraceLensesFromScene(const Ray &rCamera, Ray *rOut) const { +PBRT_CPU_GPU Float RealisticCamera::TraceLensesFromScene(const Ray &rCamera, Ray *rOut) const { Float elementZ = -LensFrontZ(); // Transform _rCamera_ from camera to lens system space const Transform LensFromCamera = Scale(1, 1, -1); diff --git a/src/pbrt/cameras.h b/src/pbrt/cameras.h index 73ffd67a5..53ecd2ef0 100644 --- a/src/pbrt/cameras.h +++ b/src/pbrt/cameras.h @@ -584,30 +584,30 @@ class RealisticCamera : public CameraBase { pstd::vector exitPupilBounds; }; -inline pstd::optional Camera::GenerateRay(CameraSample sample, +PBRT_CPU_GPU inline pstd::optional Camera::GenerateRay(CameraSample sample, SampledWavelengths &lambda) const { auto generate = [&](auto ptr) { return ptr->GenerateRay(sample, lambda); }; return Dispatch(generate); } -inline Film Camera::GetFilm() const { +PBRT_CPU_GPU inline Film Camera::GetFilm() const { auto getfilm = [&](auto ptr) { return ptr->GetFilm(); }; return Dispatch(getfilm); } -inline Float Camera::SampleTime(Float u) const { +PBRT_CPU_GPU inline Float Camera::SampleTime(Float u) const { auto sample = [&](auto ptr) { return ptr->SampleTime(u); }; return Dispatch(sample); } -inline const CameraTransform &Camera::GetCameraTransform() const { +PBRT_CPU_GPU inline const CameraTransform &Camera::GetCameraTransform() const { auto gtc = [&](auto ptr) -> const CameraTransform & { return ptr->GetCameraTransform(); }; return Dispatch(gtc); } -inline void Camera::Approximate_dp_dxy(Point3f p, Normal3f n, Float time, +PBRT_CPU_GPU inline void Camera::Approximate_dp_dxy(Point3f p, Normal3f n, Float time, int samplesPerPixel, Vector3f *dpdx, Vector3f *dpdy) const { if constexpr (AllInheritFrom(Camera::Types())) { diff --git a/src/pbrt/film.cpp b/src/pbrt/film.cpp index df9c9dcba..a5700a1c2 100644 --- a/src/pbrt/film.cpp +++ b/src/pbrt/film.cpp @@ -34,7 +34,7 @@ namespace pbrt { -void Film::AddSplat(Point2f p, SampledSpectrum v, const SampledWavelengths &lambda) { +PBRT_CPU_GPU void Film::AddSplat(Point2f p, SampledSpectrum v, const SampledWavelengths &lambda) { auto splat = [&](auto ptr) { return ptr->AddSplat(p, v, lambda); }; return Dispatch(splat); } @@ -175,7 +175,7 @@ FilmBaseParameters::FilmBaseParameters(const ParameterDictionary ¶meters, } // FilmBase Method Definitions -Bounds2f FilmBase::SampleBounds() const { +PBRT_CPU_GPU Bounds2f FilmBase::SampleBounds() const { Vector2f radius = filter.Radius(); return Bounds2f(pixelBounds.pMin - radius + Vector2f(0.5f, 0.5f), pixelBounds.pMax + radius - Vector2f(0.5f, 0.5f)); @@ -188,7 +188,7 @@ std::string FilmBase::BaseToString() const { } // VisibleSurface Method Definitions -VisibleSurface::VisibleSurface(const SurfaceInteraction &si, SampledSpectrum albedo, +PBRT_CPU_GPU VisibleSurface::VisibleSurface(const SurfaceInteraction &si, SampledSpectrum albedo, const SampledWavelengths &lambda) : albedo(albedo) { set = true; @@ -496,7 +496,7 @@ RGBFilm::RGBFilm(FilmBaseParameters p, const RGBColorSpace *colorSpace, outputRGBFromSensorRGB = colorSpace->RGBFromXYZ * sensor->XYZFromSensorRGB; } -void RGBFilm::AddSplat(Point2f p, SampledSpectrum L, const SampledWavelengths &lambda) { +PBRT_CPU_GPU void RGBFilm::AddSplat(Point2f p, SampledSpectrum L, const SampledWavelengths &lambda) { CHECK(!L.HasNaNs()); // Convert sample radiance to _PixelSensor_ RGB RGB rgb = sensor->ToSensorRGB(L, lambda); @@ -585,7 +585,7 @@ RGBFilm *RGBFilm::Create(const ParameterDictionary ¶meters, Float exposureTi } // GBufferFilm Method Definitions -void GBufferFilm::AddSample(Point2i pFilm, SampledSpectrum L, +PBRT_CPU_GPU void GBufferFilm::AddSample(Point2i pFilm, SampledSpectrum L, const SampledWavelengths &lambda, const VisibleSurface *visibleSurface, Float weight) { RGB rgb = sensor->ToSensorRGB(L, lambda); @@ -656,7 +656,7 @@ GBufferFilm::GBufferFilm(FilmBaseParameters p, const AnimatedTransform &outputFr outputRGBFromSensorRGB = colorSpace->RGBFromXYZ * sensor->XYZFromSensorRGB; } -void GBufferFilm::AddSplat(Point2f p, SampledSpectrum v, +PBRT_CPU_GPU void GBufferFilm::AddSplat(Point2f p, SampledSpectrum v, const SampledWavelengths &lambda) { // NOTE: same code as RGBFilm::AddSplat()... CHECK(!v.HasNaNs()); @@ -886,7 +886,7 @@ SpectralFilm::SpectralFilm(FilmBaseParameters p, Float lambdaMin, Float lambdaMa } } -RGB SpectralFilm::GetPixelRGB(Point2i p, Float splatScale) const { +PBRT_CPU_GPU RGB SpectralFilm::GetPixelRGB(Point2i p, Float splatScale) const { // Note: this is effectively the same as RGBFilm::GetPixelRGB const Pixel &pixel = pixels[p]; @@ -906,7 +906,7 @@ RGB SpectralFilm::GetPixelRGB(Point2i p, Float splatScale) const { return rgb; } -void SpectralFilm::AddSplat(Point2f p, SampledSpectrum L, +PBRT_CPU_GPU void SpectralFilm::AddSplat(Point2f p, SampledSpectrum L, const SampledWavelengths &lambda) { // This, too, is similar to RGBFilm::AddSplat(), with additions for // spectra. diff --git a/src/pbrt/film.h b/src/pbrt/film.h index 9f964e2fc..1476ccbc5 100644 --- a/src/pbrt/film.h +++ b/src/pbrt/film.h @@ -295,7 +295,7 @@ class RGBFilm : public FilmBase { return outputRGBFromSensorRGB * sensorRGB; } - PBRT_CPU_GPU void ResetPixel(Point2i p) { std::memset(&pixels[p], 0, sizeof(Pixel)); } + PBRT_CPU_GPU void ResetPixel(Point2i p) { memset(&pixels[p], 0, sizeof(Pixel)); } private: // RGBFilm::Pixel Definition @@ -369,7 +369,7 @@ class GBufferFilm : public FilmBase { std::string ToString() const; - PBRT_CPU_GPU void ResetPixel(Point2i p) { std::memset(&pixels[p], 0, sizeof(Pixel)); } + PBRT_CPU_GPU void ResetPixel(Point2i p) { memset(&pixels[p], 0, sizeof(Pixel)); } private: // GBufferFilm::Pixel Definition @@ -490,9 +490,9 @@ class SpectralFilm : public FilmBase { pix.rgbSum[0] = pix.rgbSum[1] = pix.rgbSum[2] = 0.; pix.rgbWeightSum = 0.; pix.rgbSplat[0] = pix.rgbSplat[1] = pix.rgbSplat[2] = 0.; - std::memset(pix.bucketSums, 0, nBuckets * sizeof(double)); - std::memset(pix.weightSums, 0, nBuckets * sizeof(double)); - std::memset(pix.bucketSplats, 0, nBuckets * sizeof(AtomicDouble)); + memset(pix.bucketSums, 0, nBuckets * sizeof(double)); + memset(pix.weightSums, 0, nBuckets * sizeof(double)); + memset(pix.bucketSplats, 0, nBuckets * sizeof(AtomicDouble)); } private: diff --git a/src/pbrt/filters.cpp b/src/pbrt/filters.cpp index c34bf0eec..cf612beba 100644 --- a/src/pbrt/filters.cpp +++ b/src/pbrt/filters.cpp @@ -63,7 +63,7 @@ MitchellFilter *MitchellFilter::Create(const ParameterDictionary ¶meters, } // Sinc Filter Method Definitions -Float LanczosSincFilter::Integral() const { +PBRT_CPU_GPU Float LanczosSincFilter::Integral() const { Float sum = 0; int sqrtSamples = 64; int nSamples = sqrtSamples * sqrtSamples; diff --git a/src/pbrt/filters.h b/src/pbrt/filters.h index ce3a50caa..af6b97ba3 100644 --- a/src/pbrt/filters.h +++ b/src/pbrt/filters.h @@ -233,22 +233,22 @@ class TriangleFilter { Vector2f radius; }; -inline Float Filter::Evaluate(Point2f p) const { +PBRT_CPU_GPU inline Float Filter::Evaluate(Point2f p) const { auto eval = [&](auto ptr) { return ptr->Evaluate(p); }; return Dispatch(eval); } -inline FilterSample Filter::Sample(Point2f u) const { +PBRT_CPU_GPU inline FilterSample Filter::Sample(Point2f u) const { auto sample = [&](auto ptr) { return ptr->Sample(u); }; return Dispatch(sample); } -inline Vector2f Filter::Radius() const { +PBRT_CPU_GPU inline Vector2f Filter::Radius() const { auto radius = [&](auto ptr) { return ptr->Radius(); }; return Dispatch(radius); } -inline Float Filter::Integral() const { +PBRT_CPU_GPU inline Float Filter::Integral() const { auto integral = [&](auto ptr) { return ptr->Integral(); }; return Dispatch(integral); } diff --git a/src/pbrt/interaction.cpp b/src/pbrt/interaction.cpp index 49ab2b20b..a0c7ecfcb 100644 --- a/src/pbrt/interaction.cpp +++ b/src/pbrt/interaction.cpp @@ -38,7 +38,7 @@ std::string MediumInteraction::ToString() const { } // SurfaceInteraction Method Definitions -void SurfaceInteraction::ComputeDifferentials(const RayDifferential &ray, Camera camera, +PBRT_CPU_GPU void SurfaceInteraction::ComputeDifferentials(const RayDifferential &ray, Camera camera, int samplesPerPixel) { if (GetOptions().disableTextureFiltering) { dudx = dudy = dvdx = dvdy = 0; @@ -88,7 +88,7 @@ void SurfaceInteraction::ComputeDifferentials(const RayDifferential &ray, Camera dvdy = IsFinite(dvdy) ? Clamp(dvdy, -1e8f, 1e8f) : 0.f; } -void SurfaceInteraction::SkipIntersection(RayDifferential *ray, Float t) const { +PBRT_CPU_GPU void SurfaceInteraction::SkipIntersection(RayDifferential *ray, Float t) const { *((Ray *)ray) = SpawnRay(ray->d); if (ray->hasDifferentials) { ray->rxOrigin = ray->rxOrigin + t * ray->rxDirection; @@ -96,7 +96,7 @@ void SurfaceInteraction::SkipIntersection(RayDifferential *ray, Float t) const { } } -RayDifferential SurfaceInteraction::SpawnRay(const RayDifferential &rayi, +PBRT_CPU_GPU RayDifferential SurfaceInteraction::SpawnRay(const RayDifferential &rayi, const BSDF &bsdf, Vector3f wi, int flags, Float eta) const { RayDifferential rd(SpawnRay(wi)); @@ -210,7 +210,7 @@ BSSRDF SurfaceInteraction::GetBSSRDF(const RayDifferential &ray, return material.GetBSSRDF(UniversalTextureEvaluator(), *this, lambda, scratchBuffer); } -SampledSpectrum SurfaceInteraction::Le(Vector3f w, +PBRT_CPU_GPU SampledSpectrum SurfaceInteraction::Le(Vector3f w, const SampledWavelengths &lambda) const { return areaLight ? areaLight.L(p(), n, uv, w, lambda) : SampledSpectrum(0.f); } diff --git a/src/pbrt/lights.cpp b/src/pbrt/lights.cpp index b66128f99..a182ccdc0 100644 --- a/src/pbrt/lights.cpp +++ b/src/pbrt/lights.cpp @@ -104,7 +104,7 @@ std::string LightBounds::ToString() const { } // LightBounds Method Definitions -Float LightBounds::Importance(Point3f p, Normal3f n) const { +PBRT_CPU_GPU Float LightBounds::Importance(Point3f p, Normal3f n) const { // Return importance for light bounds at reference point // Compute clamped squared distance to reference point Point3f pc = (bounds.pMin + bounds.pMax) / 2; @@ -172,7 +172,7 @@ pstd::optional PointLight::Bounds() const { std::cos(Pi / 2), false); } -pstd::optional PointLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional PointLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { Point3f p = renderFromLight(Point3f(0, 0, 0)); @@ -180,7 +180,7 @@ pstd::optional PointLight::SampleLe(Point2f u1, Point2f u2, return LightLeSample(scale * I->Sample(lambda), ray, 1, UniformSpherePDF()); } -void PointLight::PDF_Le(const Ray &, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void PointLight::PDF_Le(const Ray &, Float *pdfPos, Float *pdfDir) const { *pdfPos = 0; *pdfDir = UniformSpherePDF(); } @@ -217,7 +217,7 @@ SampledSpectrum DistantLight::Phi(SampledWavelengths lambda) const { return scale * Lemit->Sample(lambda) * Pi * Sqr(sceneRadius); } -pstd::optional DistantLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional DistantLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Choose point on disk oriented toward infinite light direction @@ -233,7 +233,7 @@ pstd::optional DistantLight::SampleLe(Point2f u1, Point2f u2, 1); } -void DistantLight::PDF_Le(const Ray &, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void DistantLight::PDF_Le(const Ray &, Float *pdfPos, Float *pdfDir) const { *pdfPos = 1 / (Pi * sceneRadius * sceneRadius); *pdfDir = 0; } @@ -317,7 +317,7 @@ ProjectionLight::ProjectionLight(Transform renderFromLight, imageBytes += image.BytesUsed() + distrib.BytesUsed(); } -pstd::optional ProjectionLight::SampleLi(LightSampleContext ctx, Point2f u, +PBRT_CPU_GPU pstd::optional ProjectionLight::SampleLi(LightSampleContext ctx, Point2f u, SampledWavelengths lambda, bool allowIncompletePDF) const { // Return sample for incident radiance from _ProjectionLight_ @@ -330,7 +330,7 @@ pstd::optional ProjectionLight::SampleLi(LightSampleContext ctx, return LightLiSample(Li, wi, 1, Interaction(p, &mediumInterface)); } -Float ProjectionLight::PDF_Li(LightSampleContext, Vector3f, +PBRT_CPU_GPU Float ProjectionLight::PDF_Li(LightSampleContext, Vector3f, bool allowIncompletePDF) const { return 0.f; } @@ -340,7 +340,7 @@ std::string ProjectionLight::ToString() const { A); } -SampledSpectrum ProjectionLight::I(Vector3f w, const SampledWavelengths &lambda) const { +PBRT_CPU_GPU SampledSpectrum ProjectionLight::I(Vector3f w, const SampledWavelengths &lambda) const { // Discard directions behind projection light if (w.z < hither) return SampledSpectrum(0.f); @@ -398,7 +398,7 @@ pstd::optional ProjectionLight::Bounds() const { return LightBounds(Bounds3f(p, p), w, phi, std::cos(0.f), cosTotalWidth, false); } -pstd::optional ProjectionLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional ProjectionLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Sample light space ray direction for projection light @@ -425,7 +425,7 @@ pstd::optional ProjectionLight::SampleLe(Point2f u1, Point2f u2, return LightLeSample(L, ray, 1, pdfDir); } -void ProjectionLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void ProjectionLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { *pdfPos = 0; // Transform ray direction to light space and reject invalid ones Vector3f w = Normalize(renderFromLight.ApplyInverse(ray.d)); @@ -535,7 +535,7 @@ GoniometricLight::GoniometricLight(const Transform &renderFromLight, imageBytes += image.BytesUsed() + distrib.BytesUsed(); } -pstd::optional GoniometricLight::SampleLi(LightSampleContext ctx, +PBRT_CPU_GPU pstd::optional GoniometricLight::SampleLi(LightSampleContext ctx, Point2f u, SampledWavelengths lambda, bool allowIncompletePDF) const { @@ -546,7 +546,7 @@ pstd::optional GoniometricLight::SampleLi(LightSampleContext ctx, return LightLiSample(L, wi, 1, Interaction(p, &mediumInterface)); } -Float GoniometricLight::PDF_Li(LightSampleContext, Vector3f, +PBRT_CPU_GPU Float GoniometricLight::PDF_Li(LightSampleContext, Vector3f, bool allowIncompletePDF) const { return 0.f; } @@ -574,7 +574,7 @@ pstd::optional GoniometricLight::Bounds() const { std::cos(Pi / 2), false); } -pstd::optional GoniometricLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional GoniometricLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Sample direction and PDF for ray leaving goniometric light @@ -588,7 +588,7 @@ pstd::optional GoniometricLight::SampleLe(Point2f u1, Point2f u2, return LightLeSample(I(wLight, lambda), ray, 1, pdfDir); } -void GoniometricLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void GoniometricLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { *pdfPos = 0.f; Vector3f wLight = Normalize(renderFromLight.ApplyInverse(ray.d)); Point2f uv = EqualAreaSphereToSquare(wLight); @@ -736,7 +736,7 @@ DiffuseAreaLight::DiffuseAreaLight(const Transform &renderFromLight, "Proceed at your own risk; your image may have errors."); } -pstd::optional DiffuseAreaLight::SampleLi(LightSampleContext ctx, +PBRT_CPU_GPU pstd::optional DiffuseAreaLight::SampleLi(LightSampleContext ctx, Point2f u, SampledWavelengths lambda, bool allowIncompletePDF) const { @@ -760,7 +760,7 @@ pstd::optional DiffuseAreaLight::SampleLi(LightSampleContext ctx, return LightLiSample(Le, wi, ss->pdf, ss->intr); } -Float DiffuseAreaLight::PDF_Li(LightSampleContext ctx, Vector3f wi, +PBRT_CPU_GPU Float DiffuseAreaLight::PDF_Li(LightSampleContext ctx, Vector3f wi, bool allowIncompletePDF) const { ShapeSampleContext shapeCtx(ctx.pi, ctx.n, ctx.ns, 0 /* time */); return shape.PDF(shapeCtx, wi); @@ -806,7 +806,7 @@ pstd::optional DiffuseAreaLight::Bounds() const { twoSided); } -pstd::optional DiffuseAreaLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional DiffuseAreaLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Sample a point on the area light's _Shape_ @@ -850,7 +850,7 @@ pstd::optional DiffuseAreaLight::SampleLe(Point2f u1, Point2f u2, return LightLeSample(Le, intr.SpawnRay(w), intr, ss->pdf, pdfDir); } -void DiffuseAreaLight::PDF_Le(const Interaction &intr, Vector3f w, Float *pdfPos, +PBRT_CPU_GPU void DiffuseAreaLight::PDF_Le(const Interaction &intr, Vector3f w, Float *pdfPos, Float *pdfDir) const { CHECK_NE(intr.n, Normal3f(0, 0, 0)); *pdfPos = shape.PDF(intr); @@ -947,12 +947,12 @@ UniformInfiniteLight::UniformInfiniteLight(const Transform &renderFromLight, Lemit(LookupSpectrum(Lemit)), scale(scale) {} -SampledSpectrum UniformInfiniteLight::Le(const Ray &ray, +PBRT_CPU_GPU SampledSpectrum UniformInfiniteLight::Le(const Ray &ray, const SampledWavelengths &lambda) const { return scale * Lemit->Sample(lambda); } -pstd::optional UniformInfiniteLight::SampleLi( +PBRT_CPU_GPU pstd::optional UniformInfiniteLight::SampleLi( LightSampleContext ctx, Point2f u, SampledWavelengths lambda, bool allowIncompletePDF) const { if (allowIncompletePDF) @@ -964,7 +964,7 @@ pstd::optional UniformInfiniteLight::SampleLi( Interaction(ctx.p() + wi * (2 * sceneRadius), &mediumInterface)); } -Float UniformInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, +PBRT_CPU_GPU Float UniformInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, bool allowIncompletePDF) const { if (allowIncompletePDF) return 0; @@ -975,7 +975,7 @@ SampledSpectrum UniformInfiniteLight::Phi(SampledWavelengths lambda) const { return 4 * Pi * Pi * Sqr(sceneRadius) * scale * Lemit->Sample(lambda); } -pstd::optional UniformInfiniteLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional UniformInfiniteLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Sample direction for uniform infinite light ray @@ -994,7 +994,7 @@ pstd::optional UniformInfiniteLight::SampleLe(Point2f u1, Point2f return LightLeSample(scale * Lemit->Sample(lambda), ray, pdfPos, pdfDir); } -void UniformInfiniteLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void UniformInfiniteLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { *pdfDir = UniformSpherePDF(); *pdfPos = 1 / (Pi * Sqr(sceneRadius)); } @@ -1039,7 +1039,7 @@ ImageInfiniteLight::ImageInfiniteLight(Transform renderFromLight, Image im, compensatedDistribution = PiecewiseConstant2D(d, domain, alloc); } -Float ImageInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, +PBRT_CPU_GPU Float ImageInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, bool allowIncompletePDF) const { Vector3f wLight = renderFromLight.ApplyInverse(w); Point2f uv = EqualAreaSphereToSquare(wLight); @@ -1070,7 +1070,7 @@ SampledSpectrum ImageInfiniteLight::Phi(SampledWavelengths lambda) const { return 4 * Pi * Pi * Sqr(sceneRadius) * scale * sumL / (width * height); } -pstd::optional ImageInfiniteLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional ImageInfiniteLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Sample infinite light image and compute ray direction _w_ @@ -1094,7 +1094,7 @@ pstd::optional ImageInfiniteLight::SampleLe(Point2f u1, Point2f u return LightLeSample(ImageLe(*uv, lambda), ray, pdfPos, pdfDir); } -void ImageInfiniteLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void ImageInfiniteLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { Vector3f wl = -renderFromLight.ApplyInverse(ray.d); Float mapPDF = distribution.PDF(EqualAreaSphereToSquare(wl)); *pdfDir = mapPDF / (4 * Pi); @@ -1205,7 +1205,7 @@ SampledSpectrum PortalImageInfiniteLight::Phi(SampledWavelengths lambda) const { return scale * Area() * sumL / (image.Resolution().x * image.Resolution().y); } -SampledSpectrum PortalImageInfiniteLight::Le(const Ray &ray, +PBRT_CPU_GPU SampledSpectrum PortalImageInfiniteLight::Le(const Ray &ray, const SampledWavelengths &lambda) const { pstd::optional uv = ImageFromRender(Normalize(ray.d)); pstd::optional b = ImageBounds(ray.o); @@ -1214,7 +1214,7 @@ SampledSpectrum PortalImageInfiniteLight::Le(const Ray &ray, return ImageLookup(*uv, lambda); } -SampledSpectrum PortalImageInfiniteLight::ImageLookup( +PBRT_CPU_GPU SampledSpectrum PortalImageInfiniteLight::ImageLookup( Point2f uv, const SampledWavelengths &lambda) const { RGB rgb; for (int c = 0; c < 3; ++c) @@ -1223,7 +1223,7 @@ SampledSpectrum PortalImageInfiniteLight::ImageLookup( return scale * spec.Sample(lambda); } -pstd::optional PortalImageInfiniteLight::SampleLi( +PBRT_CPU_GPU pstd::optional PortalImageInfiniteLight::SampleLi( LightSampleContext ctx, Point2f u, SampledWavelengths lambda, bool allowIncompletePDF) const { // Sample $(u,v)$ in potentially visible region of light image @@ -1249,7 +1249,7 @@ pstd::optional PortalImageInfiniteLight::SampleLi( return LightLiSample(L, wi, pdf, Interaction(pl, &mediumInterface)); } -Float PortalImageInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, +PBRT_CPU_GPU Float PortalImageInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, bool allowIncompletePDF) const { // Find image $(u,v)$ coordinates corresponding to direction _w_ Float duv_dw; @@ -1265,7 +1265,7 @@ Float PortalImageInfiniteLight::PDF_Li(LightSampleContext ctx, Vector3f w, return pdf / duv_dw; } -pstd::optional PortalImageInfiniteLight::SampleLe( +PBRT_CPU_GPU pstd::optional PortalImageInfiniteLight::SampleLe( Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { Float mapPDF; Bounds2f b(Point2f(0, 0), Point2f(1, 1)); @@ -1310,7 +1310,7 @@ pstd::optional PortalImageInfiniteLight::SampleLe( return LightLeSample(L, ray, pdfPos, pdfDir); } -void PortalImageInfiniteLight::PDF_Le(const Ray &ray, Float *pdfPos, +PBRT_CPU_GPU void PortalImageInfiniteLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { // TODO: negate here or??? Vector3f w = -Normalize(ray.d); @@ -1353,11 +1353,11 @@ SpotLight::SpotLight(const Transform &renderFromLight, CHECK_LE(falloffStart, totalWidth); } -Float SpotLight::PDF_Li(LightSampleContext, Vector3f, bool allowIncompletePDF) const { +PBRT_CPU_GPU Float SpotLight::PDF_Li(LightSampleContext, Vector3f, bool allowIncompletePDF) const { return 0.f; } -SampledSpectrum SpotLight::I(Vector3f w, SampledWavelengths lambda) const { +PBRT_CPU_GPU SampledSpectrum SpotLight::I(Vector3f w, SampledWavelengths lambda) const { return SmoothStep(CosTheta(w), cosFalloffEnd, cosFalloffStart) * scale * Iemit->Sample(lambda); } @@ -1379,7 +1379,7 @@ pstd::optional SpotLight::Bounds() const { return LightBounds(Bounds3f(p, p), w, phi, cosFalloffStart, cosTheta_e, false); } -pstd::optional SpotLight::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional SpotLight::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { // Choose whether to sample spotlight center cone or falloff region @@ -1412,7 +1412,7 @@ pstd::optional SpotLight::SampleLe(Point2f u1, Point2f u2, return LightLeSample(I(wLight, lambda), ray, 1, pdfDir); } -void SpotLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void SpotLight::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { Float p[2] = {1 - cosFalloffStart, (cosFalloffStart - cosFalloffEnd) / 2}; *pdfPos = 0; // Find spotlight directional PDF based on $\cos \theta$ @@ -1473,14 +1473,14 @@ void Light::Preprocess(const Bounds3f &sceneBounds) { return DispatchCPU(preprocess); } -pstd::optional Light::SampleLe(Point2f u1, Point2f u2, +PBRT_CPU_GPU pstd::optional Light::SampleLe(Point2f u1, Point2f u2, SampledWavelengths &lambda, Float time) const { auto sample = [&](auto ptr) { return ptr->SampleLe(u1, u2, lambda, time); }; return Dispatch(sample); } -void Light::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { +PBRT_CPU_GPU void Light::PDF_Le(const Ray &ray, Float *pdfPos, Float *pdfDir) const { auto pdf = [&](auto ptr) { return ptr->PDF_Le(ray, pdfPos, pdfDir); }; return Dispatch(pdf); } @@ -1498,7 +1498,7 @@ std::string Light::ToString() const { return DispatchCPU(str); } -void Light::PDF_Le(const Interaction &intr, Vector3f w, Float *pdfPos, +PBRT_CPU_GPU void Light::PDF_Le(const Interaction &intr, Vector3f w, Float *pdfPos, Float *pdfDir) const { auto pdf = [&](auto ptr) { return ptr->PDF_Le(intr, w, pdfPos, pdfDir); }; return Dispatch(pdf); diff --git a/src/pbrt/lights.h b/src/pbrt/lights.h index e6cce0555..76caca782 100644 --- a/src/pbrt/lights.h +++ b/src/pbrt/lights.h @@ -788,7 +788,7 @@ class SpotLight : public LightBase { Float scale, cosFalloffStart, cosFalloffEnd; }; -inline pstd::optional Light::SampleLi(LightSampleContext ctx, Point2f u, +PBRT_CPU_GPU inline pstd::optional Light::SampleLi(LightSampleContext ctx, Point2f u, SampledWavelengths lambda, bool allowIncompletePDF) const { auto sample = [&](auto ptr) { @@ -797,25 +797,26 @@ inline pstd::optional Light::SampleLi(LightSampleContext ctx, Poi return Dispatch(sample); } -inline Float Light::PDF_Li(LightSampleContext ctx, Vector3f wi, +PBRT_CPU_GPU inline Float Light::PDF_Li(LightSampleContext ctx, Vector3f wi, bool allowIncompletePDF) const { auto pdf = [&](auto ptr) { return ptr->PDF_Li(ctx, wi, allowIncompletePDF); }; return Dispatch(pdf); } -inline SampledSpectrum Light::L(Point3f p, Normal3f n, Point2f uv, Vector3f w, +PBRT_CPU_GPU inline SampledSpectrum Light::L(Point3f p, Normal3f n, Point2f uv, Vector3f w, const SampledWavelengths &lambda) const { CHECK(Type() == LightType::Area); auto l = [&](auto ptr) { return ptr->L(p, n, uv, w, lambda); }; return Dispatch(l); } -inline SampledSpectrum Light::Le(const Ray &ray, const SampledWavelengths &lambda) const { +PBRT_CPU_GPU inline SampledSpectrum Light::Le(const Ray &ray, + const SampledWavelengths &lambda) const { auto le = [&](auto ptr) { return ptr->Le(ray, lambda); }; return Dispatch(le); } -inline LightType Light::Type() const { +PBRT_CPU_GPU inline LightType Light::Type() const { auto t = [&](auto ptr) { return ptr->Type(); }; return Dispatch(t); } diff --git a/src/pbrt/lightsamplers.cpp b/src/pbrt/lightsamplers.cpp index af6fbd909..1592d35aa 100644 --- a/src/pbrt/lightsamplers.cpp +++ b/src/pbrt/lightsamplers.cpp @@ -265,7 +265,7 @@ ExhaustiveLightSampler::ExhaustiveLightSampler(pstd::span lights, } } -pstd::optional ExhaustiveLightSampler::Sample(const LightSampleContext &ctx, +PBRT_CPU_GPU pstd::optional ExhaustiveLightSampler::Sample(const LightSampleContext &ctx, Float u) const { Float pInfinite = Float(infiniteLights.size()) / Float(infiniteLights.size() + (!lightBounds.empty() ? 1 : 0)); @@ -293,7 +293,7 @@ pstd::optional ExhaustiveLightSampler::Sample(const LightSampleCon } } -Float ExhaustiveLightSampler::PMF(const LightSampleContext &ctx, Light light) const { +PBRT_CPU_GPU Float ExhaustiveLightSampler::PMF(const LightSampleContext &ctx, Light light) const { if (!lightToBoundedIndex.HasKey(light)) return 1.f / (infiniteLights.size() + (!lightBounds.empty() ? 1 : 0)); diff --git a/src/pbrt/lightsamplers.h b/src/pbrt/lightsamplers.h index e98e4bdca..d572de4db 100644 --- a/src/pbrt/lightsamplers.h +++ b/src/pbrt/lightsamplers.h @@ -438,23 +438,23 @@ class ExhaustiveLightSampler { HashMap lightToBoundedIndex; }; -inline pstd::optional LightSampler::Sample(const LightSampleContext &ctx, +PBRT_CPU_GPU inline pstd::optional LightSampler::Sample(const LightSampleContext &ctx, Float u) const { auto s = [&](auto ptr) { return ptr->Sample(ctx, u); }; return Dispatch(s); } -inline Float LightSampler::PMF(const LightSampleContext &ctx, Light light) const { +PBRT_CPU_GPU inline Float LightSampler::PMF(const LightSampleContext &ctx, Light light) const { auto pdf = [&](auto ptr) { return ptr->PMF(ctx, light); }; return Dispatch(pdf); } -inline pstd::optional LightSampler::Sample(Float u) const { +PBRT_CPU_GPU inline pstd::optional LightSampler::Sample(Float u) const { auto sample = [&](auto ptr) { return ptr->Sample(u); }; return Dispatch(sample); } -inline Float LightSampler::PMF(Light light) const { +PBRT_CPU_GPU inline Float LightSampler::PMF(Light light) const { auto pdf = [&](auto ptr) { return ptr->PMF(light); }; return Dispatch(pdf); } diff --git a/src/pbrt/materials.cpp b/src/pbrt/materials.cpp index 20c87b618..450b98222 100644 --- a/src/pbrt/materials.cpp +++ b/src/pbrt/materials.cpp @@ -252,7 +252,7 @@ ConductorMaterial *ConductorMaterial::Create(const TextureParameterDictionary &p // CoatedDiffuseMaterial Method Definitions template -CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF(TextureEvaluator texEval, +PBRT_CPU_GPU CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF(TextureEvaluator texEval, const MaterialEvalContext &ctx, SampledWavelengths &lambda) const { // Initialize diffuse component of plastic material @@ -283,10 +283,10 @@ CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF(TextureEvaluator texEval, } // Explicit template instantiation -template CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF( +template PBRT_CPU_GPU CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF( BasicTextureEvaluator, const MaterialEvalContext &ctx, SampledWavelengths &lambda) const; -template CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF( +template PBRT_CPU_GPU CoatedDiffuseBxDF CoatedDiffuseMaterial::GetBxDF( UniversalTextureEvaluator, const MaterialEvalContext &ctx, SampledWavelengths &lambda) const; @@ -343,7 +343,7 @@ CoatedDiffuseMaterial *CoatedDiffuseMaterial::Create( } template -CoatedConductorBxDF CoatedConductorMaterial::GetBxDF(TextureEvaluator texEval, +PBRT_CPU_GPU CoatedConductorBxDF CoatedConductorMaterial::GetBxDF(TextureEvaluator texEval, const MaterialEvalContext &ctx, SampledWavelengths &lambda) const { Float iurough = texEval(interfaceURoughness, ctx); @@ -391,10 +391,10 @@ CoatedConductorBxDF CoatedConductorMaterial::GetBxDF(TextureEvaluator texEval, maxDepth, nSamples); } -template CoatedConductorBxDF CoatedConductorMaterial::GetBxDF( +template PBRT_CPU_GPU CoatedConductorBxDF CoatedConductorMaterial::GetBxDF( BasicTextureEvaluator, const MaterialEvalContext &ctx, SampledWavelengths &lambda) const; -template CoatedConductorBxDF CoatedConductorMaterial::GetBxDF( +template PBRT_CPU_GPU CoatedConductorBxDF CoatedConductorMaterial::GetBxDF( UniversalTextureEvaluator, const MaterialEvalContext &ctx, SampledWavelengths &lambda) const; diff --git a/src/pbrt/materials.h b/src/pbrt/materials.h index 5cc1df28c..5a86cf1a9 100644 --- a/src/pbrt/materials.h +++ b/src/pbrt/materials.h @@ -914,7 +914,7 @@ inline BSDF Material::GetBSDF(TextureEvaluator texEval, MaterialEvalContext ctx, } template -inline bool Material::CanEvaluateTextures(TextureEvaluator texEval) const { +PBRT_CPU_GPU inline bool Material::CanEvaluateTextures(TextureEvaluator texEval) const { auto eval = [&](auto ptr) { return ptr->CanEvaluateTextures(texEval); }; return Dispatch(eval); } @@ -937,17 +937,17 @@ inline BSSRDF Material::GetBSSRDF(TextureEvaluator texEval, MaterialEvalContext return DispatchCPU(get); } -inline bool Material::HasSubsurfaceScattering() const { +PBRT_CPU_GPU inline bool Material::HasSubsurfaceScattering() const { auto has = [&](auto ptr) { return ptr->HasSubsurfaceScattering(); }; return Dispatch(has); } -inline FloatTexture Material::GetDisplacement() const { +PBRT_CPU_GPU inline FloatTexture Material::GetDisplacement() const { auto disp = [&](auto ptr) { return ptr->GetDisplacement(); }; return Dispatch(disp); } -inline const Image *Material::GetNormalMap() const { +PBRT_CPU_GPU inline const Image *Material::GetNormalMap() const { auto nmap = [&](auto ptr) { return ptr->GetNormalMap(); }; return Dispatch(nmap); } diff --git a/src/pbrt/media.h b/src/pbrt/media.h index 19abb9873..8a5677a74 100644 --- a/src/pbrt/media.h +++ b/src/pbrt/media.h @@ -24,7 +24,7 @@ #include #include #include -#ifdef PBRT_BUILD_GPU_RENDERER +#if defined(PBRT_BUILD_GPU_RENDERER) && defined(__NVCC__) #include #endif // PBRT_BUILD_GPU_RENDERER @@ -678,28 +678,28 @@ class NanoVDBMedium { Float LeScale, temperatureOffset, temperatureScale; }; -inline Float PhaseFunction::p(Vector3f wo, Vector3f wi) const { +PBRT_CPU_GPU inline Float PhaseFunction::p(Vector3f wo, Vector3f wi) const { auto p = [&](auto ptr) { return ptr->p(wo, wi); }; return Dispatch(p); } -inline pstd::optional PhaseFunction::Sample_p(Vector3f wo, +PBRT_CPU_GPU inline pstd::optional PhaseFunction::Sample_p(Vector3f wo, Point2f u) const { auto sample = [&](auto ptr) { return ptr->Sample_p(wo, u); }; return Dispatch(sample); } -inline Float PhaseFunction::PDF(Vector3f wo, Vector3f wi) const { +PBRT_CPU_GPU inline Float PhaseFunction::PDF(Vector3f wo, Vector3f wi) const { auto pdf = [&](auto ptr) { return ptr->PDF(wo, wi); }; return Dispatch(pdf); } -inline pstd::optional RayMajorantIterator::Next() { +PBRT_CPU_GPU inline pstd::optional RayMajorantIterator::Next() { auto next = [](auto ptr) { return ptr->Next(); }; return Dispatch(next); } -inline MediumProperties Medium::SamplePoint(Point3f p, +PBRT_CPU_GPU inline MediumProperties Medium::SamplePoint(Point3f p, const SampledWavelengths &lambda) const { auto sample = [&](auto ptr) { return ptr->SamplePoint(p, lambda); }; return Dispatch(sample); diff --git a/src/pbrt/samplers.cpp b/src/pbrt/samplers.cpp index 5953f1711..b103d789d 100644 --- a/src/pbrt/samplers.cpp +++ b/src/pbrt/samplers.cpp @@ -319,16 +319,16 @@ StratifiedSampler *StratifiedSampler::Create(const ParameterDictionary ¶mete } // MLTSampler Method Definitions -Float MLTSampler::Get1D() { +PBRT_CPU_GPU Float MLTSampler::Get1D() { int index = GetNextIndex(); EnsureReady(index); return X[index].value; } -Point2f MLTSampler::Get2D() { +PBRT_CPU_GPU Point2f MLTSampler::Get2D() { return {Get1D(), Get1D()}; } -Point2f MLTSampler::GetPixel2D() { +PBRT_CPU_GPU Point2f MLTSampler::GetPixel2D() { return Get2D(); } @@ -337,17 +337,17 @@ Sampler MLTSampler::Clone(Allocator alloc) { return {}; } -void MLTSampler::StartIteration() { +PBRT_CPU_GPU void MLTSampler::StartIteration() { currentIteration++; largeStep = rng.Uniform() < largeStepProbability; } -void MLTSampler::Accept() { +PBRT_CPU_GPU void MLTSampler::Accept() { if (largeStep) lastLargeStepIteration = currentIteration; } -void MLTSampler::EnsureReady(int index) { +PBRT_CPU_GPU void MLTSampler::EnsureReady(int index) { #ifdef PBRT_IS_GPU_CODE LOG_FATAL("MLTSampler not supported on GPU--needs vector resize..."); return; @@ -380,14 +380,14 @@ void MLTSampler::EnsureReady(int index) { #endif } -void MLTSampler::Reject() { +PBRT_CPU_GPU void MLTSampler::Reject() { for (auto &X_i : X) if (X_i.lastModificationIteration == currentIteration) X_i.Restore(); --currentIteration; } -void MLTSampler::StartStream(int index) { +PBRT_CPU_GPU void MLTSampler::StartStream(int index) { DCHECK_LT(index, streamCount); streamIndex = index; sampleIndex = 0; diff --git a/src/pbrt/samplers.h b/src/pbrt/samplers.h index 2d13f94f7..f81141eb9 100644 --- a/src/pbrt/samplers.h +++ b/src/pbrt/samplers.h @@ -765,29 +765,29 @@ class DebugMLTSampler : public MLTSampler { std::vector u; }; -inline void Sampler::StartPixelSample(Point2i p, int sampleIndex, int dimension) { +PBRT_CPU_GPU inline void Sampler::StartPixelSample(Point2i p, int sampleIndex, int dimension) { auto start = [&](auto ptr) { return ptr->StartPixelSample(p, sampleIndex, dimension); }; return Dispatch(start); } -inline int Sampler::SamplesPerPixel() const { +PBRT_CPU_GPU inline int Sampler::SamplesPerPixel() const { auto spp = [&](auto ptr) { return ptr->SamplesPerPixel(); }; return Dispatch(spp); } -inline Float Sampler::Get1D() { +PBRT_CPU_GPU inline Float Sampler::Get1D() { auto get = [&](auto ptr) { return ptr->Get1D(); }; return Dispatch(get); } -inline Point2f Sampler::Get2D() { +PBRT_CPU_GPU inline Point2f Sampler::Get2D() { auto get = [&](auto ptr) { return ptr->Get2D(); }; return Dispatch(get); } -inline Point2f Sampler::GetPixel2D() { +PBRT_CPU_GPU inline Point2f Sampler::GetPixel2D() { auto get = [&](auto ptr) { return ptr->GetPixel2D(); }; return Dispatch(get); } diff --git a/src/pbrt/shapes.cpp b/src/pbrt/shapes.cpp index f765baef9..7867d5c4d 100644 --- a/src/pbrt/shapes.cpp +++ b/src/pbrt/shapes.cpp @@ -25,21 +25,17 @@ #include #include -#if defined(PBRT_BUILD_GPU_RENDERER) -#include -#endif - #include namespace pbrt { // Sphere Method Definitions -Bounds3f Sphere::Bounds() const { +PBRT_CPU_GPU Bounds3f Sphere::Bounds() const { return (*renderFromObject)( Bounds3f(Point3f(-radius, -radius, zMin), Point3f(radius, radius, zMax))); } -pstd::optional Sphere::Sample(Point2f u) const { +PBRT_CPU_GPU pstd::optional Sphere::Sample(Point2f u) const { Point3f pObj = Point3f(0, 0, 0) + radius * SampleUniformSphere(u); // Reproject _pObj_ to sphere surface and compute _pObjError_ pObj *= radius / Distance(pObj, Point3f(0, 0, 0)); @@ -85,12 +81,12 @@ Sphere *Sphere::Create(const Transform *renderFromObject, } // Disk Method Definitions -Bounds3f Disk::Bounds() const { +PBRT_CPU_GPU Bounds3f Disk::Bounds() const { return (*renderFromObject)( Bounds3f(Point3f(-radius, -radius, height), Point3f(radius, radius, height))); } -DirectionCone Disk::NormalBounds() const { +PBRT_CPU_GPU DirectionCone Disk::NormalBounds() const { Normal3f n = (*renderFromObject)(Normal3f(0, 0, 1)); if (reverseOrientation) n = -n; @@ -119,7 +115,7 @@ Disk *Disk::Create(const Transform *renderFromObject, const Transform *objectFro } // Cylinder Method Definitions -Bounds3f Cylinder::Bounds() const { +PBRT_CPU_GPU Bounds3f Cylinder::Bounds() const { return (*renderFromObject)( Bounds3f({-radius, -radius, zMin}, {radius, radius, zMax})); } @@ -161,15 +157,15 @@ void Triangle::Init(Allocator alloc) { allMeshes = alloc.new_object>(alloc); #if defined(PBRT_BUILD_GPU_RENDERER) if (Options->useGPU) - CUDA_CHECK( - cudaMemcpyToSymbol(allTriangleMeshesGPU, &allMeshes, sizeof(allMeshes))); + CUDA_CHECK(cudaMemcpyToSymbol((const void *)&allTriangleMeshesGPU, + (const void *)&allMeshes, sizeof(allMeshes))); #endif } STAT_MEMORY_COUNTER("Memory/Triangles", triangleBytes); // Triangle Functions -pstd::optional IntersectTriangle(const Ray &ray, Float tMax, +PBRT_CPU_GPU pstd::optional IntersectTriangle(const Ray &ray, Float tMax, Point3f p0, Point3f p1, Point3f p2) { // Return no intersection if triangle is degenerate @@ -291,7 +287,7 @@ pstd::vector Triangle::CreateTriangles(const TriangleMesh *mesh, Allocato return tris; } -Bounds3f Triangle::Bounds() const { +PBRT_CPU_GPU Bounds3f Triangle::Bounds() const { // Get triangle vertices in _p0_, _p1_, and _p2_ const TriangleMesh *mesh = GetMesh(); const int *v = &mesh->vertexIndices[3 * triIndex]; @@ -300,7 +296,7 @@ Bounds3f Triangle::Bounds() const { return Union(Bounds3f(p0, p1), p2); } -DirectionCone Triangle::NormalBounds() const { +PBRT_CPU_GPU DirectionCone Triangle::NormalBounds() const { // Get triangle vertices in _p0_, _p1_, and _p2_ const TriangleMesh *mesh = GetMesh(); const int *v = &mesh->vertexIndices[3 * triIndex]; @@ -317,7 +313,7 @@ DirectionCone Triangle::NormalBounds() const { return DirectionCone(Vector3f(n)); } -pstd::optional Triangle::Intersect(const Ray &ray, Float tMax) const { +PBRT_CPU_GPU pstd::optional Triangle::Intersect(const Ray &ray, Float tMax) const { #ifndef PBRT_IS_GPU_CODE ++nTriTests; #endif @@ -338,7 +334,7 @@ pstd::optional Triangle::Intersect(const Ray &ray, Float tMax return ShapeIntersection{intr, triIsect->t}; } -bool Triangle::IntersectP(const Ray &ray, Float tMax) const { +PBRT_CPU_GPU bool Triangle::IntersectP(const Ray &ray, Float tMax) const { #ifndef PBRT_IS_GPU_CODE ++nTriTests; #endif @@ -520,7 +516,7 @@ pstd::vector CreateCurve(const Transform *renderFromObject, } // Curve Method Definitions -Bounds3f Curve::Bounds() const { +PBRT_CPU_GPU Bounds3f Curve::Bounds() const { pstd::span cpSpan(common->cpObj); Bounds3f objBounds = BoundCubicBezier(cpSpan, uMin, uMax); // Expand _objBounds_ by maximum curve width over $u$ range @@ -531,7 +527,7 @@ Bounds3f Curve::Bounds() const { return (*common->renderFromObject)(objBounds); } -Float Curve::Area() const { +PBRT_CPU_GPU Float Curve::Area() const { pstd::array cpObj = CubicBezierControlPoints(pstd::MakeConstSpan(common->cpObj), uMin, uMax); Float width0 = Lerp(uMin, common->width[0], common->width[1]); @@ -737,23 +733,23 @@ bool Curve::RecursiveIntersect(const Ray &ray, Float tMax, pstd::span Curve::Sample(Point2f u) const { +PBRT_CPU_GPU pstd::optional Curve::Sample(Point2f u) const { LOG_FATAL("Curve::Sample not implemented."); return {}; } -Float Curve::PDF(const Interaction &) const { +PBRT_CPU_GPU Float Curve::PDF(const Interaction &) const { LOG_FATAL("Curve::PDF not implemented."); return {}; } -pstd::optional Curve::Sample(const ShapeSampleContext &ctx, +PBRT_CPU_GPU pstd::optional Curve::Sample(const ShapeSampleContext &ctx, Point2f u) const { LOG_FATAL("Curve::Sample not implemented."); return {}; } -Float Curve::PDF(const ShapeSampleContext &ctx, Vector3f wi) const { +PBRT_CPU_GPU Float Curve::PDF(const ShapeSampleContext &ctx, Vector3f wi) const { LOG_FATAL("Curve::PDF not implemented."); return {}; } @@ -1030,8 +1026,9 @@ void BilinearPatch::Init(Allocator alloc) { allMeshes = alloc.new_object>(alloc); #if defined(PBRT_BUILD_GPU_RENDERER) if (Options->useGPU) - CUDA_CHECK( - cudaMemcpyToSymbol(allBilinearMeshesGPU, &allMeshes, sizeof(allMeshes))); + CUDA_CHECK(cudaMemcpyToSymbol((const void *)&allBilinearMeshesGPU, + (const void *)&allMeshes, + sizeof(allMeshes))); #endif } @@ -1070,7 +1067,7 @@ BilinearPatch::BilinearPatch(const BilinearPatchMesh *mesh, int meshIndex, int b } } -Bounds3f BilinearPatch::Bounds() const { +PBRT_CPU_GPU Bounds3f BilinearPatch::Bounds() const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ const int *v = &mesh->vertexIndices[4 * blpIndex]; @@ -1080,7 +1077,7 @@ Bounds3f BilinearPatch::Bounds() const { return Union(Bounds3f(p00, p01), Bounds3f(p10, p11)); } -DirectionCone BilinearPatch::NormalBounds() const { +PBRT_CPU_GPU DirectionCone BilinearPatch::NormalBounds() const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ const int *v = &mesh->vertexIndices[4 * blpIndex]; @@ -1128,7 +1125,7 @@ DirectionCone BilinearPatch::NormalBounds() const { return DirectionCone(n, Clamp(cosTheta, -1, 1)); } -pstd::optional BilinearPatch::Intersect(const Ray &ray, +PBRT_CPU_GPU pstd::optional BilinearPatch::Intersect(const Ray &ray, Float tMax) const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ @@ -1145,7 +1142,7 @@ pstd::optional BilinearPatch::Intersect(const Ray &ray, return ShapeIntersection{intr, blpIsect->t}; } -bool BilinearPatch::IntersectP(const Ray &ray, Float tMax) const { +PBRT_CPU_GPU bool BilinearPatch::IntersectP(const Ray &ray, Float tMax) const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ const int *v = &mesh->vertexIndices[4 * blpIndex]; @@ -1155,7 +1152,7 @@ bool BilinearPatch::IntersectP(const Ray &ray, Float tMax) const { return IntersectBilinearPatch(ray, tMax, p00, p10, p01, p11).has_value(); } -pstd::optional BilinearPatch::Sample(Point2f u) const { +PBRT_CPU_GPU pstd::optional BilinearPatch::Sample(Point2f u) const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ const int *v = &mesh->vertexIndices[4 * blpIndex]; @@ -1216,7 +1213,7 @@ pstd::optional BilinearPatch::Sample(Point2f u) const { pdf / Length(Cross(dpdu, dpdv))}; } -Float BilinearPatch::PDF(const Interaction &intr) const { +PBRT_CPU_GPU Float BilinearPatch::PDF(const Interaction &intr) const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ const int *v = &mesh->vertexIndices[4 * blpIndex]; @@ -1254,7 +1251,7 @@ Float BilinearPatch::PDF(const Interaction &intr) const { return pdf / Length(Cross(dpdu, dpdv)); } -pstd::optional BilinearPatch::Sample(const ShapeSampleContext &ctx, +PBRT_CPU_GPU pstd::optional BilinearPatch::Sample(const ShapeSampleContext &ctx, Point2f u) const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ @@ -1329,7 +1326,7 @@ pstd::optional BilinearPatch::Sample(const ShapeSampleContext &ctx, return ShapeSample{Interaction(p, n, ctx.time, st), pdf}; } -Float BilinearPatch::PDF(const ShapeSampleContext &ctx, Vector3f wi) const { +PBRT_CPU_GPU Float BilinearPatch::PDF(const ShapeSampleContext &ctx, Vector3f wi) const { const BilinearPatchMesh *mesh = GetMesh(); // Get bilinear patch vertices in _p00_, _p01_, _p10_, and _p11_ const int *v = &mesh->vertexIndices[4 * blpIndex]; diff --git a/src/pbrt/shapes.h b/src/pbrt/shapes.h index fc984b9d8..f74df4ba9 100644 --- a/src/pbrt/shapes.h +++ b/src/pbrt/shapes.h @@ -812,7 +812,7 @@ inline Cylinder::Cylinder(const Transform *renderFromObject, phiMax(Radians(Clamp(phiMax, 0, 360))) {} // Triangle Declarations -#if defined(PBRT_BUILD_GPU_RENDERER) && defined(__CUDACC__) +#if defined(PBRT_BUILD_GPU_RENDERER) && (defined(__HIPCC__) || defined(__CUDACC__)) extern PBRT_GPU pstd::vector *allTriangleMeshesGPU; #endif @@ -1264,7 +1264,7 @@ class Curve { }; // BilinearPatch Declarations -#if defined(PBRT_BUILD_GPU_RENDERER) && defined(__CUDACC__) +#if defined(PBRT_BUILD_GPU_RENDERER) && (defined(__HIPCC__) || defined(__CUDACC__)) extern PBRT_GPU pstd::vector *allBilinearMeshesGPU; #endif @@ -1538,49 +1538,49 @@ class BilinearPatch { static constexpr Float MinSphericalSampleArea = 1e-4; }; -inline Bounds3f Shape::Bounds() const { +PBRT_CPU_GPU inline Bounds3f Shape::Bounds() const { auto bounds = [&](auto ptr) { return ptr->Bounds(); }; return Dispatch(bounds); } -inline pstd::optional Shape::Intersect(const Ray &ray, +PBRT_CPU_GPU inline pstd::optional Shape::Intersect(const Ray &ray, Float tMax) const { auto intr = [&](auto ptr) { return ptr->Intersect(ray, tMax); }; return Dispatch(intr); } -inline bool Shape::IntersectP(const Ray &ray, Float tMax) const { +PBRT_CPU_GPU inline bool Shape::IntersectP(const Ray &ray, Float tMax) const { auto intr = [&](auto ptr) { return ptr->IntersectP(ray, tMax); }; return Dispatch(intr); } -inline Float Shape::Area() const { +PBRT_CPU_GPU inline Float Shape::Area() const { auto area = [&](auto ptr) { return ptr->Area(); }; return Dispatch(area); } -inline pstd::optional Shape::Sample(Point2f u) const { +PBRT_CPU_GPU inline pstd::optional Shape::Sample(Point2f u) const { auto sample = [&](auto ptr) { return ptr->Sample(u); }; return Dispatch(sample); } -inline Float Shape::PDF(const Interaction &in) const { +PBRT_CPU_GPU inline Float Shape::PDF(const Interaction &in) const { auto pdf = [&](auto ptr) { return ptr->PDF(in); }; return Dispatch(pdf); } -inline pstd::optional Shape::Sample(const ShapeSampleContext &ctx, +PBRT_CPU_GPU inline pstd::optional Shape::Sample(const ShapeSampleContext &ctx, Point2f u) const { auto sample = [&](auto ptr) { return ptr->Sample(ctx, u); }; return Dispatch(sample); } -inline Float Shape::PDF(const ShapeSampleContext &ctx, Vector3f wi) const { +PBRT_CPU_GPU inline Float Shape::PDF(const ShapeSampleContext &ctx, Vector3f wi) const { auto pdf = [&](auto ptr) { return ptr->PDF(ctx, wi); }; return Dispatch(pdf); } -inline DirectionCone Shape::NormalBounds() const { +PBRT_CPU_GPU inline DirectionCone Shape::NormalBounds() const { auto nb = [&](auto ptr) { return ptr->NormalBounds(); }; return Dispatch(nb); } diff --git a/src/pbrt/util/color.cpp b/src/pbrt/util/color.cpp index 6cf9e1003..33fbfcf5a 100644 --- a/src/pbrt/util/color.cpp +++ b/src/pbrt/util/color.cpp @@ -5,8 +5,7 @@ #include #if defined(PBRT_BUILD_GPU_RENDERER) -#include -#include +#include #endif #include @@ -33,7 +32,7 @@ std::string RGBSigmoidPolynomial::ToString() const { } // RGBToSpectrumTable Method Definitions -RGBSigmoidPolynomial RGBToSpectrumTable::operator()(RGB rgb) const { +PBRT_CPU_GPU RGBSigmoidPolynomial RGBToSpectrumTable::operator()(RGB rgb) const { DCHECK(rgb[0] >= 0.f && rgb[1] >= 0.f && rgb[2] >= 0.f && rgb[0] <= 1.f && rgb[1] <= 1.f && rgb[2] <= 1.f); @@ -187,21 +186,21 @@ std::string XYZ::ToString() const { } // ColorEncoding Method Definitions -void sRGBColorEncoding::FromLinear(pstd::span vin, +PBRT_CPU_GPU void sRGBColorEncoding::FromLinear(pstd::span vin, pstd::span vout) const { DCHECK_EQ(vin.size(), vout.size()); for (size_t i = 0; i < vin.size(); ++i) vout[i] = LinearToSRGB8(vin[i]); } -void sRGBColorEncoding::ToLinear(pstd::span vin, +PBRT_CPU_GPU void sRGBColorEncoding::ToLinear(pstd::span vin, pstd::span vout) const { DCHECK_EQ(vin.size(), vout.size()); for (size_t i = 0; i < vin.size(); ++i) vout[i] = SRGB8ToLinear(vin[i]); } -Float sRGBColorEncoding::ToFloatLinear(Float v) const { +PBRT_CPU_GPU Float sRGBColorEncoding::ToFloatLinear(Float v) const { return SRGBToLinear(v); } @@ -249,7 +248,7 @@ const ColorEncoding ColorEncoding::Get(const std::string &name, Allocator alloc) } } -GammaColorEncoding::GammaColorEncoding(Float gamma) : gamma(gamma) { +PBRT_CPU_GPU GammaColorEncoding::GammaColorEncoding(Float gamma) : gamma(gamma) { for (int i = 0; i < 256; ++i) { Float v = Float(i) / 255.f; applyLUT[i] = std::pow(v, gamma); @@ -260,18 +259,18 @@ GammaColorEncoding::GammaColorEncoding(Float gamma) : gamma(gamma) { } } -void GammaColorEncoding::ToLinear(pstd::span vin, +PBRT_CPU_GPU void GammaColorEncoding::ToLinear(pstd::span vin, pstd::span vout) const { DCHECK_EQ(vin.size(), vout.size()); for (size_t i = 0; i < vin.size(); ++i) vout[i] = applyLUT[vin[i]]; } -Float GammaColorEncoding::ToFloatLinear(Float v) const { +PBRT_CPU_GPU Float GammaColorEncoding::ToFloatLinear(Float v) const { return std::pow(v, gamma); } -void GammaColorEncoding::FromLinear(pstd::span vin, +PBRT_CPU_GPU void GammaColorEncoding::FromLinear(pstd::span vin, pstd::span vout) const { DCHECK_EQ(vin.size(), vout.size()); for (size_t i = 0; i < vin.size(); ++i) diff --git a/src/pbrt/util/image.cpp b/src/pbrt/util/image.cpp index 5ee5862c1..4929bbebe 100644 --- a/src/pbrt/util/image.cpp +++ b/src/pbrt/util/image.cpp @@ -66,7 +66,7 @@ std::string ToString(PixelFormat format) { } } -int TexelBytes(PixelFormat format) { +PBRT_CPU_GPU int TexelBytes(PixelFormat format) { switch (format) { case PixelFormat::U256: return 1; @@ -1019,7 +1019,7 @@ bool Image::Write(std::string name, const ImageMetadata &metadata) const { /////////////////////////////////////////////////////////////////////////// // OpenEXR - +#ifndef PBRT_IS_GPU_CODE static Imf::FrameBuffer imageToFrameBuffer(const Image &image, const ImageChannelDesc &desc, const Imath::Box2i &dataWindow) { @@ -1812,5 +1812,6 @@ bool Image::WritePFM(const std::string &filename, const ImageMetadata &metadata) fclose(fp); return false; } +#endif } // namespace pbrt diff --git a/src/pbrt/util/math.cpp b/src/pbrt/util/math.cpp index d65a555bb..4f7a136bf 100644 --- a/src/pbrt/util/math.cpp +++ b/src/pbrt/util/math.cpp @@ -43,7 +43,7 @@ std::string SquareMatrix::ToString() const { // General case template -pstd::optional> Inverse(const SquareMatrix &m) { +PBRT_CPU_GPU pstd::optional> Inverse(const SquareMatrix &m) { int indxc[N], indxr[N]; int ipiv[N] = {0}; Float minv[N][N]; @@ -106,8 +106,8 @@ pstd::optional> Inverse(const SquareMatrix &m) { } template class SquareMatrix<2>; -template pstd::optional> Inverse(const SquareMatrix<2> &); -template SquareMatrix<2> operator*(const SquareMatrix<2> &m1, const SquareMatrix<2> &m2); +template PBRT_CPU_GPU pstd::optional> Inverse(const SquareMatrix<2> &); +template PBRT_CPU_GPU SquareMatrix<2> operator*(const SquareMatrix<2> &m1, const SquareMatrix<2> &m2); template class SquareMatrix<3>; template class SquareMatrix<4>; @@ -154,7 +154,7 @@ std::string Interval::ToString() const { } // Spline Interpolation Function Definitions -bool CatmullRomWeights(pstd::span nodes, Float x, int *offset, +PBRT_CPU_GPU bool CatmullRomWeights(pstd::span nodes, Float x, int *offset, pstd::span weights) { CHECK_GE(weights.size(), 4); // Return _false_ if _x_ is out of bounds @@ -200,7 +200,7 @@ bool CatmullRomWeights(pstd::span nodes, Float x, int *offset, return true; } -Float CatmullRom(pstd::span nodes, pstd::span f, Float x) { +PBRT_CPU_GPU Float CatmullRom(pstd::span nodes, pstd::span f, Float x) { CHECK_EQ(nodes.size(), f.size()); if (!(x >= nodes.front() && x <= nodes.back())) return 0; @@ -224,7 +224,7 @@ Float CatmullRom(pstd::span nodes, pstd::span f, Float (t3 - t2) * d1; } -Float InvertCatmullRom(pstd::span nodes, pstd::span f, +PBRT_CPU_GPU Float InvertCatmullRom(pstd::span nodes, pstd::span f, Float u) { // Stop when _u_ is out of bounds if (!(u > f.front())) @@ -264,7 +264,7 @@ Float InvertCatmullRom(pstd::span nodes, pstd::span f, return x0 + t * width; } -Float IntegrateCatmullRom(pstd::span nodes, pstd::span f, +PBRT_CPU_GPU Float IntegrateCatmullRom(pstd::span nodes, pstd::span f, pstd::span cdf) { CHECK_EQ(nodes.size(), f.size()); Float sum = 0; @@ -289,7 +289,7 @@ Float IntegrateCatmullRom(pstd::span nodes, pstd::span // Square--Sphere Mapping Function Definitions // Via source code from Clarberg: Fast Equal-Area Mapping of the (Hemi)Sphere using SIMD -Vector3f EqualAreaSquareToSphere(Point2f p) { +PBRT_CPU_GPU Vector3f EqualAreaSquareToSphere(Point2f p) { CHECK(p.x >= 0 && p.x <= 1 && p.y >= 0 && p.y <= 1); // Transform _p_ to $[-1,1]^2$ and compute absolute values Float u = 2 * p.x - 1, v = 2 * p.y - 1; @@ -314,7 +314,7 @@ Vector3f EqualAreaSquareToSphere(Point2f p) { } // Via source code from Clarberg: Fast Equal-Area Mapping of the (Hemi)Sphere using SIMD -Point2f EqualAreaSphereToSquare(Vector3f d) { +PBRT_CPU_GPU Point2f EqualAreaSphereToSquare(Vector3f d) { DCHECK(LengthSquared(d) > .999 && LengthSquared(d) < 1.001); Float x = std::abs(d.x), y = std::abs(d.y), z = std::abs(d.z); @@ -360,7 +360,7 @@ Point2f EqualAreaSphereToSquare(Vector3f d) { return Point2f(0.5f * (u + 1), 0.5f * (v + 1)); } -Point2f WrapEqualAreaSquare(Point2f uv) { +PBRT_CPU_GPU Point2f WrapEqualAreaSquare(Point2f uv) { if (uv[0] < 0) { uv[0] = -uv[0]; // mirror across u = 0 uv[1] = 1 - uv[1]; // mirror across v = 0.5 diff --git a/src/pbrt/util/math.h b/src/pbrt/util/math.h index 3574318a1..864967256 100644 --- a/src/pbrt/util/math.h +++ b/src/pbrt/util/math.h @@ -25,7 +25,7 @@ namespace pbrt { -#ifdef PBRT_IS_GPU_CODE +#if defined(PBRT_IS_GPU_CODE) && defined(__HIP_PLATFORM_NVIDIA__) #define ShadowEpsilon 0.0001f #define Pi Float(3.14159265358979323846) @@ -471,7 +471,6 @@ PBRT_CPU_GPU inline float FastExp(float x) { bits &= 0b10000000011111111111111111111111u; bits |= (exponent + 127) << 23; return BitsToFloat(bits); - #endif } @@ -957,7 +956,7 @@ PBRT_CPU_GPU inline bool InRange(Interval a, Interval b) { return a.LowerBound() <= b.UpperBound() && a.UpperBound() >= b.LowerBound(); } -inline Interval Interval::operator/(Interval i) const { +PBRT_CPU_GPU inline Interval Interval::operator/(Interval i) const { if (InRange(0, i)) // The interval we're dividing by straddles zero, so just // return an interval of everything. @@ -1384,7 +1383,7 @@ class SquareMatrix { // SquareMatrix Inline Methods template -inline bool SquareMatrix::IsIdentity() const { +PBRT_CPU_GPU inline bool SquareMatrix::IsIdentity() const { for (int i = 0; i < N; ++i) for (int j = 0; j < N; ++j) { if (i == j) { diff --git a/src/pbrt/util/rng.h b/src/pbrt/util/rng.h index 2c035dfae..4dd4e602f 100644 --- a/src/pbrt/util/rng.h +++ b/src/pbrt/util/rng.h @@ -71,15 +71,15 @@ class RNG { // RNG Inline Method Definitions template -inline T RNG::Uniform() { +PBRT_CPU_GPU inline T RNG::Uniform() { return T::unimplemented; } template <> -inline uint32_t RNG::Uniform(); +PBRT_CPU_GPU inline uint32_t RNG::Uniform(); template <> -inline uint32_t RNG::Uniform() { +PBRT_CPU_GPU inline uint32_t RNG::Uniform() { uint64_t oldstate = state; state = oldstate * PCG32_MULT + inc; uint32_t xorshifted = (uint32_t)(((oldstate >> 18u) ^ oldstate) >> 27u); @@ -88,13 +88,13 @@ inline uint32_t RNG::Uniform() { } template <> -inline uint64_t RNG::Uniform() { +PBRT_CPU_GPU inline uint64_t RNG::Uniform() { uint64_t v0 = Uniform(), v1 = Uniform(); return (v0 << 32) | v1; } template <> -inline int32_t RNG::Uniform() { +PBRT_CPU_GPU inline int32_t RNG::Uniform() { // https://stackoverflow.com/a/13208789 uint32_t v = Uniform(); if (v <= (uint32_t)std::numeric_limits::max()) @@ -105,7 +105,7 @@ inline int32_t RNG::Uniform() { } template <> -inline int64_t RNG::Uniform() { +PBRT_CPU_GPU inline int64_t RNG::Uniform() { // https://stackoverflow.com/a/13208789 uint64_t v = Uniform(); if (v <= (uint64_t)std::numeric_limits::max()) @@ -116,7 +116,7 @@ inline int64_t RNG::Uniform() { std::numeric_limits::min(); } -inline void RNG::SetSequence(uint64_t sequenceIndex, uint64_t seed) { +PBRT_CPU_GPU inline void RNG::SetSequence(uint64_t sequenceIndex, uint64_t seed) { state = 0u; inc = (sequenceIndex << 1u) | 1u; Uniform(); @@ -125,16 +125,16 @@ inline void RNG::SetSequence(uint64_t sequenceIndex, uint64_t seed) { } template <> -inline float RNG::Uniform() { +PBRT_CPU_GPU inline float RNG::Uniform() { return std::min(OneMinusEpsilon, Uniform() * 0x1p-32f); } template <> -inline double RNG::Uniform() { +PBRT_CPU_GPU inline double RNG::Uniform() { return std::min(OneMinusEpsilon, Uniform() * 0x1p-64); } -inline void RNG::Advance(int64_t idelta) { +PBRT_CPU_GPU inline void RNG::Advance(int64_t idelta) { uint64_t curMult = PCG32_MULT, curPlus = inc, accMult = 1u; uint64_t accPlus = 0u, delta = (uint64_t)idelta; while (delta > 0) { diff --git a/src/pbrt/util/sampling.cpp b/src/pbrt/util/sampling.cpp index 78d71d42a..1043880eb 100644 --- a/src/pbrt/util/sampling.cpp +++ b/src/pbrt/util/sampling.cpp @@ -25,7 +25,7 @@ namespace pbrt { // Sampling Function Definitions -pstd::array SampleSphericalTriangle(const pstd::array &v, Point3f p, +PBRT_CPU_GPU pstd::array SampleSphericalTriangle(const pstd::array &v, Point3f p, Point2f u, Float *pdf) { if (pdf) *pdf = 0; @@ -107,7 +107,7 @@ pstd::array SampleSphericalTriangle(const pstd::array &v, } // Via Jim Arvo's SphTri.C -Point2f InvertSphericalTriangleSample(const pstd::array &v, Point3f p, +PBRT_CPU_GPU Point2f InvertSphericalTriangleSample(const pstd::array &v, Point3f p, Vector3f w) { // Compute vectors _a_, _b_, and _c_ to spherical triangle vertices Vector3f a(v[0] - p), b(v[1] - p), c(v[2] - p); @@ -160,7 +160,7 @@ Point2f InvertSphericalTriangleSample(const pstd::array &v, Point3f return Point2f(Clamp(u0, 0, 1), Clamp(u1, 0, 1)); } -Point3f SampleSphericalRectangle(Point3f pRef, Point3f s, Vector3f ex, Vector3f ey, +PBRT_CPU_GPU Point3f SampleSphericalRectangle(Point3f pRef, Point3f s, Vector3f ex, Vector3f ey, Point2f u, Float *pdf) { // Compute local reference frame and transform rectangle coordinates Float exl = Length(ex), eyl = Length(ey); @@ -219,7 +219,8 @@ Point3f SampleSphericalRectangle(Point3f pRef, Point3f s, Vector3f ex, Vector3f return pRef + R.FromLocal(Vector3f(xu, yv, z0)); } -Point2f InvertSphericalRectangleSample(Point3f pRef, Point3f s, Vector3f ex, Vector3f ey, +PBRT_CPU_GPU Point2f InvertSphericalRectangleSample(Point3f pRef, Point3f s, Vector3f ex, + Vector3f ey, Point3f pRect) { // TODO: Delete anything unused in the below... @@ -344,7 +345,7 @@ Point2f InvertSphericalRectangleSample(Point3f pRef, Point3f s, Vector3f ex, Vec return u; } -Vector3f SampleHenyeyGreenstein(Vector3f wo, Float g, Point2f u, Float *pdf) { +PBRT_CPU_GPU Vector3f SampleHenyeyGreenstein(Vector3f wo, Float g, Point2f u, Float *pdf) { // When g \approx -1 and u[0] \approx 0 or with g \approx 1 and u[0] // \approx 1, the computation of cosTheta below is unstable and can // give, leading to NaNs. For now we limit g to the range where it is @@ -381,7 +382,7 @@ Point2f RejectionSampleDisk(RNG &rng) { return p; } -Float SampleCatmullRom(pstd::span nodes, pstd::span f, +PBRT_CPU_GPU Float SampleCatmullRom(pstd::span nodes, pstd::span f, pstd::span F, Float u, Float *fval, Float *pdf) { CHECK_EQ(nodes.size(), f.size()); CHECK_EQ(f.size(), F.size()); @@ -421,7 +422,7 @@ Float SampleCatmullRom(pstd::span nodes, pstd::span f, return x0 + width * t; } -Float SampleCatmullRom2D(pstd::span nodes1, pstd::span nodes2, +PBRT_CPU_GPU Float SampleCatmullRom2D(pstd::span nodes1, pstd::span nodes2, pstd::span values, pstd::span cdf, Float alpha, Float u, Float *fval, Float *pdf) { // Determine offset and coefficients for the _alpha_ parameter @@ -616,7 +617,7 @@ AliasTable::AliasTable(pstd::span weights, Allocator alloc) } } -int AliasTable::Sample(Float u, Float *pmf, Float *uRemapped) const { +PBRT_CPU_GPU int AliasTable::Sample(Float u, Float *pmf, Float *uRemapped) const { // Compute alias table _offset_ and remapped random sample _up_ int offset = std::min(u * bins.size(), bins.size() - 1); Float up = std::min(u * bins.size() - offset, OneMinusEpsilon); diff --git a/src/pbrt/util/sampling.h b/src/pbrt/util/sampling.h index e5caf7831..f0f844c1f 100644 --- a/src/pbrt/util/sampling.h +++ b/src/pbrt/util/sampling.h @@ -1300,7 +1300,7 @@ class PiecewiseLinear2D { private: using FloatStorage = pstd::vector; -#if !defined(_MSC_VER) && !defined(__CUDACC__) +#if !defined(_MSC_VER) && !(defined(__CUDACC__) || defined(__HIPCC__)) static constexpr size_t ArraySize = Dimension; #else static constexpr size_t ArraySize = (Dimension != 0) ? Dimension : 1; diff --git a/src/pbrt/util/spectrum.cpp b/src/pbrt/util/spectrum.cpp index d59f4f898..03e439d5b 100644 --- a/src/pbrt/util/spectrum.cpp +++ b/src/pbrt/util/spectrum.cpp @@ -65,7 +65,7 @@ std::string Spectrum::ToString() const { } // Spectrum Method Definitions -Float PiecewiseLinearSpectrum::operator()(Float lambda) const { +PBRT_CPU_GPU Float PiecewiseLinearSpectrum::operator()(Float lambda) const { // Handle _PiecewiseLinearSpectrum_ corner cases if (lambdas.empty() || lambda < lambdas.front() || lambda > lambdas.back()) return 0; @@ -77,7 +77,7 @@ Float PiecewiseLinearSpectrum::operator()(Float lambda) const { return Lerp(t, values[o], values[o + 1]); } -Float PiecewiseLinearSpectrum::MaxValue() const { +PBRT_CPU_GPU Float PiecewiseLinearSpectrum::MaxValue() const { if (values.empty()) return 0; return *std::max_element(values.begin(), values.end()); @@ -166,7 +166,7 @@ std::string BlackbodySpectrum::ToString() const { return StringPrintf("[ BlackbodySpectrum T: %f ]", T); } -SampledSpectrum ConstantSpectrum::Sample(const SampledWavelengths &) const { +PBRT_CPU_GPU SampledSpectrum ConstantSpectrum::Sample(const SampledWavelengths &) const { return SampledSpectrum(c); } @@ -206,7 +206,7 @@ std::string SampledWavelengths::ToString() const { return r; } -XYZ SampledSpectrum::ToXYZ(const SampledWavelengths &lambda) const { +PBRT_CPU_GPU XYZ SampledSpectrum::ToXYZ(const SampledWavelengths &lambda) const { // Sample the $X$, $Y$, and $Z$ matching curves at _lambda_ SampledSpectrum X = Spectra::X().Sample(lambda); SampledSpectrum Y = Spectra::Y().Sample(lambda); @@ -219,31 +219,31 @@ XYZ SampledSpectrum::ToXYZ(const SampledWavelengths &lambda) const { CIE_Y_integral; } -Float SampledSpectrum::y(const SampledWavelengths &lambda) const { +PBRT_CPU_GPU Float SampledSpectrum::y(const SampledWavelengths &lambda) const { SampledSpectrum Ys = Spectra::Y().Sample(lambda); SampledSpectrum pdf = lambda.PDF(); return SafeDiv(Ys * *this, pdf).Average() / CIE_Y_integral; } -RGB SampledSpectrum::ToRGB(const SampledWavelengths &lambda, +PBRT_CPU_GPU RGB SampledSpectrum::ToRGB(const SampledWavelengths &lambda, const RGBColorSpace &cs) const { XYZ xyz = ToXYZ(lambda); return cs.ToRGB(xyz); } -RGBAlbedoSpectrum::RGBAlbedoSpectrum(const RGBColorSpace &cs, RGB rgb) { +PBRT_CPU_GPU RGBAlbedoSpectrum::RGBAlbedoSpectrum(const RGBColorSpace &cs, RGB rgb) { DCHECK_LE(std::max({rgb.r, rgb.g, rgb.b}), 1); DCHECK_GE(std::min({rgb.r, rgb.g, rgb.b}), 0); rsp = cs.ToRGBCoeffs(rgb); } -RGBUnboundedSpectrum::RGBUnboundedSpectrum(const RGBColorSpace &cs, RGB rgb) { +PBRT_CPU_GPU RGBUnboundedSpectrum::RGBUnboundedSpectrum(const RGBColorSpace &cs, RGB rgb) { Float m = std::max({rgb.r, rgb.g, rgb.b}); scale = 2 * m; rsp = cs.ToRGBCoeffs(scale ? rgb / scale : RGB(0, 0, 0)); } -RGBIlluminantSpectrum::RGBIlluminantSpectrum(const RGBColorSpace &cs, RGB rgb) +PBRT_CPU_GPU RGBIlluminantSpectrum::RGBIlluminantSpectrum(const RGBColorSpace &cs, RGB rgb) : illuminant(&cs.illuminant) { Float m = std::max({rgb.r, rgb.g, rgb.b}); scale = 2 * m; @@ -2595,9 +2595,9 @@ void Init(Allocator alloc) { #ifdef PBRT_BUILD_GPU_RENDERER if (Options->useGPU) { - CUDA_CHECK(cudaMemcpyToSymbol(xGPU, &x, sizeof(x))); - CUDA_CHECK(cudaMemcpyToSymbol(yGPU, &y, sizeof(y))); - CUDA_CHECK(cudaMemcpyToSymbol(zGPU, &z, sizeof(z))); + CUDA_CHECK(cudaMemcpyToSymbol((const void *)&xGPU, (const void *)&x, sizeof(x))); + CUDA_CHECK(cudaMemcpyToSymbol((const void *)&yGPU, (const void *)&y, sizeof(y))); + CUDA_CHECK(cudaMemcpyToSymbol((const void *)&zGPU, (const void *)&z, sizeof(z))); } #endif diff --git a/src/pbrt/util/spectrum.h b/src/pbrt/util/spectrum.h index 943a9993c..660523b4a 100644 --- a/src/pbrt/util/spectrum.h +++ b/src/pbrt/util/spectrum.h @@ -756,13 +756,13 @@ Spectrum GetNamedSpectrum(std::string name); std::string FindMatchingNamedSpectrum(Spectrum s); namespace Spectra { -inline const DenselySampledSpectrum &X(); -inline const DenselySampledSpectrum &Y(); -inline const DenselySampledSpectrum &Z(); +PBRT_CPU_GPU inline const DenselySampledSpectrum &X(); +PBRT_CPU_GPU inline const DenselySampledSpectrum &Y(); +PBRT_CPU_GPU inline const DenselySampledSpectrum &Z(); } // namespace Spectra // Spectrum Inline Functions -inline Float InnerProduct(Spectrum f, Spectrum g) { +PBRT_CPU_GPU inline Float InnerProduct(Spectrum f, Spectrum g) { Float integral = 0; for (Float lambda = Lambda_min; lambda <= Lambda_max; ++lambda) integral += f(lambda) * g(lambda); @@ -770,17 +770,17 @@ inline Float InnerProduct(Spectrum f, Spectrum g) { } // Spectrum Inline Method Definitions -inline Float Spectrum::operator()(Float lambda) const { +PBRT_CPU_GPU inline Float Spectrum::operator()(Float lambda) const { auto op = [&](auto ptr) { return (*ptr)(lambda); }; return Dispatch(op); } -inline SampledSpectrum Spectrum::Sample(const SampledWavelengths &lambda) const { +PBRT_CPU_GPU inline SampledSpectrum Spectrum::Sample(const SampledWavelengths &lambda) const { auto samp = [&](auto ptr) { return ptr->Sample(lambda); }; return Dispatch(samp); } -inline Float Spectrum::MaxValue() const { +PBRT_CPU_GPU inline Float Spectrum::MaxValue() const { auto max = [&](auto ptr) { return ptr->MaxValue(); }; return Dispatch(max); } diff --git a/src/pbrt/util/taggedptr.h b/src/pbrt/util/taggedptr.h index 9231de0f9..f616fb810 100644 --- a/src/pbrt/util/taggedptr.h +++ b/src/pbrt/util/taggedptr.h @@ -358,19 +358,19 @@ PBRT_CPU_GPU R Dispatch(F &&func, void *ptr, int index) { } template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_EQ(0, index); return func((const T *)ptr); } template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_EQ(0, index); return func((T *)ptr); } template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 2); @@ -381,7 +381,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { } template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 2); @@ -392,7 +392,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { } template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 3); @@ -407,7 +407,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { } template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 3); @@ -422,7 +422,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { } template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 4); @@ -439,7 +439,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { } template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 4); @@ -457,7 +457,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 5); @@ -477,7 +477,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 5); @@ -497,7 +497,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 6); @@ -519,7 +519,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 6); @@ -541,7 +541,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 7); @@ -565,7 +565,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 7); @@ -589,7 +589,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { template -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 8); @@ -615,7 +615,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { template -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); DCHECK_LT(index, 8); @@ -642,7 +642,7 @@ auto DispatchCPU(F &&func, void *ptr, int index) { template 0)>> -auto DispatchCPU(F &&func, const void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, const void *ptr, int index) { DCHECK_GE(index, 0); switch (index) { @@ -670,7 +670,7 @@ auto DispatchCPU(F &&func, const void *ptr, int index) { template 0)>> -auto DispatchCPU(F &&func, void *ptr, int index) { +PBRT_CPU_GPU auto DispatchCPU(F &&func, void *ptr, int index) { DCHECK_GE(index, 0); switch (index) { @@ -743,9 +743,14 @@ class TaggedPointer { template PBRT_CPU_GPU TaggedPointer(T *ptr) { uint64_t iptr = reinterpret_cast(ptr); - DCHECK_EQ(iptr & ptrMask, iptr); constexpr unsigned int type = TypeIndex(); +#if defined(PBRT_IS_WINDOWS) && defined(__HIP_PLATFORM_AMD__) + uint64_t aptr = (iptr & tagMask) >> 5ull; + bits = (iptr & ptrMask) | aptr | ((uint64_t)type << (tagShift + 2ull)); +#else + DCHECK_EQ(iptr & ptrMask, iptr); bits = iptr | ((uint64_t)type << tagShift); +#endif } PBRT_CPU_GPU @@ -769,7 +774,14 @@ class TaggedPointer { } PBRT_CPU_GPU - unsigned int Tag() const { return ((bits & tagMask) >> tagShift); } + unsigned int Tag() const { +#if defined(PBRT_IS_WINDOWS) && defined(__HIP_PLATFORM_AMD__) + return ((bits & tagMask) >> (tagShift + 2ull)); +#else + return ((bits & tagMask) >> tagShift); +#endif + } + template PBRT_CPU_GPU bool Is() const { return Tag() == TypeIndex(); @@ -824,10 +836,42 @@ class TaggedPointer { bool operator!=(const TaggedPointer &tp) const { return bits != tp.bits; } PBRT_CPU_GPU - void *ptr() { return reinterpret_cast(bits & ptrMask); } + void *ptr() { +#if defined(PBRT_IS_WINDOWS) && defined(__HIP_PLATFORM_AMD__) + unsigned int aptr = (bits >> tagShift) & 3; + uint64_t iptr = bits & ptrMask; + // lld crashes on Windows + // iptr |= (uint64_t)aptr << 60; + if (aptr == 1) + iptr |= 1ull << 60; + if (aptr == 2) + iptr |= 2ull << 60; + if (aptr == 3) + iptr |= 3ull << 60; + return reinterpret_cast(iptr); +#else + return reinterpret_cast(bits & ptrMask); +#endif + } PBRT_CPU_GPU - const void *ptr() const { return reinterpret_cast(bits & ptrMask); } + const void *ptr() const { +#if defined(PBRT_IS_WINDOWS) && defined(__HIP_PLATFORM_AMD__) + unsigned int aptr = (bits >> tagShift) & 3; + uint64_t iptr = bits & ptrMask; + // lld crashes on Windows + // iptr |= (uint64_t)aptr << 60; + if (aptr == 1) + iptr |= 1ull << 60; + if (aptr == 2) + iptr |= 2ull << 60; + if (aptr == 3) + iptr |= 3ull << 60; + return reinterpret_cast(iptr); +#else + return reinterpret_cast(bits & ptrMask); +#endif + } template PBRT_CPU_GPU decltype(auto) Dispatch(F &&func) { @@ -851,7 +895,7 @@ class TaggedPointer { } template - decltype(auto) DispatchCPU(F &&func) const { + PBRT_CPU_GPU decltype(auto) DispatchCPU(F &&func) const { DCHECK(ptr()); using R = typename detail::ReturnTypeConst::type; return detail::DispatchCPU(func, ptr(), Tag() - 1); @@ -861,7 +905,11 @@ class TaggedPointer { static_assert(sizeof(uintptr_t) <= sizeof(uint64_t), "Expected pointer size to be <= 64 bits"); // TaggedPointer Private Members +#if defined(PBRT_IS_WINDOWS) && defined(__HIP_PLATFORM_AMD__) + static constexpr int tagShift = 55; +#else static constexpr int tagShift = 57; +#endif static constexpr int tagBits = 64 - tagShift; static constexpr uint64_t tagMask = ((1ull << tagBits) - 1) << tagShift; static constexpr uint64_t ptrMask = ~tagMask; diff --git a/src/pbrt/util/transform.cpp b/src/pbrt/util/transform.cpp index 1c7d77dfe..9ebf809f9 100644 --- a/src/pbrt/util/transform.cpp +++ b/src/pbrt/util/transform.cpp @@ -18,7 +18,7 @@ namespace pbrt { // Transform Function Definitions // clang-format off -Transform Translate(Vector3f delta) { +PBRT_CPU_GPU Transform Translate(Vector3f delta) { SquareMatrix<4> m(1, 0, 0, delta.x, 0, 1, 0, delta.y, 0, 0, 1, delta.z, @@ -32,7 +32,7 @@ Transform Translate(Vector3f delta) { // clang-format on // clang-format off -Transform Scale(Float x, Float y, Float z) { +PBRT_CPU_GPU Transform Scale(Float x, Float y, Float z) { SquareMatrix<4> m(x, 0, 0, 0, 0, y, 0, 0, 0, 0, z, 0, @@ -46,7 +46,7 @@ Transform Scale(Float x, Float y, Float z) { // clang-format on // clang-format off -Transform RotateX(Float theta) { +PBRT_CPU_GPU Transform RotateX(Float theta) { Float sinTheta = std::sin(Radians(theta)); Float cosTheta = std::cos(Radians(theta)); SquareMatrix<4> m(1, 0, 0, 0, @@ -58,7 +58,7 @@ Transform RotateX(Float theta) { // clang-format on // clang-format off -Transform RotateY(Float theta) { +PBRT_CPU_GPU Transform RotateY(Float theta) { Float sinTheta = std::sin(Radians(theta)); Float cosTheta = std::cos(Radians(theta)); SquareMatrix<4> m( cosTheta, 0, sinTheta, 0, @@ -67,7 +67,7 @@ Transform RotateY(Float theta) { 0, 0, 0, 1); return Transform(m, Transpose(m)); } -Transform RotateZ(Float theta) { +PBRT_CPU_GPU Transform RotateZ(Float theta) { Float sinTheta = std::sin(Radians(theta)); Float cosTheta = std::cos(Radians(theta)); SquareMatrix<4> m(cosTheta, -sinTheta, 0, 0, @@ -78,7 +78,7 @@ Transform RotateZ(Float theta) { } // clang-format on -Transform LookAt(Point3f pos, Point3f look, Vector3f up) { +PBRT_CPU_GPU Transform LookAt(Point3f pos, Point3f look, Vector3f up) { SquareMatrix<4> worldFromCamera; // Initialize fourth column of viewing matrix worldFromCamera[0][3] = pos.x; @@ -112,11 +112,11 @@ Transform LookAt(Point3f pos, Point3f look, Vector3f up) { return Transform(cameraFromWorld, worldFromCamera); } -Transform Orthographic(Float zNear, Float zFar) { +PBRT_CPU_GPU Transform Orthographic(Float zNear, Float zFar) { return Scale(1, 1, 1 / (zFar - zNear)) * Translate(Vector3f(0, 0, -zNear)); } -Transform Perspective(Float fov, Float n, Float f) { +PBRT_CPU_GPU Transform Perspective(Float fov, Float n, Float f) { // Perform projective divide for perspective projection // clang-format off SquareMatrix<4> persp(1, 0, 0, 0, @@ -131,18 +131,18 @@ SquareMatrix<4> persp(1, 0, 0, 0, } // Transform Method Definitions -Bounds3f Transform::operator()(const Bounds3f &b) const { +PBRT_CPU_GPU Bounds3f Transform::operator()(const Bounds3f &b) const { Bounds3f bt; for (int i = 0; i < 8; ++i) bt = Union(bt, (*this)(b.Corner(i))); return bt; } -Transform Transform::operator*(const Transform &t2) const { +PBRT_CPU_GPU Transform Transform::operator*(const Transform &t2) const { return Transform(m * t2.m, t2.mInv * mInv); } -bool Transform::SwapsHandedness() const { +PBRT_CPU_GPU bool Transform::SwapsHandedness() const { // clang-format off SquareMatrix<3> s(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], @@ -151,7 +151,7 @@ bool Transform::SwapsHandedness() const { return Determinant(s) < 0; } -Transform::operator Quaternion() const { +PBRT_CPU_GPU Transform::operator Quaternion() const { Float trace = m[0][0] + m[1][1] + m[2][2]; Quaternion quat; if (trace > 0.f) { @@ -226,7 +226,7 @@ void Transform::Decompose(Vector3f *T, SquareMatrix<4> *R, SquareMatrix<4> *S) c *S = InvertOrExit(*R) * M; } -SurfaceInteraction Transform::operator()(const SurfaceInteraction &si) const { +PBRT_CPU_GPU SurfaceInteraction Transform::operator()(const SurfaceInteraction &si) const { SurfaceInteraction ret; const Transform &t = *this; ret.pi = t(si.pi); @@ -260,7 +260,7 @@ SurfaceInteraction Transform::operator()(const SurfaceInteraction &si) const { return ret; } -Point3fi Transform::ApplyInverse(const Point3fi &p) const { +PBRT_CPU_GPU Point3fi Transform::ApplyInverse(const Point3fi &p) const { Float x = Float(p.x), y = Float(p.y), z = Float(p.z); // Compute transformed coordinates from point _pt_ Float xp = (mInv[0][0] * x + mInv[0][1] * y) + (mInv[0][2] * z + mInv[0][3]); @@ -302,7 +302,7 @@ Point3fi Transform::ApplyInverse(const Point3fi &p) const { return Point3fi(Point3f(xp, yp, zp), pOutError) / wp; } -Interaction Transform::operator()(const Interaction &in) const { +PBRT_CPU_GPU Interaction Transform::operator()(const Interaction &in) const { Interaction ret; ret.pi = (*this)(in.pi); ret.n = (*this)(in.n); @@ -317,7 +317,7 @@ Interaction Transform::operator()(const Interaction &in) const { return ret; } -Interaction Transform::ApplyInverse(const Interaction &in) const { +PBRT_CPU_GPU Interaction Transform::ApplyInverse(const Interaction &in) const { Interaction ret; Transform t = Inverse(*this); ret.pi = t(in.pi); @@ -333,7 +333,7 @@ Interaction Transform::ApplyInverse(const Interaction &in) const { return ret; } -SurfaceInteraction Transform::ApplyInverse(const SurfaceInteraction &si) const { +PBRT_CPU_GPU SurfaceInteraction Transform::ApplyInverse(const SurfaceInteraction &si) const { SurfaceInteraction ret; ret.pi = (*this)(si.pi); @@ -961,7 +961,7 @@ AnimatedTransform::AnimatedTransform(const Transform &startTransform, Float star } } -Ray AnimatedTransform::operator()(const Ray &r, Float *tMax) const { +PBRT_CPU_GPU Ray AnimatedTransform::operator()(const Ray &r, Float *tMax) const { if (!actuallyAnimated || r.time <= startTime) return startTransform(r, tMax); else if (r.time >= endTime) @@ -972,7 +972,7 @@ Ray AnimatedTransform::operator()(const Ray &r, Float *tMax) const { } } -Ray AnimatedTransform::ApplyInverse(const Ray &r, Float *tMax) const { +PBRT_CPU_GPU Ray AnimatedTransform::ApplyInverse(const Ray &r, Float *tMax) const { if (!actuallyAnimated || r.time <= startTime) return startTransform.ApplyInverse(r, tMax); else if (r.time >= endTime) @@ -983,7 +983,7 @@ Ray AnimatedTransform::ApplyInverse(const Ray &r, Float *tMax) const { } } -RayDifferential AnimatedTransform::operator()(const RayDifferential &r, +PBRT_CPU_GPU RayDifferential AnimatedTransform::operator()(const RayDifferential &r, Float *tMax) const { if (!actuallyAnimated || r.time <= startTime) return startTransform(r, tMax); @@ -995,7 +995,7 @@ RayDifferential AnimatedTransform::operator()(const RayDifferential &r, } } -Point3f AnimatedTransform::operator()(Point3f p, Float time) const { +PBRT_CPU_GPU Point3f AnimatedTransform::operator()(Point3f p, Float time) const { if (!actuallyAnimated || time <= startTime) return startTransform(p); else if (time >= endTime) @@ -1004,7 +1004,7 @@ Point3f AnimatedTransform::operator()(Point3f p, Float time) const { return t(p); } -Vector3f AnimatedTransform::operator()(Vector3f v, Float time) const { +PBRT_CPU_GPU Vector3f AnimatedTransform::operator()(Vector3f v, Float time) const { if (!actuallyAnimated || time <= startTime) return startTransform(v); else if (time >= endTime) @@ -1013,7 +1013,7 @@ Vector3f AnimatedTransform::operator()(Vector3f v, Float time) const { return t(v); } -Normal3f AnimatedTransform::operator()(Normal3f n, Float time) const { +PBRT_CPU_GPU Normal3f AnimatedTransform::operator()(Normal3f n, Float time) const { if (!actuallyAnimated || time <= startTime) return startTransform(n); else if (time >= endTime) @@ -1022,28 +1022,28 @@ Normal3f AnimatedTransform::operator()(Normal3f n, Float time) const { return t(n); } -Interaction AnimatedTransform::operator()(const Interaction &it) const { +PBRT_CPU_GPU Interaction AnimatedTransform::operator()(const Interaction &it) const { if (!actuallyAnimated) return startTransform(it); Transform t = Interpolate(it.time); return t(it); } -Interaction AnimatedTransform::ApplyInverse(const Interaction &it) const { +PBRT_CPU_GPU Interaction AnimatedTransform::ApplyInverse(const Interaction &it) const { if (!actuallyAnimated) return startTransform.ApplyInverse(it); Transform t = Interpolate(it.time); return t.ApplyInverse(it); } -SurfaceInteraction AnimatedTransform::operator()(const SurfaceInteraction &it) const { +PBRT_CPU_GPU SurfaceInteraction AnimatedTransform::operator()(const SurfaceInteraction &it) const { if (!actuallyAnimated) return startTransform(it); Transform t = Interpolate(it.time); return t(it); } -SurfaceInteraction AnimatedTransform::ApplyInverse(const SurfaceInteraction &it) const { +PBRT_CPU_GPU SurfaceInteraction AnimatedTransform::ApplyInverse(const SurfaceInteraction &it) const { if (!actuallyAnimated) return startTransform.ApplyInverse(it); Transform t = Interpolate(it.time); @@ -1059,7 +1059,7 @@ std::string AnimatedTransform::ToString() const { hasRotation); } -Transform AnimatedTransform::Interpolate(Float time) const { +PBRT_CPU_GPU Transform AnimatedTransform::Interpolate(Float time) const { // Handle boundary conditions for matrix interpolation if (!actuallyAnimated || time <= startTime) return startTransform; @@ -1080,7 +1080,7 @@ Transform AnimatedTransform::Interpolate(Float time) const { return Translate(trans) * Transform(rotate) * Transform(scale); } -Bounds3f AnimatedTransform::MotionBounds(const Bounds3f &b) const { +PBRT_CPU_GPU Bounds3f AnimatedTransform::MotionBounds(const Bounds3f &b) const { // Handle easy cases for _Bounds3f_ motion bounds if (!actuallyAnimated) return startTransform(b); @@ -1094,7 +1094,7 @@ Bounds3f AnimatedTransform::MotionBounds(const Bounds3f &b) const { return bounds; } -Bounds3f AnimatedTransform::BoundPointMotion(Point3f p) const { +PBRT_CPU_GPU Bounds3f AnimatedTransform::BoundPointMotion(Point3f p) const { if (!actuallyAnimated) return Bounds3f(startTransform(p)); Bounds3f bounds(startTransform(p), endTransform(p)); @@ -1117,7 +1117,7 @@ Bounds3f AnimatedTransform::BoundPointMotion(Point3f p) const { return bounds; } -void AnimatedTransform::FindZeros(Float c1, Float c2, Float c3, Float c4, Float c5, +PBRT_CPU_GPU void AnimatedTransform::FindZeros(Float c1, Float c2, Float c3, Float c4, Float c5, Float theta, Interval tInterval, pstd::span zeros, int *nZeros, int depth) { // Evaluate motion derivative in interval form, return if no zeros diff --git a/src/pbrt/util/transform.h b/src/pbrt/util/transform.h index 8e5b17d42..abd9fcb66 100644 --- a/src/pbrt/util/transform.h +++ b/src/pbrt/util/transform.h @@ -269,7 +269,7 @@ PBRT_CPU_GPU inline Transform RotateFromTo(Vector3f from, Vector3f to) { return Transform(r, Transpose(r)); } -inline Vector3fi Transform::operator()(const Vector3fi &v) const { +PBRT_CPU_GPU inline Vector3fi Transform::operator()(const Vector3fi &v) const { Float x = Float(v.x), y = Float(v.y), z = Float(v.z); Vector3f vOutError; if (v.IsExact()) { @@ -307,7 +307,7 @@ inline Vector3fi Transform::operator()(const Vector3fi &v) const { // Transform Inline Methods template -inline Point3 Transform::operator()(Point3 p) const { +PBRT_CPU_GPU inline Point3 Transform::operator()(Point3 p) const { T xp = m[0][0] * p.x + m[0][1] * p.y + m[0][2] * p.z + m[0][3]; T yp = m[1][0] * p.x + m[1][1] * p.y + m[1][2] * p.z + m[1][3]; T zp = m[2][0] * p.x + m[2][1] * p.y + m[2][2] * p.z + m[2][3]; @@ -319,21 +319,21 @@ inline Point3 Transform::operator()(Point3 p) const { } template -inline Vector3 Transform::operator()(Vector3 v) const { +PBRT_CPU_GPU inline Vector3 Transform::operator()(Vector3 v) const { return Vector3(m[0][0] * v.x + m[0][1] * v.y + m[0][2] * v.z, m[1][0] * v.x + m[1][1] * v.y + m[1][2] * v.z, m[2][0] * v.x + m[2][1] * v.y + m[2][2] * v.z); } template -inline Normal3 Transform::operator()(Normal3 n) const { +PBRT_CPU_GPU inline Normal3 Transform::operator()(Normal3 n) const { T x = n.x, y = n.y, z = n.z; return Normal3(mInv[0][0] * x + mInv[1][0] * y + mInv[2][0] * z, mInv[0][1] * x + mInv[1][1] * y + mInv[2][1] * z, mInv[0][2] * x + mInv[1][2] * y + mInv[2][2] * z); } -inline Ray Transform::operator()(const Ray &r, Float *tMax) const { +PBRT_CPU_GPU inline Ray Transform::operator()(const Ray &r, Float *tMax) const { Point3fi o = (*this)(Point3fi(r.o)); Vector3f d = (*this)(r.d); // Offset ray origin to edge of error bounds and compute _tMax_ @@ -347,7 +347,7 @@ inline Ray Transform::operator()(const Ray &r, Float *tMax) const { return Ray(Point3f(o), d, r.time, r.medium); } -inline RayDifferential Transform::operator()(const RayDifferential &r, +PBRT_CPU_GPU inline RayDifferential Transform::operator()(const RayDifferential &r, Float *tMax) const { Ray tr = (*this)(Ray(r), tMax); RayDifferential ret(tr.o, tr.d, tr.time, tr.medium); @@ -359,12 +359,12 @@ inline RayDifferential Transform::operator()(const RayDifferential &r, return ret; } -inline Transform::Transform(const Frame &frame) +PBRT_CPU_GPU inline Transform::Transform(const Frame &frame) : Transform(SquareMatrix<4>(frame.x.x, frame.x.y, frame.x.z, 0, frame.y.x, frame.y.y, frame.y.z, 0, frame.z.x, frame.z.y, frame.z.z, 0, 0, 0, 0, 1)) {} -inline Transform::Transform(Quaternion q) { +PBRT_CPU_GPU inline Transform::Transform(Quaternion q) { Float xx = q.v.x * q.v.x, yy = q.v.y * q.v.y, zz = q.v.z * q.v.z; Float xy = q.v.x * q.v.y, xz = q.v.x * q.v.z, yz = q.v.y * q.v.z; Float wx = q.v.x * q.w, wy = q.v.y * q.w, wz = q.v.z * q.w; @@ -384,7 +384,7 @@ inline Transform::Transform(Quaternion q) { } template -inline Point3 Transform::ApplyInverse(Point3 p) const { +PBRT_CPU_GPU inline Point3 Transform::ApplyInverse(Point3 p) const { T x = p.x, y = p.y, z = p.z; T xp = (mInv[0][0] * x + mInv[0][1] * y) + (mInv[0][2] * z + mInv[0][3]); T yp = (mInv[1][0] * x + mInv[1][1] * y) + (mInv[1][2] * z + mInv[1][3]); @@ -398,7 +398,7 @@ inline Point3 Transform::ApplyInverse(Point3 p) const { } template -inline Vector3 Transform::ApplyInverse(Vector3 v) const { +PBRT_CPU_GPU inline Vector3 Transform::ApplyInverse(Vector3 v) const { T x = v.x, y = v.y, z = v.z; return Vector3(mInv[0][0] * x + mInv[0][1] * y + mInv[0][2] * z, mInv[1][0] * x + mInv[1][1] * y + mInv[1][2] * z, @@ -406,14 +406,14 @@ inline Vector3 Transform::ApplyInverse(Vector3 v) const { } template -inline Normal3 Transform::ApplyInverse(Normal3 n) const { +PBRT_CPU_GPU inline Normal3 Transform::ApplyInverse(Normal3 n) const { T x = n.x, y = n.y, z = n.z; return Normal3(m[0][0] * x + m[1][0] * y + m[2][0] * z, m[0][1] * x + m[1][1] * y + m[2][1] * z, m[0][2] * x + m[1][2] * y + m[2][2] * z); } -inline Ray Transform::ApplyInverse(const Ray &r, Float *tMax) const { +PBRT_CPU_GPU inline Ray Transform::ApplyInverse(const Ray &r, Float *tMax) const { Point3fi o = ApplyInverse(Point3fi(r.o)); Vector3f d = ApplyInverse(r.d); // Offset ray origin to edge of error bounds and compute _tMax_ @@ -428,7 +428,7 @@ inline Ray Transform::ApplyInverse(const Ray &r, Float *tMax) const { return Ray(Point3f(o), d, r.time, r.medium); } -inline RayDifferential Transform::ApplyInverse(const RayDifferential &r, +PBRT_CPU_GPU inline RayDifferential Transform::ApplyInverse(const RayDifferential &r, Float *tMax) const { Ray tr = ApplyInverse(Ray(r), tMax); RayDifferential ret(tr.o, tr.d, tr.time, tr.medium); diff --git a/src/pbrt/util/vecmath.cpp b/src/pbrt/util/vecmath.cpp index d4b587aa5..31e0f05ad 100644 --- a/src/pbrt/util/vecmath.cpp +++ b/src/pbrt/util/vecmath.cpp @@ -53,7 +53,7 @@ std::string Quaternion::ToString() const { } // DirectionCone Function Definitions -DirectionCone Union(const DirectionCone &a, const DirectionCone &b) { +PBRT_CPU_GPU DirectionCone Union(const DirectionCone &a, const DirectionCone &b) { // Handle the cases where one or both cones are empty if (a.IsEmpty()) return b; diff --git a/src/pbrt/util/vecmath.h b/src/pbrt/util/vecmath.h index 60f691dbc..09cedb139 100644 --- a/src/pbrt/util/vecmath.h +++ b/src/pbrt/util/vecmath.h @@ -879,7 +879,7 @@ class Quaternion { // Vector2 Inline Functions template template -Vector2::Vector2(Point2 p) : Tuple2(T(p.x), T(p.y)) {} +PBRT_CPU_GPU Vector2::Vector2(Point2 p) : Tuple2(T(p.x), T(p.y)) {} template PBRT_CPU_GPU inline auto Dot(Vector2 v1, Vector2 v2) -> @@ -926,7 +926,7 @@ PBRT_CPU_GPU inline auto DistanceSquared(Point2 p1, Point2 p2) -> // Vector3 Inline Functions template template -Vector3::Vector3(Point3 p) : Tuple3(T(p.x), T(p.y), T(p.z)) {} +PBRT_CPU_GPU Vector3::Vector3(Point3 p) : Tuple3(T(p.x), T(p.y), T(p.z)) {} template PBRT_CPU_GPU inline Vector3 Cross(Vector3 v1, Normal3 v2) { @@ -1023,7 +1023,7 @@ PBRT_CPU_GPU inline void CoordinateSystem(Normal3 v1, Vector3 *v2, Vector3 template template -Vector3::Vector3(Normal3 n) : Tuple3(T(n.x), T(n.y), T(n.z)) {} +PBRT_CPU_GPU Vector3::Vector3(Normal3 n) : Tuple3(T(n.x), T(n.y), T(n.z)) {} // Point3 Inline Functions template @@ -1926,7 +1926,7 @@ class Frame { }; // Frame Inline Functions -inline Frame::Frame(Vector3f x, Vector3f y, Vector3f z) : x(x), y(y), z(z) { +PBRT_CPU_GPU inline Frame::Frame(Vector3f x, Vector3f y, Vector3f z) : x(x), y(y), z(z) { DCHECK_LT(std::abs(LengthSquared(x) - 1), 1e-4); DCHECK_LT(std::abs(LengthSquared(y) - 1), 1e-4); DCHECK_LT(std::abs(LengthSquared(z) - 1), 1e-4); diff --git a/src/pbrt/wavefront/workitems.h b/src/pbrt/wavefront/workitems.h index a4befe04e..a5ed18adf 100644 --- a/src/pbrt/wavefront/workitems.h +++ b/src/pbrt/wavefront/workitems.h @@ -343,7 +343,7 @@ class RayQueue : public WorkQueue { }; // RayQueue Inline Methods -inline int RayQueue::PushCameraRay(const Ray &ray, const SampledWavelengths &lambda, +PBRT_CPU_GPU inline int RayQueue::PushCameraRay(const Ray &ray, const SampledWavelengths &lambda, int pixelIndex) { int index = AllocateEntry(); DCHECK(!ray.HasNaN()); @@ -397,7 +397,7 @@ class EscapedRayQueue : public WorkQueue { using WorkQueue::Push; }; -inline int EscapedRayQueue::Push(RayWorkItem r) { +PBRT_CPU_GPU inline int EscapedRayQueue::Push(RayWorkItem r) { return Push(EscapedRayWorkItem{r.ray.o, r.ray.d, r.depth, r.lambda, r.pixelIndex, r.beta, (int)r.specularBounce, r.r_u, r.r_l, r.prevIntrCtx});