AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

added inter thread data movement intrinsics

Browse Source 
Change-Id: I2a8a8ed49429cb7f96439bd28c4b83b5142737df


[ROCm/clr commit: f097b6ef81]
This commit is contained in:
Aditya Atluri
2016-11-01 16:37:33 -05:00
parent 5040e8bcc3
commit 2d299543bf
2 changed files with 35 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+9
View File
@@ -560,6 +560,15 @@ extern "C" __device__ void __threadfence(void);
*/
__device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
__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
/**
projects/clr/hipamd/src/hip_hcc.cpp
+26
View File
@@ -180,7 +180,33 @@ __device__ void* __hip_hc_free(void *ptr)
return nullptr;
}
__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
return hc::__amdgcn_ds_bpermute(index, src);
}
__device__ float __hip_ds_bpermutef(int index, float src) {
return hc::__amdgcn_ds_bpermute(index, src);
}
__device__ unsigned __hip_ds_permute(int index, unsigned src) {
return hc::__amdgcn_ds_permute(index, src);
}
__device__ float __hip_ds_permutef(int index, float src) {
return hc::__amdgcn_ds_permute(index, src);
}
__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
return hc::__amdgcn_ds_swizzle(src, pattern);
}
__device__ float __hip_ds_swizzlef(float src, int pattern) {
return hc::__amdgcn_ds_swizzle(src, pattern);
}
__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
}
//=================================================================================================
// Thread-local storage:
//=================================================================================================