#include "Sampler.h" #include "thread_index.h" __device__ glm::vec3 _raySphereIntersect(glm::vec3 p, glm::vec3 v, float r, float &o_depth) { float pp = glm::dot(p, p); float vv = glm::dot(v, v); float pv = glm::dot(p, v); o_depth = (sqrtf(pv * pv - vv * (pp - r * r)) - pv) / vv; return p + o_depth * v; } __device__ float _getAngle(float x, float y) { return -atan(x / y) + (y < 0) * (float)M_PI + 0.5f * (float)M_PI; } /** * Dispatch with block_size=(n_samples, 1024) */ __global__ void cu_sampleOnRays(glm::vec3 *o_sphericalCoords, float *o_depths, glm::vec3 *rays, glm::vec3 rayCenter, float range0, float rangeStep, uint n) { glm::uvec3 idx3 = IDX3; uint rayIdx = flattenIdx({idx3.y, idx3.z, 0}); if (rayIdx >= n) return; uint idx = flattenIdx(idx3); float r_reciprocal = rangeStep * idx3.x + range0; glm::vec3 p = _raySphereIntersect(rayCenter, rays[rayIdx], 1.0f / r_reciprocal, o_depths[idx]); o_sphericalCoords[idx] = glm::vec3(r_reciprocal, _getAngle(p.x, p.z), acos(p.y * r_reciprocal)); } void Sampler::sampleOnRays(sptr> o_sphericalCoords, sptr> o_depths, sptr> rays, glm::vec3 rayCenter) { dim3 blockSize(_samples, 1024 / _samples); dim3 gridSize(1, (uint)ceil(rays->n() / (float)blockSize.y)); cu_sampleOnRays<<>>(*o_sphericalCoords, *o_depths, *rays, rayCenter, _dispRange.x, (_dispRange.y - _dispRange.x) / (_samples - 1), rays->n()); CHECK_EX(cudaGetLastError()); }