asc_reduce_min
产品支持情况
产品 |
是否支持 |
|---|---|
Atlas 350 加速卡 |
√ |
x |
|
x |
|
x |
|
x |
|
x |
|
x |
功能说明
对Warp内所有活跃线程输入val求最小值。Warp内所有活跃线程返回相同的结果。
函数原型
1 | inline int32_t asc_reduce_min(int32_t val) |
1 | inline uint32_t asc_reduce_min(uint32_t val) |
1 | inline float asc_reduce_min(float val) |
1 | inline half asc_reduce_min(half val) |
参数说明
参数名 |
输入/输出 |
描述 |
|---|---|---|
val |
输入 |
源操作数。 |
返回值说明
Warp内所有活跃线程输入val的最小值。
约束说明
SIMT编程场景当前不支持使用该接口。
需要包含的头文件
使用除half类型之外的接口需要包含"simt_api/device_warp_functions.h"头文件，使用half类型接口需要包含"simt_api/asc_fp16.h"头文件。
1 | #include "simt_api/device_warp_functions.h" |
1 | #include "simt_api/asc_fp16.h" |
调用示例
SIMD与SIMT混合编程场景：
1 2 3 4 5 6 7 8 | __simt_vf__ __launch_bounds__(1024) inline void KernelReduceMin(__gm__ int32_t* dst) { // asc_vf_call参数：dim3{1024, 1, 1} int idx = threadIdx.x + blockIdx.x * blockDim.x; int32_t laneId = idx % 32; int32_t result = asc_reduce_min(laneId); // 返回值为0 dst[idx] = result; } |
父主题： Warp函数