| 149 | |
| 150 | template <typename T> |
| 151 | DEVICE_FUNCTION void ldgStream(int32_t (&dst)[1], T const* gmem) |
| 152 | { |
| 153 | uint32_t tmp; |
| 154 | asm volatile("ld.global.cs.nc.s32 %0, [%1];" : "=r"(tmp) : "l"((uint32_t const*) gmem)); |
| 155 | dst[0] = tmp; |
| 156 | } |
| 157 | |
| 158 | template <typename T> |
| 159 | DEVICE_FUNCTION void ldg(int32_t (&dst)[2], T const* gmem) |