Searched refs:GetThreadIdInBlock (Results 1 – 11 of 11) sorted by relevance
| /external/llvm-project/openmp/libomptarget/deviceRTLs/common/src/ |
| D | data_sharing.cu | 17 return !isSPMDExecutionMode && GetMasterThreadID() == GetThreadIdInBlock(); in IsMasterThread()
57 if (GetThreadIdInBlock() == 0) in __kmpc_data_sharing_init_stack_spmd()
67 bool IsWarpMaster = (GetThreadIdInBlock() % WARPSIZE) == 0; in data_sharing_push_stack_common()
179 if (GetThreadIdInBlock() % WARPSIZE == 0) { in __kmpc_data_sharing_pop_stack()
249 if (GetThreadIdInBlock() == 0) { in __kmpc_get_team_static_memory()
255 ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(), in __kmpc_get_team_static_memory()
267 if (GetThreadIdInBlock() == 0) { in __kmpc_restore_team_static_memory()
273 ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(), in __kmpc_restore_team_static_memory()
|
| D | omptarget.cu | 37 int threadIdInBlock = GetThreadIdInBlock(); in __kmpc_kernel_init()
86 int threadId = GetThreadIdInBlock(); in __kmpc_spmd_kernel_init()
146 int threadId = GetThreadIdInBlock(); in __kmpc_spmd_kernel_deinit_v2()
|
| D | reduction.cu | 61 physical_lane_id = GetThreadIdInBlock() % WARPSIZE; in gpu_irregular_simd_reduce()
111 /*LaneId=*/GetThreadIdInBlock() % WARPSIZE); in nvptx_parallel_reduce_nowait()
134 /*LaneId=*/GetThreadIdInBlock() % WARPSIZE); in nvptx_parallel_reduce_nowait()
196 if (checkGenericMode(loc) && GetThreadIdInBlock() != GetMasterThreadID()) in __kmpc_nvptx_teams_reduce_nowait_v2()
|
| D | parallel.cu | 107 ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(), in __kmpc_kernel_prepare_parallel()
135 int threadId = GetThreadIdInBlock(); in __kmpc_kernel_parallel()
178 int threadId = GetThreadIdInBlock(); in __kmpc_kernel_end_parallel()
|
| D | support.cu | 125 int tid = GetThreadIdInBlock(); in GetLogicalThreadIdInBlock()
144 rc = GetThreadIdInBlock(); in GetOmpThreadId()
|
| D | libcall.cu | 147 return level == 1 ? GetThreadIdInBlock() : 0; in omp_get_ancestor_thread_num()
|
| /external/llvm-project/openmp/libomptarget/deviceRTLs/nvptx/src/ |
| D | target_impl.h | 206 INLINE int GetThreadIdInBlock() { return threadIdx.x; } in GetThreadIdInBlock() function
210 INLINE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; } in GetWarpId()
211 INLINE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); } in GetLaneId()
|
| /external/llvm-project/openmp/libomptarget/deviceRTLs/common/ |
| D | debug.h | 134 printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(), in log()
145 printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(), in check()
|
| D | omptargeti.h | 61 GetThreadIdInBlock(); // get ids from cuda (only called for 1st level) in InitLevelOneTaskDescr()
113 GetThreadIdInBlock(); // get ids from cuda (only called for 1st level) in CopyFromWorkDescr()
|
| /external/llvm-project/openmp/libomptarget/deviceRTLs/amdgcn/src/ |
| D | target_impl.h | 131 INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); } in GetThreadIdInBlock() function
|
| D | target_impl.hip | 142 DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
|