补充SIMT shfl类接口 bf16/bf16x2 类型重载-asc-devkit-AtomGit

Merged

补充SIMT shfl类接口 bf16/bf16x2 类型重载 #3278

louyuxuanAdd on Jun 16Merged fromlouyuxuan/asc-devkit： master into cann/asc-devkit： 9.1.0

补充SIMT shfl类接口 bf16/bf16x2 类型重载 #3278

Merged

louyuxuanAdd on Jun 16

Merged fromlouyuxuan/asc-devkit masterintocann/asc-devkit 9.1.0

Total: 17 files changed+159-25

Mdocs/api/SIMT-API/Warp函数/Warp-Shfl类函数/asc_shfl.md+13-1

	@@ -64,6 +64,14 @@ inline half asc_shfl(half var, int32_t src_lane, int32_t width = warpSize)
64	inline half2 asc_shfl(half2 var, int32_t src_lane, int32_t width = warpSize)	64	inline half2 asc_shfl(half2 var, int32_t src_lane, int32_t width = warpSize)
65	```	65	```
66		66
		67	+```
		68	+inline bfloat16_t asc_shfl(bfloat16_t var, int32_t src_lane, int32_t width = warpSize)
		69	+```
		70	+
		71	+```
		72	+inline bfloat16x2_t asc_shfl(bfloat16x2_t var, int32_t src_lane, int32_t width = warpSize)
		73	+```
		74	+
67	## 参数说明	75	## 参数说明
68		76
69	表 1 参数说明	77	表 1 参数说明
	@@ -85,7 +93,7 @@ Warp内指定线程的var值。
85		93
86	## 需要包含的头文件	94	## 需要包含的头文件
87		95
88	-使用除half、half2类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件。	96	+使用除half、half2、bfloat16\_t、bfloat16x2\_t类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件，使用bfloat16\_t和bfloat16x2\_t类型接口需要包含"simt\_api/asc\_bf16.h"头文件。
89		97
90	```	98	```
91	#include "simt_api/device_warp_functions.h"	99	#include "simt_api/device_warp_functions.h"
	@@ -95,6 +103,10 @@ Warp内指定线程的var值。
95	#include "simt_api/asc_fp16.h"	103	#include "simt_api/asc_fp16.h"
96	```	104	```
97		105
		106	+```
		107	+#include "simt_api/asc_bf16.h"
		108	+```
		109	+
98	## 调用示例	110	## 调用示例
99		111
100	完整样例请参考[InsertHashTable算子样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/02_features/01_api_features/00_memory_access/insert_hash_table/README.md)。	112	完整样例请参考[InsertHashTable算子样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/02_features/01_api_features/00_memory_access/insert_hash_table/README.md)。

Mdocs/api/SIMT-API/Warp函数/Warp-Shfl类函数/asc_shfl_down.md+13-1

	@@ -54,6 +54,14 @@ inline half asc_shfl_down(half var, uint32_t delta, int32_t width = warpSize)
54	inline half2 asc_shfl_down(half2 var, uint32_t delta, int32_t width = warpSize)	54	inline half2 asc_shfl_down(half2 var, uint32_t delta, int32_t width = warpSize)
55	```	55	```
56		56
		57	+```
		58	+inline bfloat16_t asc_shfl_down(bfloat16_t var, uint32_t delta, int32_t width = warpSize)
		59	+```
		60	+
		61	+```
		62	+inline bfloat16x2_t asc_shfl_down(bfloat16x2_t var, uint32_t delta, int32_t width = warpSize)
		63	+```
		64	+
57	## 参数说明	65	## 参数说明
58		66
59	表 1 参数说明	67	表 1 参数说明
	@@ -76,7 +84,7 @@ Warp内指定线程的var值。
76		84
77	## 需要包含的头文件	85	## 需要包含的头文件
78		86
79	-使用除half、half2类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件。	87	+使用除half、half2、bfloat16\_t、bfloat16x2\_t类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件，使用bfloat16\_t和bfloat16x2\_t类型接口需要包含"simt\_api/asc\_bf16.h"头文件。
80		88
81	```	89	```
82	#include "simt_api/device_warp_functions.h"	90	#include "simt_api/device_warp_functions.h"
	@@ -86,6 +94,10 @@ Warp内指定线程的var值。
