Home
last modified time | relevance | path

Searched refs:logical_lane_id (Results 1 – 2 of 2) sorted by relevance

/external/llvm-project/openmp/libomptarget/deviceRTLs/common/src/
Dreduction.cu64 uint32_t logical_lane_id = __kmpc_impl_popc(Liveness & lanemask_lt) * 2; in gpu_irregular_simd_reduce() local
70 logical_lane_id /= 2; in gpu_irregular_simd_reduce()
71 shflFct(reduce_data, /*LaneId =*/logical_lane_id, in gpu_irregular_simd_reduce()
73 } while (logical_lane_id % 2 == 0 && size > 1); in gpu_irregular_simd_reduce()
74 return (logical_lane_id == 0); in gpu_irregular_simd_reduce()
/external/llvm-project/openmp/libomptarget/deviceRTLs/nvptx/docs/
DReductionDesign.txt183 int logical_lane_id = find_number_of_dispersed_active_lanes_before_me() * 2;
189 logical_lane_id /= 2;
190 ShuffleReduceFn(reduce_data, logical_lane_id, remote_id-1-threadIdx.x, 2);
191 } while (logical_lane_id % 2 == 0 && size > 1);
212 concept of logical_lane_id, and it is important to distinguish it
214 1. //logical_lane_id//: as the name suggests, it refers to the calculated
216 our algorithm logically concise. A thread with logical_lane_id k means
233 f(x): logical_lane_id -> physical_lane_id;
235 f^-1(x): physical_lane_id -> logical_lane_id
243 logical_lane_id is recalculated; so is the remote_id.