| 153 | } |
| 154 | |
| 155 | DEVICE static void invalidate(ValueType const *smem_ptr) { |
| 156 | uint32_t smem_addr = cast_smem_ptr_to_uint(smem_ptr); |
| 157 | asm volatile( |
| 158 | "{\n\t" |
| 159 | "mbarrier.ival.shared.b64 [%0]; \n\t" |
| 160 | "}" |
| 161 | : |
| 162 | : "r"(smem_addr)); |
| 163 | } |
| 164 | }; |
| 165 | |
| 166 | template <uint32_t Stages_> |
nothing calls this directly
no test coverage detected