summaryrefslogtreecommitdiff
path: root/Build/source/utils/asymptote/cudareflect/kernel.cu
diff options
context:
space:
mode:
Diffstat (limited to 'Build/source/utils/asymptote/cudareflect/kernel.cu')
-rw-r--r--Build/source/utils/asymptote/cudareflect/kernel.cu148
1 files changed, 148 insertions, 0 deletions
diff --git a/Build/source/utils/asymptote/cudareflect/kernel.cu b/Build/source/utils/asymptote/cudareflect/kernel.cu
new file mode 100644
index 00000000000..9104ed3b64d
--- /dev/null
+++ b/Build/source/utils/asymptote/cudareflect/kernel.cu
@@ -0,0 +1,148 @@
+/**
+* @file kernel.cu
+* @author Supakorn "Jamie" Rassameemasmuang <jamievlin@outlook.com>
+* CUDA Kernel for computing irradiance by solid angle integration
+
+* Partially based on:
+* https://cdn2.unrealengine.com/Resources/files/2013SiggraphPresentationsNotes-26915738.pdf
+*/
+
+#include "kernel.h"
+#include "helper.cuh"
+#include "utils.cuh"
+
+#include "simpson.cuh"
+
+#include <cuda.h>
+#include <texture_indirect_functions.h>
+#include <device_launch_parameters.h>
+
+#define GLM_FORCE_CUDA
+#include <glm/glm.hpp>
+
+#include <functional>
+
+class IntegrateSampler
+{
+public:
+ __device__
+ IntegrateSampler(
+ cudaTextureObject_t tObjin,
+ glm::mat3 normalOrthBasis,
+ size_t const& inWidth, size_t const& inHeight) :
+ normalOrthBasis(normalOrthBasis), width(inWidth), height(inHeight),
+ tObj(tObjin)
+
+ {
+ }
+
+ __device__ ~IntegrateSampler() {}
+
+ __device__
+ glm::vec3 integrand(float const& sampled_phi, float const& sampled_theta)
+ {
+ // vec3 is the world space coordinate
+ glm::vec2 sphcoord = to_sphcoord(angleToBasis(normalOrthBasis, sampled_phi, sampled_theta));
+ float4 frag = tex2D<float4>(tObj,
+ sphcoord.x * PI_RECR * 0.5*width,
+ sphcoord.y * PI_RECR * height);
+
+ return glm::vec3(frag.x, frag.y, frag.z);
+ }
+
+ __device__
+ glm::vec3 inner(float const& sampled_theta)
+ {
+ return simpson(
+ [this, &sampled_theta](float const& phi) {return this->integrand(phi,sampled_theta); },
+ 0, TAU, acc)*0.5f*__sinf(2 * sampled_theta);
+ }
+
+ __device__
+ glm::vec3 integrate()
+ {
+ return PI_RECR * simpson(
+ [this](float const& theta) {return this->inner(theta); },
+ 0, HALFPI, acc);
+ }
+
+private:
+ glm::mat3 normalOrthBasis;
+ size_t width, height;
+ cudaTextureObject_t tObj;
+};
+
+
+
+__global__
+void irradiate(cudaTextureObject_t tObjin, float3* out, size_t width, size_t height)
+{
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ int idx_y = blockIdx.y * blockDim.y + threadIdx.y;
+
+
+ if (idx < width && idx_y < height)
+ {
+ int access_idx = to_idx(width, idx, idx_y);
+
+ float target_phi = TAU * ((idx + 0.5f) / width);
+ float target_theta = PI * ((idx_y + 0.5f) / height);
+
+ const glm::vec3 N = from_sphcoord(target_phi, target_theta);
+ const glm::vec3 N1(
+ __cosf(target_theta) * __cosf(target_phi),
+ __cosf(target_theta) * __sinf(target_phi),
+ -1*__sinf(target_theta));
+ const glm::vec3 N2(-1 * __sinf(target_phi), __cosf(target_phi), 0);
+
+ glm::mat3 normalBasisMat(N1,N2,N);
+
+ IntegrateSampler integrator(tObjin, normalBasisMat, width, height);
+ glm::vec3 out_val = integrator.integrate();
+ out[access_idx] = make_float3(out_val.x, out_val.y, out_val.z);
+ }
+}
+
+const size_t blkSz = 8;
+void irradiate_ker(float4* in, float3* out, size_t width, size_t height)
+{
+ float4* d_ptr;
+ size_t pitch;
+ cudaErrorCheck(cudaMallocPitch(
+ &d_ptr, &pitch, width * sizeof(float4), height));
+ cudaErrorCheck(cudaMemcpy2D(d_ptr, pitch, in,
+ width * sizeof(float4), width*sizeof(float4),
+ height, cudaMemcpyHostToDevice));
+
+ cudaResourceDesc cRD;
+ memset(&cRD, 0, sizeof(cudaResourceDesc));
+ cRD.resType = cudaResourceTypePitch2D;
+ cRD.res.pitch2D.devPtr = d_ptr;
+ cRD.res.pitch2D.width = width;
+ cRD.res.pitch2D.height = height;
+ cRD.res.pitch2D.desc = cudaCreateChannelDesc<float4>();
+ cRD.res.pitch2D.pitchInBytes = pitch;
+
+ cudaTextureDesc texDesc;
+ memset(&texDesc, 0, sizeof(texDesc));
+ texDesc.filterMode = cudaFilterModeLinear;
+ texDesc.sRGB = 0;
+ texDesc.readMode = cudaReadModeElementType;
+
+ cudaTextureObject_t t_obj;
+ cudaErrorCheck(cudaCreateTextureObject(
+ &t_obj, &cRD, &texDesc, nullptr));
+
+ // out source
+ float3* d_out;
+ cudaErrorCheck(cudaMalloc(
+ (void**)&d_out, static_cast<size_t>(sizeof(float3) * width * height)));
+ dim3 blockSz((width / blkSz) + 1, (height / blkSz) + 1);
+ dim3 kerSz(blkSz, blkSz);
+ irradiate KERNEL_ARGS(blockSz, kerSz) (t_obj, d_out, width, height);
+
+ cudaErrorCheck(cudaMemcpy(
+ out, d_out, sizeof(float3) * width * height, cudaMemcpyDeviceToHost));
+ cudaErrorCheck(cudaDestroyTextureObject(t_obj));
+ cudaFree(d_ptr);
+}