AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Merge branch 'master' into move-memcpy

Telusuri Sumber
This commit is contained in:
Maneesh Gupta
2018-07-17 10:51:42 +05:30
committed by GitHub
orang tua 2531815dfe 753aa1fbff
melakukan 1fc3ef3cd4
35 mengubah file dengan 2155 tambahan dan 16857 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/device_util.cpp
-63
Melihat File
@@ -92,70 +92,7 @@ __device__ void* __hip_hc_free(void* ptr) {
return nullptr;
}
__device__ long long int clock64() { return (long long int)hc::__cycle_u64(); };
__device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); };
// abort
__device__ void abort() { return hc::abort(); }
// warp vote function __all __any __ballot
__device__ int __all(int input) { return hc::__all(input); }
__device__ int __any(int input) {
#ifdef NVCC_COMPAT
if (hc::__any(input) != 0)
return 1;
else
return 0;
#else
return hc::__any(input);
#endif
}
__device__ unsigned long long int __ballot(int input) { return hc::__ballot(input); }
// warp shuffle functions
__device__ int __shfl(int input, int lane, int width) { return hc::__shfl(input, lane, width); }
__device__ int __shfl_up(int input, unsigned int lane_delta, int width) {
return hc::__shfl_up(input, lane_delta, width);
}
__device__ int __shfl_down(int input, unsigned int lane_delta, int width) {
return hc::__shfl_down(input, lane_delta, width);
}
__device__ int __shfl_xor(int input, int lane_mask, int width) {
return hc::__shfl_xor(input, lane_mask, width);
}
__device__ float __shfl(float input, int lane, int width) { return hc::__shfl(input, lane, width); }
__device__ float __shfl_up(float input, unsigned int lane_delta, int width) {
return hc::__shfl_up(input, lane_delta, width);
}
__device__ float __shfl_down(float input, unsigned int lane_delta, int width) {
return hc::__shfl_down(input, lane_delta, width);
}
__device__ float __shfl_xor(float input, int lane_mask, int width) {
return hc::__shfl_xor(input, lane_mask, width);
}
__host__ __device__ int min(int arg1, int arg2) {
return (int)(hc::precise_math::fmin((float)arg1, (float)arg2));
}
__host__ __device__ int max(int arg1, int arg2) {
return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));
}
__device__ void* __get_dynamicgroupbaseptr() {
return hc::get_dynamic_group_segment_base_pointer();
}
__host__ void* __get_dynamicgroupbaseptr() { return nullptr; }
__device__ void __threadfence_system(void) { std::atomic_thread_fence(std::memory_order_seq_cst); }