86	#include "simt_api/asc_fp16.h"	94	#include "simt_api/asc_fp16.h"
87	```	95	```
88		96
		97	+```
		98	+#include "simt_api/asc_bf16.h"
		99	+```
		100	+
89	## 调用示例	101	## 调用示例
90		102
91	完整样例请参考[Sobel边缘检测样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/02_features/01_api_features/03_warp_instruction/sobel_warp_shfl/README.md)。	103	完整样例请参考[Sobel边缘检测样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/02_features/01_api_features/03_warp_instruction/sobel_warp_shfl/README.md)。

Mdocs/api/SIMT-API/Warp函数/Warp-Shfl类函数/asc_shfl_up.md+13-1

	@@ -54,6 +54,14 @@ inline half asc_shfl_up(half var, uint32_t delta, int32_t width = warpSize)
54	inline half2 asc_shfl_up(half2 var, uint32_t delta, int32_t width = warpSize)	54	inline half2 asc_shfl_up(half2 var, uint32_t delta, int32_t width = warpSize)
55	```	55	```
56		56
		57	+```
		58	+inline bfloat16_t asc_shfl_up(bfloat16_t var, uint32_t delta, int32_t width = warpSize)
		59	+```
		60	+
		61	+```
		62	+inline bfloat16x2_t asc_shfl_up(bfloat16x2_t var, uint32_t delta, int32_t width = warpSize)
		63	+```
		64	+
57	## 参数说明	65	## 参数说明
58		66
59	表 1 参数说明	67	表 1 参数说明
	@@ -76,7 +84,7 @@ Warp内指定线程的var值。
76		84
77	## 需要包含的头文件	85	## 需要包含的头文件
78		86
79	-使用除half、half2类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件。	87	+使用除half、half2、bfloat16\_t、bfloat16x2\_t类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件，使用bfloat16\_t和bfloat16x2\_t类型接口需要包含"simt\_api/asc\_bf16.h"头文件。
80		88
81	```	89	```
82	#include "simt_api/device_warp_functions.h"	90	#include "simt_api/device_warp_functions.h"
	@@ -86,6 +94,10 @@ Warp内指定线程的var值。
86	#include "simt_api/asc_fp16.h"	94	#include "simt_api/asc_fp16.h"
87	```	95	```
88		96
		97	+```
		98	+#include "simt_api/asc_bf16.h"
		99	+```
		100	+
89	## 调用示例	101	## 调用示例
90		102
91	完整样例请参考[Sobel边缘检测样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/02_features/01_api_features/03_warp_instruction/sobel_warp_shfl/README.md)。	103	完整样例请参考[Sobel边缘检测样例](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/02_features/01_api_features/03_warp_instruction/sobel_warp_shfl/README.md)。

Mdocs/api/SIMT-API/Warp函数/Warp-Shfl类函数/asc_shfl_xor.md+13-1

	@@ -54,6 +54,14 @@ inline half asc_shfl_xor(half var, int32_t lane_mask, int32_t width = warpSize)
54	inline half2 asc_shfl_xor(half2 var, int32_t lane_mask, int32_t width = warpSize)	54	inline half2 asc_shfl_xor(half2 var, int32_t lane_mask, int32_t width = warpSize)
55	```	55	```
56		56
		57	+```
		58	+inline bfloat16_t asc_shfl_xor(bfloat16_t var, int32_t lane_mask, int32_t width = warpSize)
		59	+```
		60	+
		61	+```
		62	+inline bfloat16x2_t asc_shfl_xor(bfloat16x2_t var, int32_t lane_mask, int32_t width = warpSize)
		63	+```
		64	+
57	## 参数说明	65	## 参数说明
58		66
59	表 1 参数说明	67	表 1 参数说明
	@@ -74,7 +82,7 @@ Warp内指定线程的var值。
74		82
75	## 需要包含的头文件	83	## 需要包含的头文件
76		84
77	-使用除half、half2类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件。	85	+使用除half、half2、bfloat16\_t、bfloat16x2\_t类型之外的接口需要包含"simt\_api/device\_warp\_functions.h"头文件，使用half和half2类型接口需要包含"simt\_api/asc\_fp16.h"头文件，使用bfloat16\_t和bfloat16x2\_t类型接口需要包含"simt\_api/asc\_bf16.h"头文件。
78		86
79	```	87	```
80	#include "simt_api/device_warp_functions.h"	88	#include "simt_api/device_warp_functions.h"
	@@ -84,6 +92,10 @@ Warp内指定线程的var值。
