Fermat
texture_view.h
1 /*
2  * Fermat
3  *
4  * Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
5  *
6  * Redistribution and use in source and binary forms, with or without
7  * modification, are permitted provided that the following conditions are met:
8  * * Redistributions of source code must retain the above copyright
9  * notice, this list of conditions and the following disclaimer.
10  * * Redistributions in binary form must reproduce the above copyright
11  * notice, this list of conditions and the following disclaimer in the
12  * documentation and/or other materials provided with the distribution.
13  * * Neither the name of the NVIDIA CORPORATION nor the
14  * names of its contributors may be used to endorse or promote products
15  * derived from this software without specific prior written permission.
16  *
17  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
18  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
19  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
21  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
22  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
24  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
26  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27  */
28 
29 #pragma once
30 
31 #include "types.h"
32 #include "texture_reference.h"
33 
34 #include <cugar/basic/types.h>
35 #include <cugar/linalg/vector.h>
36 
37 #define BILINEAR_TEXTURE_LOOKUPS 1
38 
41 
44 
45 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
46 float4 texture_load(const float4* tex)
47 {
48 #ifdef FERMAT_DEVICE_COMPILATION
49  return __ldg(tex);
50 #else
51  return *tex;
52 #endif
53 }
54 
57 struct TextureView
58 {
59  FERMAT_HOST_DEVICE float4& operator()(const uint32 pixel) { return c[pixel]; }
60  FERMAT_HOST_DEVICE const float4& operator()(const uint32 pixel) const { return texture_load(c + pixel); }
61  FERMAT_HOST_DEVICE float4& operator()(const uint32 x, const uint32 y) { return c[y*res_x + x]; }
62  FERMAT_HOST_DEVICE const float4& operator()(const uint32 x, const uint32 y) const { return texture_load(c + y*res_x + x); }
63  FERMAT_HOST_DEVICE const float4* ptr() const { return c; }
64  FERMAT_HOST_DEVICE float4* ptr() { return c; }
65 
66  float4* c;
67  uint32 res_x;
68  uint32 res_y;
69 };
70 
73 struct MipMapView
74 {
75  FERMAT_HOST_DEVICE float4& operator()(const uint32 pixel, const uint32 lod) { return levels[lod](pixel); }
76  FERMAT_HOST_DEVICE const float4& operator()(const uint32 pixel, const uint32 lod) const { return reinterpret_cast<const TextureView*>(levels)[lod](pixel); }
77  FERMAT_HOST_DEVICE float4& operator()(const uint32 x, const uint32 y, const uint32 lod) { return levels[lod](x,y); }
78  FERMAT_HOST_DEVICE const float4& operator()(const uint32 x, const uint32 y, const uint32 lod) const { return reinterpret_cast<const TextureView*>(levels)[lod](x,y); }
79 
80  TextureView* levels;
81  uint32 n_levels;
82  uint32 res_x;
83  uint32 res_y;
84 };
85 
86 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
87 cugar::Vector4f unfiltered_texture_lookup(float4 st, const TextureReference texture_ref, const MipMapView* textures, const float4 default_value)
88 {
89  if (texture_ref.texture == uint32(-1) || textures[texture_ref.texture].n_levels == 0)
90  return default_value;
91 
92  st.x *= texture_ref.scaling.x;
93  st.y *= texture_ref.scaling.y;
94 
95  st.x = cugar::mod(st.x, 1.0f);
96  st.y = cugar::mod(st.y, 1.0f);
97 
98  const MipMapView texture = textures[texture_ref.texture];
99  if (texture.n_levels > 0)
100  {
101  const uint32 x = cugar::min(uint32(st.x * texture.res_x), texture.res_x - 1);
102  const uint32 y = cugar::min(uint32(st.y * texture.res_y), texture.res_y - 1);
103  return texture(x, y, 0); // access LOD 0
104  }
105  else
106  return default_value;
107 }
108 
109 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
110 void unchecked_quad_texture_lookup(uint2 texel, const MipMapView texture, const uint32 lod, const float4 default_value, cugar::Vector4f quad[4])
111 {
112  const uint32 x = texel.x;
113  const uint32 y = texel.y;
114  const uint32 xx = (x + 1) % texture.res_x;
115  const uint32 yy = (y + 1) % texture.res_y;
116 
117  quad[0] = texture(x, y, lod);
118  quad[1] = texture(xx, y, lod);
119  quad[2] = texture(x, yy, lod);
120  quad[3] = texture(xx, yy, lod);
121 }
122 
123 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
124 void quad_texture_lookup(uint2 texel, const TextureReference texture_ref, const MipMapView* textures, const uint32 lod, const float4 default_value, cugar::Vector4f quad[4])
125 {
126  if (texture_ref.texture == uint32(-1) || textures[texture_ref.texture].n_levels == 0)
127  {
128  quad[0] = quad[1] = quad[2] = quad[3] = default_value;
129  return;
130  }
131 
132  const MipMapView texture = textures[texture_ref.texture];
133  if (texture.n_levels > lod)
134  {
135  // fetch the quad at the given lod
136  unchecked_quad_texture_lookup( texel, texture, lod, default_value, quad );
137  }
138  else
139  quad[0] = quad[1] = quad[2] = quad[3] = default_value;
140 }
141 
142 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
143 cugar::Vector4f bilinear_texture_lookup_unscaled(cugar::Vector4f st, const TextureReference texture_ref, const MipMapView* textures, const float4 default_value)
144 {
145  if (texture_ref.texture == uint32(-1) || textures[texture_ref.texture].n_levels == 0)
146  return default_value;
147 
148  const MipMapView texture = textures[texture_ref.texture];
149  if (texture.n_levels > 0)
150  {
151  const uint32 x = uint32(st.x) % texture.res_x;
152  const uint32 y = uint32(st.y) % texture.res_y;
153 
154  // fetch the quad at LOD 0
155  cugar::Vector4f quad[4];
156  const uint32 lod = 0u;
157  unchecked_quad_texture_lookup( make_uint2(x,y), texture, lod, default_value, quad );
158 
159  const float u = cugar::mod(st.x, 1.0f);
160  const float v = cugar::mod(st.y, 1.0f);
161 
162  return
163  (quad[0] * (1-u) + quad[1] * u) * (1-v) +
164  (quad[2] * (1-u) + quad[3] * u) * v;
165  }
166  else
167  return default_value;
168 }
169 
170 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
171 cugar::Vector4f bilinear_texture_lookup(cugar::Vector4f st, const TextureReference texture_ref, const MipMapView* textures, const float4 default_value)
172 {
173  if (texture_ref.texture == uint32(-1) || textures[texture_ref.texture].n_levels == 0)
174  return default_value;
175 
176  st.x *= texture_ref.scaling.x;
177  st.y *= texture_ref.scaling.y;
178 
179  st.x = cugar::mod(st.x, 1.0f);
180  st.y = cugar::mod(st.y, 1.0f);
181 
182  const MipMapView texture = textures[texture_ref.texture];
183  if (texture.n_levels > 0)
184  {
185  const uint32 x = cugar::min(uint32(st.x * texture.res_x), texture.res_x - 1);
186  const uint32 y = cugar::min(uint32(st.y * texture.res_y), texture.res_y - 1);
187 
188  // fetch the quad at LOD 0
189  cugar::Vector4f quad[4];
190  const uint32 lod = 0u;
191  unchecked_quad_texture_lookup( make_uint2(x,y), texture, lod, default_value, quad );
192 
193  const float u = cugar::mod(st.x * texture.res_x, 1.0f);
194  const float v = cugar::mod(st.y * texture.res_y, 1.0f);
195 
196  return
197  (quad[0] * (1-u) + quad[1] * u) * (1-v) +
198  (quad[2] * (1-u) + quad[3] * u) * v;
199  }
200  else
201  return default_value;
202 }
203 
204 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
205 cugar::Vector4f texture_lookup(float4 st, const TextureReference texture_ref, const MipMapView* textures, const float4 default_value)
206 {
207  #if BILINEAR_TEXTURE_LOOKUPS
208  return bilinear_texture_lookup( st, texture_ref, textures, default_value );
209  #else
210  return unfiltered_texture_lookup( st, texture_ref, textures, default_value );
211  #endif
212 }
213 
214 FERMAT_HOST_DEVICE FERMAT_FORCEINLINE
215 cugar::Vector2f diff_texture_lookup(cugar::Vector4f st, const TextureReference texture_ref, const MipMapView* textures, const float2 default_value)
216 {
217  if (texture_ref.texture == uint32(-1) || textures[texture_ref.texture].n_levels == 0)
218  return default_value;
219 
220  st.x *= texture_ref.scaling.x;
221  st.y *= texture_ref.scaling.y;
222 
223  st.x = cugar::mod(st.x, 1.0f);
224  st.y = cugar::mod(st.y, 1.0f);
225 
226  const MipMapView texture = textures[texture_ref.texture];
227  if (texture.n_levels > 0)
228  {
229  // unnormalize the coordinates
230  st *= cugar::Vector4f(float(texture.res_x), float(texture.res_y), 1, 1);
231 
232  cugar::Vector4f c = bilinear_texture_lookup_unscaled( st, texture_ref, textures, cugar::Vector4f(0.0f) );
233  cugar::Vector4f cu = bilinear_texture_lookup_unscaled( st + cugar::Vector4f(1,0,0,0), texture_ref, textures, cugar::Vector4f(0.0f) );
234  cugar::Vector4f cv = bilinear_texture_lookup_unscaled( st + cugar::Vector4f(0,1,0,0), texture_ref, textures, cugar::Vector4f(0.0f) );
235 
236  const float i = (c.x + c.y + c.z);
237  const float iu = (cu.x + cu.y + cu.z);
238  const float iv = (cv.x + cv.y + cv.z);
239  return cugar::Vector2f(iu - i, iv - i);
240  }
241  else
242  return default_value;
243 }
244 
TextureReference
Definition: texture_reference.h:41
cugar::mod
float CUGAR_HOST_DEVICE mod(const float x, const float m)
Definition: numbers.h:606
TextureView
Definition: texture_view.h:57
MipMapView
Definition: texture_view.h:73
Generated by   doxygen 1.8.13