Searched refs:WARPSIZE (Results 1 – 14 of 14) sorted by relevance
37 for (uint32_t mask = WARPSIZE / 2; mask > 0; mask /= 2) { in gpu_regular_warp_reduce()61 physical_lane_id = GetThreadIdInBlock() % WARPSIZE; in gpu_irregular_simd_reduce()99 uint32_t WarpsNeeded = (NumThreads + WARPSIZE - 1) / WARPSIZE; in nvptx_parallel_reduce_nowait()100 uint32_t WarpId = BlockThreadId / WARPSIZE; in nvptx_parallel_reduce_nowait()106 if ((NumThreads % WARPSIZE == 0) || (WarpId < WarpsNeeded - 1)) in nvptx_parallel_reduce_nowait()110 /*LaneCount=*/NumThreads % WARPSIZE, in nvptx_parallel_reduce_nowait()111 /*LaneId=*/GetThreadIdInBlock() % WARPSIZE); in nvptx_parallel_reduce_nowait()117 if (NumThreads > WARPSIZE) { in nvptx_parallel_reduce_nowait()134 /*LaneId=*/GetThreadIdInBlock() % WARPSIZE); in nvptx_parallel_reduce_nowait()144 if (NumThreads > WARPSIZE) { in nvptx_parallel_reduce_nowait()[all …]
58 WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE); in __kmpc_barrier()
67 bool IsWarpMaster = (GetThreadIdInBlock() % WARPSIZE) == 0; in data_sharing_push_stack_common()160 : WARPSIZE * DataSize; in __kmpc_data_sharing_push_stack()179 if (GetThreadIdInBlock() % WARPSIZE == 0) { in __kmpc_data_sharing_pop_stack()
64 if (NumThreads < WARPSIZE) { in determineNumberOfThreads()67 NumThreads = (NumThreads & ~((uint16_t)WARPSIZE - 1)); in determineNumberOfThreads()
35 DEVICE SHARED uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
34 for (int I = 0; I < MAX_THREADS_PER_TEAM / WARPSIZE; ++I) in __kmpc_kernel_init()
109 DEVICE int GetMasterThreadID() { return (GetNumberOfThreadsInBlock() - 1) & ~(WARPSIZE - 1); } in GetMasterThreadID()
38 #define WARPSIZE 32 macro77 DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,210 INLINE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; } in GetWarpId()211 INLINE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); } in GetLaneId()
30 if (lane == (WARPSIZE - 1))54 int width = WARPSIZE;78 uint32_t num_waves = num_threads / WARPSIZE;88 assert(num_waves * WARPSIZE == num_threads);142 DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
43 #define WARPSIZE 64 macro63 DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
103 for (int offset = WARPSIZE/2; offset > 0; offset /= 2)253 int wid = threadIdx.x/WARPSIZE;254 int lane_id = threadIdx.x%WARPSIZE;256 int warp_needed = (size+WARPSIZE-1)/WARPSIZE; //ceiling of division262 if (thread_num == WARPSIZE) {266 if (thread_num < WARPSIZE) {
299 parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
81 ; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.warpsize(), !range ![[WARPSIZE:[0-9]+]]319 ; RANGE-DAG: ![[WARPSIZE]] = !{i32 32, i32 33}
81 ; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.warpsize(), !range ![[WARPSIZE:[0-9]+]]329 ; RANGE-DAG: ![[WARPSIZE]] = !{i32 32, i32 33}