thread_block构造函数
产品支持情况
- Ascend 950PR/Ascend 950DT:支持
- Atlas A3 训练系列产品/Atlas A3 推理系列产品:不支持
- Atlas A2 训练系列产品/Atlas A2 推理系列产品:不支持
- Atlas 200I/500 A2 推理产品:不支持
- Atlas 推理系列产品AI Core:不支持
- Atlas 推理系列产品Vector Core:不支持
- Atlas 训练系列产品:不支持
功能说明
thread_block不提供默认的构造函数,用户使用this_thread_block函数获取当前线程所属的thread_block对象。
函数原型
C++
// 默认场景使用此接口创建thread_block
thread_block this_thread_block()
C++
// 后续需要创建跨warp的thread_block_tile(size > 32)时使用此接口创建thread_block
template <unsigned int MaxBlockSize>
thread_block this_thread_block(block_tile_memory<MaxBlockSize>& scratch)
参数说明
接口参数说明
表1 参数说明
| 参数名 | 输入/输出 | 描述 |
|---|---|---|
| scratch | 输入 | Global Memory或Unified Buffer上的block_tile_memory对象,用于为thread_block携带跨Warp临时存储。 |
block_tile_memory说明
block_tile_memory用于为跨Warp的thread_block_tile提供组内同步和跨Warp通信所需的临时存储。结构体大小可通过sizeof()接口获取。
C++
template <unsigned int MaxBlockSize = 1024>
struct block_tile_memory
表2 block_tile_memory模板参数说明
| 参数名 | 输入/输出 | 描述 |
|---|---|---|
| MaxBlockSize | 输入 | 模板参数,指定当前线程块中需要支持的最大线程数,必须大于0且小于等于2048,必须是Warp大小32的整数倍,默认值为1024。 |
注意
线程块实际启动线程数不能超过MaxBlockSize,否则是未定义行为。
返回值说明
返回当前线程所属的thread_block对象。
约束说明
- 线程块实际启动线程数不能超过
MaxBlockSize。 scratch对应的block_tile_memory对象必须位于Global Memory或Unified Buffer,不能是在栈空间中创建的对象。
调用示例
参考调用示例