84	#include "simt_api/asc_fp16.h"	92	#include "simt_api/asc_fp16.h"
85	```	93	```
86		94
		95	+```
		96	+#include "simt_api/asc_bf16.h"
		97	+```
		98	+
87	## 调用示例	99	## 调用示例
88		100
89	- SIMT编程场景：	101	- SIMT编程场景：

Mdocs/api/SIMT-API/协作组/coalesced_group/shfl.md+1-1

	@@ -29,7 +29,7 @@ T shfl(T var, int src_rank) const
29		29
30	\| 参数名 \| 输入/输出 \| 描述 \|	30	\| 参数名 \| 输入/输出 \| 描述 \|
31	\| --- \| --- \| --- \|	31	\| --- \| --- \| --- \|
32	-\| var \| 输入 \| 线程用于交换的数据。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	32	+\| var \| 输入 \| 线程用于交换的数据。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
33	\| src_rank \| 输入 \| 期望获取的`var`值所在的线程在组内的排名。当`src_rank`大于等于组内线程数时，获取`src_rank % num_threads()`对应rank线程的`var`值。 \|	33	\| src_rank \| 输入 \| 期望获取的`var`值所在的线程在组内的排名。当`src_rank`大于等于组内线程数时，获取`src_rank % num_threads()`对应rank线程的`var`值。 \|
34		34
35	## 返回值说明	35	## 返回值说明

Mdocs/api/SIMT-API/协作组/coalesced_group/shfl_down.md+1-1

	@@ -29,7 +29,7 @@ T shfl_down(T var, unsigned int delta) const
29		29
30	\| 参数名 \| 输入/输出 \| 描述 \|	30	\| 参数名 \| 输入/输出 \| 描述 \|
31	\| --- \| --- \| --- \|	31	\| --- \| --- \| --- \|
32	-\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	32	+\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向后偏移值。 \|	33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向后偏移值。 \|
34		34
35	## 返回值说明	35	## 返回值说明

Mdocs/api/SIMT-API/协作组/coalesced_group/shfl_up.md+1-1

	@@ -29,7 +29,7 @@ T shfl_up(T var, unsigned int delta) const
29		29
30	\| 参数名 \| 输入/输出 \| 描述 \|	30	\| 参数名 \| 输入/输出 \| 描述 \|
31	\| --- \| --- \| --- \|	31	\| --- \| --- \| --- \|
32	-\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	32	+\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向前偏移值。 \|	33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向前偏移值。 \|
34		34
35	## 返回值说明	35	## 返回值说明

Mdocs/api/SIMT-API/协作组/thread_block_tile/shfl.md+1-1

	@@ -30,7 +30,7 @@ T shfl(T var, int src_rank) const
30		30
31	\| 参数名 \| 输入/输出 \| 描述 \|	31	\| 参数名 \| 输入/输出 \| 描述 \|
32	\| --- \| --- \| --- \|	32	\| --- \| --- \| --- \|
33	-\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	33	+\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
34	\| src_rank \| 输入 \| 期望获取的`var`值所在的线程在组内的排名。当`src_rank`大于等于组内线程数时，获取`src_rank % Size`对应rank线程的`var`值。 \|	34	\| src_rank \| 输入 \| 期望获取的`var`值所在的线程在组内的排名。当`src_rank`大于等于组内线程数时，获取`src_rank % Size`对应rank线程的`var`值。 \|
35		35
36	## 返回值说明	36	## 返回值说明

Mdocs/api/SIMT-API/协作组/thread_block_tile/shfl_down.md+1-1

	@@ -29,7 +29,7 @@ T shfl_down(T var, unsigned int delta) const
29		29
30	\| 参数名 \| 输入/输出 \| 描述 \|	30	\| 参数名 \| 输入/输出 \| 描述 \|
31	\| --- \| --- \| --- \|	31	\| --- \| --- \| --- \|
32	-\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	32	+\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向后偏移值。 \|	33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向后偏移值。 \|
34		34
35	## 返回值说明	35	## 返回值说明

Mdocs/api/SIMT-API/协作组/thread_block_tile/shfl_up.md+1-1

	@@ -29,7 +29,7 @@ T shfl_up(T var, unsigned int delta) const
