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

Added architecture guards around __shfl, dpp and ds_permute device functions

Telusuri Sumber 
Change-Id: I10f9b08618fbf25b61c1932278fc5759e41c0d66


[ROCm/hip commit: 1567d20aa8]
This commit is contained in:
Aditya Atluri
2017-03-10 08:40:59 -06:00
orang tua 87dab4f207
melakukan b0c2ac0a90
1 mengubah file dengan 14 tambahan dan 10 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/hip_runtime.h
+14 -10
Melihat File
@@ -226,6 +226,8 @@ __device__ int __all( int input);
__device__ int __any( int input);
__device__ unsigned long long int __ballot( int input);
#if __HIP_ARCH_GFX803__ == 1
// warp shuffle functions
#ifdef __cplusplus
__device__ int __shfl(int input, int lane, int width=warpSize);
@@ -247,6 +249,18 @@ __device__ float __shfl_down(float input, unsigned int lane_delta, int width);
__device__ float __shfl_xor(float input, int lane_mask, int width);
#endif
__device__ unsigned __hip_ds_bpermute(int index, unsigned src);
__device__ float __hip_ds_bpermutef(int index, float src);
__device__ unsigned __hip_ds_permute(int index, unsigned src);
__device__ float __hip_ds_permutef(int index, float src);
__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern);
__device__ float __hip_ds_swizzlef(float src, int pattern);
__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl);
#endif
__host__ __device__ int min(int arg1, int arg2);
__host__ __device__ int max(int arg1, int arg2);
@@ -321,16 +335,6 @@ __device__ static inline void __threadfence(void) {
//__device__ void __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
__device__ void __threadfence_system(void) ;
__device__ unsigned __hip_ds_bpermute(int index, unsigned src);
__device__ float __hip_ds_bpermutef(int index, float src);
__device__ unsigned __hip_ds_permute(int index, unsigned src);
__device__ float __hip_ds_permutef(int index, float src);
__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern);
__device__ float __hip_ds_swizzlef(float src, int pattern);
__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl);
// doxygen end Fence Fence
/**
* @}