133   c2 rv = __nvvm_ldg_c2(reinterpret_cast<const c2 *>(ptr));  in __ldg()  local
135   ret.x = rv[0];  in __ldg()
136   ret.y = rv[1];  in __ldg()
141   c4 rv = __nvvm_ldg_c4(reinterpret_cast<const c4 *>(ptr));  in __ldg()  local
143   ret.x = rv[0];  in __ldg()
144   ret.y = rv[1];  in __ldg()
145   ret.z = rv[2];  in __ldg()
146   ret.w = rv[3];  in __ldg()
151   s2 rv = __nvvm_ldg_s2(reinterpret_cast<const s2 *>(ptr));  in __ldg()  local
153   ret.x = rv[0];  in __ldg()
154   ret.y = rv[1];  in __ldg()
159   s4 rv = __nvvm_ldg_s4(reinterpret_cast<const s4 *>(ptr));  in __ldg()  local
161   ret.x = rv[0];  in __ldg()
162   ret.y = rv[1];  in __ldg()
163   ret.z = rv[2];  in __ldg()
164   ret.w = rv[3];  in __ldg()
169   i2 rv = __nvvm_ldg_i2(reinterpret_cast<const i2 *>(ptr));  in __ldg()  local
171   ret.x = rv[0];  in __ldg()
172   ret.y = rv[1];  in __ldg()
177   i4 rv = __nvvm_ldg_i4(reinterpret_cast<const i4 *>(ptr));  in __ldg()  local
179   ret.x = rv[0];  in __ldg()
180   ret.y = rv[1];  in __ldg()
181   ret.z = rv[2];  in __ldg()
182   ret.w = rv[3];  in __ldg()
187   ll2 rv = __nvvm_ldg_ll2(reinterpret_cast<const ll2 *>(ptr));  in __ldg()  local
189   ret.x = rv[0];  in __ldg()
190   ret.y = rv[1];  in __ldg()
196   uc2 rv = __nvvm_ldg_uc2(reinterpret_cast<const uc2 *>(ptr));  in __ldg()  local
198   ret.x = rv[0];  in __ldg()
199   ret.y = rv[1];  in __ldg()
204   uc4 rv = __nvvm_ldg_uc4(reinterpret_cast<const uc4 *>(ptr));  in __ldg()  local
206   ret.x = rv[0];  in __ldg()
207   ret.y = rv[1];  in __ldg()
208   ret.z = rv[2];  in __ldg()
209   ret.w = rv[3];  in __ldg()
214   us2 rv = __nvvm_ldg_us2(reinterpret_cast<const us2 *>(ptr));  in __ldg()  local
216   ret.x = rv[0];  in __ldg()
217   ret.y = rv[1];  in __ldg()
222   us4 rv = __nvvm_ldg_us4(reinterpret_cast<const us4 *>(ptr));  in __ldg()  local
224   ret.x = rv[0];  in __ldg()
225   ret.y = rv[1];  in __ldg()
226   ret.z = rv[2];  in __ldg()
227   ret.w = rv[3];  in __ldg()
232   ui2 rv = __nvvm_ldg_ui2(reinterpret_cast<const ui2 *>(ptr));  in __ldg()  local
234   ret.x = rv[0];  in __ldg()
235   ret.y = rv[1];  in __ldg()
240   ui4 rv = __nvvm_ldg_ui4(reinterpret_cast<const ui4 *>(ptr));  in __ldg()  local
242   ret.x = rv[0];  in __ldg()
243   ret.y = rv[1];  in __ldg()
244   ret.z = rv[2];  in __ldg()
245   ret.w = rv[3];  in __ldg()
250   ull2 rv = __nvvm_ldg_ull2(reinterpret_cast<const ull2 *>(ptr));  in __ldg()  local
252   ret.x = rv[0];  in __ldg()
253   ret.y = rv[1];  in __ldg()
259   f2 rv = __nvvm_ldg_f2(reinterpret_cast<const f2 *>(ptr));  in __ldg()  local
261   ret.x = rv[0];  in __ldg()
262   ret.y = rv[1];  in __ldg()
267   f4 rv = __nvvm_ldg_f4(reinterpret_cast<const f4 *>(ptr));  in __ldg()  local
269   ret.x = rv[0];  in __ldg()
270   ret.y = rv[1];  in __ldg()
271   ret.z = rv[2];  in __ldg()
272   ret.w = rv[3];  in __ldg()
277   d2 rv = __nvvm_ldg_d2(reinterpret_cast<const d2 *>(ptr));  in __ldg()  local
279   ret.x = rv[0];  in __ldg()
280   ret.y = rv[1];  in __ldg()