29		29
30	\| 参数名 \| 输入/输出 \| 描述 \|	30	\| 参数名 \| 输入/输出 \| 描述 \|
31	\| --- \| --- \| --- \|	31	\| --- \| --- \| --- \|
32	-\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	32	+\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向前偏移值。 \|	33	\| delta \| 输入 \| 期望获取的`var`值所在线程在组内相对当前线程的向前偏移值。 \|
34		34
35	## 返回值说明	35	## 返回值说明

Mdocs/api/SIMT-API/协作组/thread_block_tile/shfl_xor.md+1-1

	@@ -29,7 +29,7 @@ T shfl_xor(T var, unsigned int lane_mask) const
29		29
30	\| 参数名 \| 输入/输出 \| 描述 \|	30	\| 参数名 \| 输入/输出 \| 描述 \|
31	\| --- \| --- \| --- \|	31	\| --- \| --- \| --- \|
32	-\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t。 \|	32	+\| var \| 输入 \| 线程用于交换的输入操作数。支持的数据类型为：half、int32_t、uint32_t、float、half2、int64_t、uint64_t、bfloat16_t、bfloat16x2_t。 \|
33	\| lane_mask \| 输入 \| 与当前线程rank做异或运算的操作数。 \|	33	\| lane_mask \| 输入 \| 与当前线程rank做异或运算的操作数。 \|
34		34
35	## 返回值说明	35	## 返回值说明

Mimpl/simt_api/asc_bf16_impl.h+48-0

	@@ -1065,6 +1065,54 @@ __SIMT_DEVICE_FUNCTIONS_DECL__ inline void asc_stwt(bfloat16x2_t* address, bfloa
1065	}	1065	}
1066	#endif	1066	#endif
1067		1067
		1068	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl(bfloat16_t var, int32_t src_lane, int32_t width)
		1069	+{
		1070	+ return __shfl(var,
		1071	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (src_lane));
		1072	+}
		1073	+
		1074	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl(bfloat16x2_t var, int32_t src_lane, int32_t width)
		1075	+{
		1076	+ return __shfl(var,
		1077	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (src_lane));
		1078	+}
		1079	+
		1080	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl_up(bfloat16_t var, uint32_t delta, int32_t width)
		1081	+{
		1082	+ return __shfl_up(var,
		1083	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_UP_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (delta));
		1084	+}
		1085	+
		1086	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl_up(bfloat16x2_t var, uint32_t delta, int32_t width)
		1087	+{
		1088	+ return __shfl_up(var,
		1089	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_UP_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (delta));
		1090	+}
		1091	+
		1092	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl_down(bfloat16_t var, uint32_t delta, int32_t width)
		1093	+{
		1094	+ return __shfl_down(var,
		1095	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (delta));
		1096	+}
		1097	+
		1098	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl_down(bfloat16x2_t var, uint32_t delta, int32_t width)
		1099	+{
		1100	+ return __shfl_down(var,
		1101	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (delta));
		1102	+}
		1103	+
		1104	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl_xor(bfloat16_t var, int32_t lane_mask, int32_t width)
		1105	+{
		1106	+ return __shfl_xor(var,
		1107	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (lane_mask));
		1108	+}
		1109	+
		1110	+__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl_xor(bfloat16x2_t var, int32_t lane_mask, int32_t width)
		1111	+{
		1112	+ return __shfl_xor(var,
		1113	+ ((warpSize - width) << __INTERNAL_LANE_MASK_START_POS) \| (__INTERNAL_MAX_OFFSET_OF_MODE << __INTERNAL_MAX_OFFSET_START_POS) \| (lane_mask));
		1114	+}
		1115	+
