diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp index 85620ff287c84ac9e8f5828784a34a8a917d48f6..4cbd760a1b0ef1e01c111dda2ae66c352cf6de30 100644 --- a/components/renderers/cpp/include/ftl/cuda/normals.hpp +++ b/components/renderers/cpp/include/ftl/cuda/normals.hpp @@ -6,7 +6,7 @@ namespace ftl { namespace cuda { -void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output, +void normals(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float4> &input, cudaStream_t stream); } diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu index 813f1aab543b81fdbcb5bdde3471ab3b0c39b0ae..7ed9bb6816972b90dd63533f80bc1b56aa685c47 100644 --- a/components/renderers/cpp/src/normals.cu +++ b/components/renderers/cpp/src/normals.cu @@ -1,6 +1,7 @@ #include <ftl/cuda/normals.hpp> #define T_PER_BLOCK 16 +#define MINF __int_as_float(0xff800000) __global__ void computeNormals_kernel(ftl::cuda::TextureObject<float4> output, ftl::cuda::TextureObject<float4> input) { @@ -12,11 +13,11 @@ __global__ void computeNormals_kernel(ftl::cuda::TextureObject<float4> output, output(x,y) = make_float4(MINF, MINF, MINF, MINF); if(x > 0 && x < input.width()-1 && y > 0 && y < input.height()-1) { - const float3 CC = make_float3(input.tex2D(x+0, y+0)); //[(y+0)*width+(x+0)]; - const float3 PC = make_float3(input.tex2D(x+0, y+1)); //[(y+1)*width+(x+0)]; - const float3 CP = make_float3(input.tex2D(x+1, y+0)); //[(y+0)*width+(x+1)]; - const float3 MC = make_float3(input.tex2D(x+0, y-1)); //[(y-1)*width+(x+0)]; - const float3 CM = make_float3(input.tex2D(x-1, y+0)); //[(y+0)*width+(x-1)]; + const float3 CC = make_float3(input.tex2D((int)x+0, (int)y+0)); //[(y+0)*width+(x+0)]; + const float3 PC = make_float3(input.tex2D((int)x+0, (int)y+1)); //[(y+1)*width+(x+0)]; + const float3 CP = make_float3(input.tex2D((int)x+1, (int)y+0)); //[(y+0)*width+(x+1)]; + const float3 MC = make_float3(input.tex2D((int)x+0, (int)y-1)); //[(y-1)*width+(x+0)]; + const float3 CM = make_float3(input.tex2D((int)x-1, (int)y+0)); //[(y+0)*width+(x-1)]; if(CC.x != MINF && PC.x != MINF && CP.x != MINF && MC.x != MINF && CM.x != MINF) { const float3 n = cross(PC-MC, CP-CM);