1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
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);
}
|