1068	__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t make_bfloat162(bfloat16_t x, bfloat16_t y)	1116	__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t make_bfloat162(bfloat16_t x, bfloat16_t y)
1069	{	1117	{
1070	bfloat16x2_t tmp;	1118	bfloat16x2_t tmp;

Mimpl/simt_api/cooperative_groups_impl.h+14-14

Minclude/simt_api/asc_bf16.h+16-0

 __SIMT_DEVICE_FUNCTIONS_DECL__ inline void asc_stwt(bfloat16x2_t* address, bfloat16x2_t val);
 #endif
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl(bfloat16_t var, int32_t src_lane, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl(bfloat16x2_t var, int32_t src_lane, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl_up(bfloat16_t var, uint32_t delta, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl_up(bfloat16x2_t var, uint32_t delta, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl_down(bfloat16_t var, uint32_t delta, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl_down(bfloat16x2_t var, uint32_t delta, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t asc_shfl_xor(bfloat16_t var, int32_t lane_mask, int32_t width = warpSize);
++
++__SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t asc_shfl_xor(bfloat16x2_t var, int32_t lane_mask, int32_t width = warpSize);
++
 __SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16x2_t make_bfloat162(bfloat16_t x, bfloat16_t y);
 __SIMT_DEVICE_FUNCTIONS_DECL__ inline bfloat16_t __ushort_as_bfloat16(const unsigned short int x);

Minclude/simt_api/cooperative_groups.h+1-0

	@@ -19,6 +19,7 @@
19	#include "simt_api/common_functions.h"	19	#include "simt_api/common_functions.h"
20	#include "simt_api/device_types.h"	20	#include "simt_api/device_types.h"
21	#include "simt_api/asc_fp16.h"	21	#include "simt_api/asc_fp16.h"
		22	+#include "simt_api/asc_bf16.h"
22		23
23	namespace cooperative_groups {	24	namespace cooperative_groups {
24		25

Mtests/api/simt_api/common_simt/device_warp_functions_debug.cpp+8-0

	@@ -93,6 +93,8 @@ REGISTER_SIMT_ASC_SHFL(__shfl, int64_t);
93	REGISTER_SIMT_ASC_SHFL(__shfl, float);	93	REGISTER_SIMT_ASC_SHFL(__shfl, float);
94	REGISTER_SIMT_ASC_SHFL(__shfl, half);	94	REGISTER_SIMT_ASC_SHFL(__shfl, half);
95	REGISTER_SIMT_ASC_SHFL(__shfl, half2);	95	REGISTER_SIMT_ASC_SHFL(__shfl, half2);
		96	+REGISTER_SIMT_ASC_SHFL(__shfl, bfloat16_t);
		97	+REGISTER_SIMT_ASC_SHFL(__shfl, bfloat16x2_t);
96	#endif	98	#endif
97		99
98	template <typename T>	100	template <typename T>
	@@ -127,6 +129,8 @@ REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, int64_t);
127	REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, float);	129	REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, float);
128	REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, half);	130	REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, half);
129	REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, half2);	131	REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, half2);
		132	+REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, bfloat16_t);
		133	+REGISTER_SIMT_ASC_SHFL_UP(__shfl_up, bfloat16x2_t);
130	#endif	134	#endif
131		135
132	template <typename T>	136	template <typename T>
	@@ -163,6 +167,8 @@ REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, int64_t);
163	REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, float);	167	REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, float);
164	REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, half);	168	REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, half);
165	REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, half2);	169	REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, half2);
		170	+REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, bfloat16_t);
		171	+REGISTER_SIMT_ASC_SHFL_DOWN(__shfl_down, bfloat16x2_t);
166	#endif	172	#endif
167		173
168	template <typename T>	174	template <typename T>
	@@ -199,6 +205,8 @@ REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, int64_t);
199	REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, float);	205	REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, float);
200	REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, half);	206	REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, half);
201	REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, half2);	207	REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, half2);
		208	+REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, bfloat16_t);
		209	+REGISTER_SIMT_ASC_SHFL_XOR(__shfl_xor, bfloat16x2_t);
202	#endif	210	#endif
203		211
204	template <typename T>	212	template <typename T>

Mtests/api/simt_api/common_simt/simt_stub.h+13-0

	@@ -48,6 +48,19 @@ T __shfl(T var, U mask)
48	return var;	48	return var;
49	}	49	}
50		50
		51	+// bfloat16_t / bfloat16x2_t warp-shuffle stubs are defined in
		52	+// device_warp_functions_debug.cpp (the released stub_fun.h only declares up to
		53	+// half/half2). Declare them here so the asc_bf16 impl resolves them in the
		54	+// CPU-debug twin build.
		55	+bfloat16_t __shfl(bfloat16_t var, int32_t num);
		56	+bfloat16x2_t __shfl(bfloat16x2_t var, int32_t num);
		57	+bfloat16_t __shfl_up(bfloat16_t var, int32_t num);
		58	+bfloat16x2_t __shfl_up(bfloat16x2_t var, int32_t num);
		59	+bfloat16_t __shfl_down(bfloat16_t var, int32_t num);
		60	+bfloat16x2_t __shfl_down(bfloat16x2_t var, int32_t num);
		61	+bfloat16_t __shfl_xor(bfloat16_t var, int32_t num);
		62	+bfloat16x2_t __shfl_xor(bfloat16x2_t var, int32_t num);
		63	+
51	static inline unsigned int __brev(unsigned int x)	64	static inline unsigned int __brev(unsigned int x)
52	{	65	{
53	unsigned int result = 0;	66	unsigned int result = 0;

	@@ -314,8 +314,8 @@ template <typename T>
314	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T coalesced_group::shfl(T var, int src_rank) const	314	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T coalesced_group::shfl(T var, int src_rank) const
315	{	315	{
316	static_assert(	316	static_assert(
317	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	317	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
318	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	318	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
319	int lane = src_rank % static_cast<int>(num_threads());	319	int lane = src_rank % static_cast<int>(num_threads());
320	if (num_threads() != warpSize) {	320	if (num_threads() != warpSize) {
321	lane = __fns_internal(_tiled_info.mask, 0, lane + 1);	321	lane = __fns_internal(_tiled_info.mask, 0, lane + 1);
	@@ -327,8 +327,8 @@ template <typename T>
327	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T coalesced_group::shfl_up(T var, unsigned int delta) const	327	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T coalesced_group::shfl_up(T var, unsigned int delta) const
328	{	328	{
329	static_assert(	329	static_assert(
330	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	330	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
331	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	331	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
332	if (num_threads() == warpSize) {	332	if (num_threads() == warpSize) {
333	return asc_shfl_up(var, delta, warpSize);	333	return asc_shfl_up(var, delta, warpSize);
334	}	334	}
	@@ -345,8 +345,8 @@ template <typename T>
345	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T coalesced_group::shfl_down(T var, unsigned int delta) const	345	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T coalesced_group::shfl_down(T var, unsigned int delta) const
346	{	346	{
347	static_assert(	347	static_assert(
348	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	348	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
349	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	349	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
350	if (num_threads() == warpSize) {	350	if (num_threads() == warpSize) {
351	return asc_shfl_down(var, delta, warpSize);	351	return asc_shfl_down(var, delta, warpSize);
352	}	352	}
	@@ -420,8 +420,8 @@ template <typename T>
420	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl(T var, int src_rank) const	420	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl(T var, int src_rank) const
421	{	421	{
422	static_assert(	422	static_assert(
423	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	423	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
424	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	424	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
425	return asc_shfl(var, src_rank, static_cast<int32_t>(numThreads));	425	return asc_shfl(var, src_rank, static_cast<int32_t>(numThreads));
426	}	426	}
427		427
	@@ -430,8 +430,8 @@ template <typename T>
430	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl_up(T var, unsigned int delta) const	430	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl_up(T var, unsigned int delta) const
431	{	431	{
432	static_assert(	432	static_assert(
433	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	433	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
434	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	434	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
435	return asc_shfl_up(var, delta, static_cast<int32_t>(numThreads));	435	return asc_shfl_up(var, delta, static_cast<int32_t>(numThreads));
436	}	436	}
437		437
	@@ -440,8 +440,8 @@ template <typename T>
440	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl_down(T var, unsigned int delta) const	440	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl_down(T var, unsigned int delta) const
441	{	441	{
442	static_assert(	442	static_assert(
443	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	443	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
444	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	444	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
445	return asc_shfl_down(var, delta, static_cast<int32_t>(numThreads));	445	return asc_shfl_down(var, delta, static_cast<int32_t>(numThreads));
446	}	446	}
447		447
	@@ -450,8 +450,8 @@ template <typename T>
450	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl_xor(T var, unsigned int lane_mask) const	450	__SIMT_DEVICE_FUNCTIONS_DECL__ inline T thread_block_tile_base<Size>::shfl_xor(T var, unsigned int lane_mask) const
451	{	451	{
452	static_assert(	452	static_assert(
453	- SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2>,	453	+ SupportTypeSimtInternel<T, int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t>,
454	- "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2.");	454	+ "Input type T only supports int32_t, uint32_t, int64_t, uint64_t, float, half, half2, bfloat16_t, bfloat16x2_t.");
455	return asc_shfl_xor(var, static_cast<int32_t>(lane_mask), static_cast<int32_t>(numThreads));	455	return asc_shfl_xor(var, static_cast<int32_t>(lane_mask), static_cast<int32_t>(numThreads));
456	}	456	}
457		457