Commit eadbf875 authored by yangzhong's avatar yangzhong
Browse files

修复版本兼容报错__ldg问题

parent c7c514c2
......@@ -17,3 +17,15 @@ __device__ __inline__ at::Half __shfl_down_sync(const unsigned mask,
const unsigned int delta) {
return __shfl_down_sync(mask, (__half)var, delta);
}
#ifdef USE_ROCM
__device__ __inline__ at::Half __ldg(const at::Half* ptr) {
return __ldg(reinterpret_cast<const __half*>(ptr));
}
#define SHFL_UP_SYNC(mask, var, delta) __shfl_up(var, delta)
#define SHFL_DOWN_SYNC(mask, var, delta) __shfl_down(var, delta)
#else
#define SHFL_UP_SYNC __shfl_up_sync
#define SHFL_DOWN_SYNC __shfl_down_sync
#endif
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment