fix: 修正矩阵转置样例类型及文档-asc-devkit-AtomGit

已合并

fix: 修正矩阵转置样例类型及文档 #5067

liuyi424创建于 25 天前从liuyi424/asc-devkit： codex/fix-simd-simt-matrix-transpose 合入到 cann/asc-devkit： master

fix: 修正矩阵转置样例类型及文档 #5067

已合并

liuyi424创建于 25 天前

从liuyi424/asc-devkit codex/fix-simd-simt-matrix-transpose合入到cann/asc-devkit master

共 3 个文件变更+74-49

Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_matrix_transpose/README.md+3-4

	@@ -233,8 +233,8 @@ SIMD与SIMT混合编程场景下，UB数据由MTE搬入。MTE搬运带padding的
233	<img src="./figures/case3_bank.png">	233	<img src="./figures/case3_bank.png">
234		234
235	```cpp	235	```cpp
236	-constexpr int TILE_PAD = 8; // MTE搬运要求32B对齐，使用32×40的UB布局	236	+constexpr uint32_t TILE_PAD = 8; // MTE搬运要求32B对齐，使用32×40的UB布局
237	-constexpr int TILE_PAD_STRIDE = TILE_DIM + TILE_PAD;	237	+constexpr uint32_t TILE_PAD_STRIDE = TILE_DIM + TILE_PAD;
238	__ubuf__ float in_tile[TILES_PER_BLOCK][TILE_DIM][TILE_PAD_STRIDE];	238	__ubuf__ float in_tile[TILES_PER_BLOCK][TILE_DIM][TILE_PAD_STRIDE];
239	__ubuf__ float out_tile[TILES_PER_BLOCK][TILE_DIM][TILE_DIM];	239	__ubuf__ float out_tile[TILES_PER_BLOCK][TILE_DIM][TILE_DIM];
240	uint32_t loop_start = block_idx * TILES_PER_BLOCK;	240	uint32_t loop_start = block_idx * TILES_PER_BLOCK;
	@@ -405,7 +405,6 @@ Case 0距离理论下限较远，主要原因是转置写回导致同一Warp内
405		405
406	执行结果如下，说明精度对比成功。	406	执行结果如下，说明精度对比成功。
407		407
408	-
409	```text	408	```text
410	[Success] Case accuracy verification passed.	409	[Success] Case accuracy verification passed.
411	```	410	```
	@@ -422,7 +421,7 @@ msOpProf工具是单算子性能分析工具。包含msopprof和msopprof simulat
422		421
423	基于可执行文件matrix_transpose通过msopprof执行算子调优：	422	基于可执行文件matrix_transpose通过msopprof执行算子调优：
424		423
425	- ```	424	+ ```bash
426	msopprof ./matrix_transpose	425	msopprof ./matrix_transpose
427	```	426	```
428		427

Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_matrix_transpose/README_en.md+64-37

	@@ -60,7 +60,7 @@ This sample constructs an optimization path through five cases. The kernel, Thre
60		60
61	Here, `core` is the hardware vector core count. The runtime queries this value by using `aclrtGetDeviceInfo(ACL_DEV_ATTR_VECTOR_CORE_NUM)`. The performance data in this sample was collected in a test environment where `core=64`. `tiles` is the total tile count `(W/32) x (H/32)`. When processing a 1024x1024 matrix, the matrix is divided into 32x32 tiles, so `tiles=1024`.	61	Here, `core` is the hardware vector core count. The runtime queries this value by using `aclrtGetDeviceInfo(ACL_DEV_ATTR_VECTOR_CORE_NUM)`. The performance data in this sample was collected in a test environment where `core=64`. `tiles` is the total tile count `(W/32) x (H/32)`. When processing a 1024x1024 matrix, the matrix is divided into 32x32 tiles, so `tiles=1024`.
62		62
63	-### Performance Metrics	63	+#### Performance Metrics
64		64
65	\| Metric \| Description \|	65	\| Metric \| Description \|
66	\| ------ \| ----------- \|	66	\| ------ \| ----------- \|
	@@ -215,9 +215,9 @@ When the SIMT VF accesses UB, bank conflicts are finer-grained subbank conflicts
215		215
216	For detailed address low-bit interleaving rules and conflict scenarios, see the [bank_conflict sample](../../../03_simt_api/02_features/01_api_features/00_memory_access/bank_conflict/README_en.md).	216	For detailed address low-bit interleaving rules and conflict scenarios, see the [bank_conflict sample](../../../03_simt_api/02_features/01_api_features/00_memory_access/bank_conflict/README_en.md).
217		217
218	-In Case 2, the UB `in_tile` array is stored in row-major order. According to the low-bit interleaving rule, the first row of `in_tile` covers bank0 through bank3, the second row covers bank4 through bank7, the third row covers bank8 through bank11, and the remaining rows follow the same pattern. Each row has 32 float elements and exactly spans 4 banks. As shown in Figure 3, the figure shows how the first 12 rows of the tile array are placed in UB. The first element of each row is marked in blue. During SIMT VF transpose, threads in one Warp read one column of `in_tile`. During the UB access, 32 threads are concentrated on subbank0 in two bank groups, which is a read-read conflict. `out_tile` is written continuously along the output tile row direction, so memory access merging can reduce write overhead. Therefore, `out_tile` does not need padding.	218	+In Case 2, the UB `in_tile` array is stored in row-major order. According to the low-bit interleaving rule, the first row of `in_tile` covers bank0 through bank3, the second row covers bank4 through bank7, the third row covers bank8 through bank11, and the remaining rows follow the same pattern. Each row has 32 float elements and exactly spans 4 banks. As shown in Figure 2, the figure shows how the first 12 rows of the tile array are placed in UB. The first element of each row is marked in blue. During SIMT VF transpose, threads in one Warp read one column of `in_tile`. During the UB access, 32 threads are concentrated on subbank0 in two bank groups, which is a read-read conflict. `out_tile` is written continuously along the output tile row direction, so memory access merging can reduce write overhead. Therefore, `out_tile` does not need padding.
219		219
220	-Figure 3: Case 2 tile layout in UB	220	+Figure 2: Case 2 tile layout in UB
221		221
222	<img src="./figures/case2_bank.png">	222	<img src="./figures/case2_bank.png">
223		223
	@@ -225,17 +225,17 @@ In SIMT scenarios, a common way to handle bank conflicts during transpose access
225		225
226	In SIMD and SIMT hybrid programming, MTE moves data into UB. When MTE moves a padded two-dimensional array, the address stride between adjacent rows in UB also needs to meet the alignment requirement. The row stride of a 32x34 layout is `34 * sizeof(float) = 136B`, which does not meet the 32B alignment requirement. Therefore, this sample sets the padding column count to 8 and uses a 32x40 layout. The row stride is `40 * sizeof(float) = 160B`, which staggers UB bank accesses and meets the row stride alignment requirement for MTE movement.	226	In SIMD and SIMT hybrid programming, MTE moves data into UB. When MTE moves a padded two-dimensional array, the address stride between adjacent rows in UB also needs to meet the alignment requirement. The row stride of a 32x34 layout is `34 * sizeof(float) = 136B`, which does not meet the 32B alignment requirement. Therefore, this sample sets the padding column count to 8 and uses a 32x40 layout. The row stride is `40 * sizeof(float) = 160B`, which staggers UB bank accesses and meets the row stride alignment requirement for MTE movement.
227		227
228	-As shown in Figure 4, Case 3 adds 8 columns of padding to the UB tile array, changing each row from 32 elements to 40 elements. The row stride of 40 float elements is 160B, which corresponds to 20 subbanks. When the same column is accessed in the transpose direction, elements in adjacent rows are staggered by a stride of 20 subbanks in the physical UB address space. The accesses of 32 threads are no longer concentrated on subbanks with the same subbank ID in the same bank group, which reduces subbank conflicts during the SIMT VF transpose access phase.	228	+As shown in Figure 3, Case 3 adds 8 columns of padding to the UB tile array, changing each row from 32 elements to 40 elements. The row stride of 40 float elements is 160B, which corresponds to 20 subbanks. When the same column is accessed in the transpose direction, elements in adjacent rows are staggered by a stride of 20 subbanks in the physical UB address space. The accesses of 32 threads are no longer concentrated on subbanks with the same subbank ID in the same bank group, which reduces subbank conflicts during the SIMT VF transpose access phase.
229		229
230	Note that the 32x40 layout is a tradeoff for SIMD and SIMT hybrid programming after considering the alignment requirements of MTE movement. Unlike the 32x34 layout formed by adding 2 columns of padding in a SIMT scenario, the 32x40 layout may still have a small number of subbank conflicts. However, its conflict intensity is much lower than that of the 32x32 layout without padding.	230	Note that the 32x40 layout is a tradeoff for SIMD and SIMT hybrid programming after considering the alignment requirements of MTE movement. Unlike the 32x34 layout formed by adding 2 columns of padding in a SIMT scenario, the 32x40 layout may still have a small number of subbank conflicts. However, its conflict intensity is much lower than that of the 32x32 layout without padding.
231		231
232	-Figure 4: Case 3 tile layout in UB	232	+Figure 3: Case 3 tile layout in UB
233		233
234	<img src="./figures/case3_bank.png">	234	<img src="./figures/case3_bank.png">
235		235
236	```cpp	236	```cpp
237	-constexpr int TILE_PAD = 8; // The hybrid MTE movement requires 32B alignment, so a 32x40 UB layout is used.	237	+constexpr uint32_t TILE_PAD = 8; // The hybrid MTE movement requires 32B alignment, so a 32x40 UB layout is used.
238	-constexpr int TILE_PAD_STRIDE = TILE_DIM + TILE_PAD;	238	+constexpr uint32_t TILE_PAD_STRIDE = TILE_DIM + TILE_PAD;
239	__ubuf__ float in_tile[TILES_PER_BLOCK][TILE_DIM][TILE_PAD_STRIDE];	239	__ubuf__ float in_tile[TILES_PER_BLOCK][TILE_DIM][TILE_PAD_STRIDE];
240	__ubuf__ float out_tile[TILES_PER_BLOCK][TILE_DIM][TILE_DIM];	240	__ubuf__ float out_tile[TILES_PER_BLOCK][TILE_DIM][TILE_DIM];
241	uint32_t loop_start = block_idx * TILES_PER_BLOCK;	241	uint32_t loop_start = block_idx * TILES_PER_BLOCK;
	@@ -297,7 +297,8 @@ for (uint32_t tile_base = curr_tile_base; tile_base < total_tiles; tile_base +=
297	copy_gm_2tile_to_padded_ub(&in_tile[curr_buffer][0][0][0], input, width, tiles_x, tile_base, total_tiles);	297	copy_gm_2tile_to_padded_ub(&in_tile[curr_buffer][0][0][0], input, width, tiles_x, tile_base, total_tiles);
298	asc_unlock(PIPE_MTE2, input_mutex);	298	asc_unlock(PIPE_MTE2, input_mutex);
299		299
300	- // The SIMT VF waits for the current input buffer. It also waits before reusing an output buffer.	300	+ // The SIMT VF waits for the current input buffer load to complete. Before reusing an output buffer,
		301	+ // it waits for the previous MTE3 store on that buffer to complete.
301	asc_lock(PIPE_V, input_mutex);	302	asc_lock(PIPE_V, input_mutex);
302	asc_lock(PIPE_V, output_mutex);	303	asc_lock(PIPE_V, output_mutex);
303	asc_vf_call<simt_transpose_2tile_pad>(	304	asc_vf_call<simt_transpose_2tile_pad>(
	@@ -323,20 +324,18 @@ for (uint32_t tile_base = curr_tile_base; tile_base < total_tiles; tile_base +=
323		324
324	Case 4 uses double buffering to provide physical isolation. This avoids MTE2 writes overwriting the input buffer that the SIMT VF is reading. It also avoids SIMT VF writes overwriting the output buffer that MTE3 is storing. Multiple mutex IDs manage the lifecycles of the two input/output buffer groups separately. The single-buffer serial barrier sequence of MTE2 load, SIMT VF transpose, and MTE3 store is changed into pipeline overlap between adjacent iterations. Compared with Case 3, Task Duration decreases from 11.596μs to 6.831μs, a latency reduction of about 41.1%. Compared with direct GM transpose in Case 0, Task Duration decreases from 36.263μs to 6.831μs, improving overall performance by about 5.31x.	325	Case 4 uses double buffering to provide physical isolation. This avoids MTE2 writes overwriting the input buffer that the SIMT VF is reading. It also avoids SIMT VF writes overwriting the output buffer that MTE3 is storing. Multiple mutex IDs manage the lifecycles of the two input/output buffer groups separately. The single-buffer serial barrier sequence of MTE2 load, SIMT VF transpose, and MTE3 store is changed into pipeline overlap between adjacent iterations. Compared with Case 3, Task Duration decreases from 11.596μs to 6.831μs, a latency reduction of about 41.1%. Compared with direct GM transpose in Case 0, Task Duration decreases from 36.263μs to 6.831μs, improving overall performance by about 5.31x.
325		326
326	-Figure 5 and Figure 6 show the simulated instruction timelines of Case 3 and Case 4. Figure 5 shows that Case 3 uses a single buffer to process each tile group serially. After the MTE2 load is complete, the SIMT VF must wait until the input buffer is readable. After the SIMT VF finishes, MTE3 can store the output buffer. Only after MTE3 completes can the same UB buffer be reused by the next MTE2 load. Therefore, the MTE2, SIMT VF, and MTE3 pipelines have clear serial waits. The gaps between SIMT VF executions mainly come from data movement and synchronization.	327	+Figure 4 and Figure 5 show the simulated instruction timelines of Case 3 and Case 4. Figure 4 shows that Case 3 uses a single buffer to process each tile group serially. After the MTE2 load is complete, the SIMT VF must wait until the input buffer is readable. After the SIMT VF finishes, MTE3 can store the output buffer. Only after MTE3 completes can the same UB buffer be reused by the next MTE2 load. Therefore, the MTE2, SIMT VF, and MTE3 pipelines have clear serial waits. The gaps between SIMT VF executions mainly come from data movement and synchronization.
327		328
328	-Figure 6 shows that Case 4 uses double buffering to overlap adjacent iterations. While the current buffer is processed by the SIMT VF and MTE3, the other buffer can start the next MTE2 load in advance. MTE2, SIMT VF, and MTE3 overlap across adjacent iterations. Therefore, Case 4 hides part of the MTE movement and synchronization overhead, further reducing Task Duration. The remaining waits mainly protect buffer reuse and cross-pipeline data dependencies.	329	+Figure 5 shows that Case 4 uses double buffering to overlap adjacent iterations. While the current buffer is processed by the SIMT VF and MTE3, the other buffer can start the next MTE2 load in advance. MTE2, SIMT VF, and MTE3 overlap across adjacent iterations. Therefore, Case 4 hides part of the MTE movement and synchronization overhead, further reducing Task Duration. The remaining waits mainly protect buffer reuse and cross-pipeline data dependencies.
329		330
330	-Figure 5: Simulated instruction timeline of Case 3	331	+Figure 4: Simulated instruction timeline of Case 3
331		332
332	<img src="./figures/case3_trace.png">	333	<img src="./figures/case3_trace.png">
333		334
334	-Figure 6: Simulated instruction timeline of Case 4	335	+Figure 5: Simulated instruction timeline of Case 4
335		336
336	<img src="./figures/case4_trace.png">	337	<img src="./figures/case4_trace.png">
337		338
338	----
339	-
340	## Performance Comparison Summary	339	## Performance Comparison Summary
341		340
342	### Ascend 950PR Performance Data	341	### Ascend 950PR Performance Data
	@@ -421,33 +420,61 @@ In the sample root directory, perform the following steps to build and run the s
421		420
422	- On-device performance collection	421	- On-device performance collection
423		422
424	- On-device performance collection directly measures the execution time of an operator on an Ascend AI Processor. This method is suitable for quickly locating operator performance issues in an on-device environment.	423	+ On-device performance collection directly measures the execution time of an operator on an Ascend AI Processor. This method is suitable for quickly locating operator performance issues in an on-device environment.
425		424
426	- Run operator tuning on the executable demo with `msopprof`:	425	+ Run operator tuning on the `matrix_transpose` executable with `msopprof`:
427		426
428	- ```	427	+ ```bash
429	- msopprof ./demo	428	+ msopprof ./matrix_transpose
		429	+ ```
		430	+
		431	+ - Performance data description
		432	+ After the command completes, a folder named "OPPROF_{timestamp}_XXX" is generated in the default directory. The performance data folder structure is as follows:
		433	+
		434	+ ```bash
		435	+ ├──dump # Raw performance data; users do not need to inspect it
		436	+ ├──ArithmeticUtilization.csv # Cube/Vector instruction cycle proportions
		437	+ ├──L2Cache.csv # L2 Cache hit rate; affects MTE2. Plan data transfer logic properly to increase the hit rate
		438	+ ├──Memory.csv # Read/write bandwidth rates of UB, L1, and main memory
		439	+ ├──MemoryL0.csv # Read/write bandwidth rates of L0A, L0B, and L0C
		440	+ ├──MemoryUB.csv # Read/write bandwidth rates from Vector and Scalar to UB
		441	+ ├──OpBasicInfo.csv # Basic operator information
		442	+ ├──PipeUtilization.csv # Durations and proportions of computation and data transfer units
		443	+ ├──ResourceConflictRatio.csv # Proportions of UB bank groups, bank conflicts, and resource conflicts among all instructions
		444	+ └──visualize_data.bin # MindStudio Insight presentation file
430	```	445	```
431		446
432	- - Performance data description	447	+ View the specific performance analysis results:
433	- After the command completes, a folder named "OPPROF_{timestamp}_XXX" will be generated in the default directory. The performance data folder structure is as follows:
434		448
435	- ```bash	449	+ ```bash
436	- ├──dump # Raw performance data; users do not need to inspect it	450	+ # View Task Duration and various metrics
437	- ├──ArithmeticUtilization.csv # Cube/Vector instruction cycle proportions	451	+ cat ./OPPROF_*/PipeUtilization.csv
438	- ├──L2Cache.csv # L2 Cache hit rate; affects MTE2. Plan data transfer logic properly to increase the hit rate	452	+ ```
439	- ├──Memory.csv # Read/write bandwidth rates of UB, L1, and main memory
440	- ├──MemoryL0.csv # Read/write bandwidth rates of L0A, L0B, and L0C
441	- ├──MemoryUB.csv # Read/write bandwidth rates from Vector and Scalar to UB
442	- ├──OpBasicInfo.csv # Basic operator information
443	- ├──PipeUtilization.csv # Durations and proportions of computation and data transfer units
444	- ├──ResourceConflictRatio.csv # Proportions of UB bank groups, bank conflicts, and resource conflicts among all instructions
445	- └──visualize_data.bin # MindStudio Insight presentation file
446	- ```
447		453
448	-View the specific performance analysis results:	454	+- Simulation performance collection
449		455
450	-```bash	456	+ Use `msopprof simulator` to perform simulation performance analysis and generate visualized instruction pipeline diagrams. Run the following commands:
451	-# View Task Duration and various metrics	457	+
452	-cat ./OPPROF_*/PipeUtilization.csv	458	+ ```bash
453	-```	459	+ SCENARIO_NUM=4 # Select an execution scenario. Valid values are 0-4.
		460	+ mkdir -p build && cd build; # Create and enter the build directory.
		461	+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO_NUM ..;make -j; # Build the project.
		462	+ msopprof simulator --soc-version=<soc_version> ./matrix_transpose
		463	+ ```
		464	+
		465	+ > Before using simulation tuning, add the `-g` compilation option to `CMakeLists.txt` to generate debug information. This allows the simulator to collect instruction pipeline diagrams. For how to obtain `soc_version` and for more simulation tuning information, see the [simulator sample](../../../01_simd_cpp_api/01_utilities/08_simulator/README_en.md).
		466	+ >
		467	+
		468	+ After the command completes, a folder named `OPPROF_{timestamp}_XXX` is generated in the current directory. The artifact structure is as follows:
		469	+
		470	+ ```text
		471	+ OPPROF_{timestamp}_XXX/
		472	+ ├── dump // Raw performance data; users do not need to inspect it
		473	+ └── simulator
		474	+ ├── core.veccore/ // Simulation instruction pipeline diagram files for each vector core
		475	+ └── visualize_data.bin // MindStudio Insight presentation file
		476	+ ```
		477	+
		478	+ After the command completes, open `visualize_data.bin` in MindStudio Insight to view the visualized instruction pipeline diagrams.
		479	+
		480	+For more information about how to use `msOpProf`, see [MindStudio Tool Tuning (msOpProf) Quick Start](https://www.hiascend.com/document/detail/zh/canncommercial/900/devaids/optool/docs/zh/quick_start/msopprof_quick_start.md).

Mexamples/05_simd_simt_hybrid/02_best_practices/simd_simt_matrix_transpose/matrix_transpose.asc+7-8

	@@ -31,12 +31,12 @@
31	#endif	31	#endif
32		32
33	namespace {	33	namespace {
34	-constexpr int TILE_DIM = 32;	34	+constexpr uint32_t TILE_DIM = 32;
35	-constexpr int TILE_PAD = 8; // MTE搬运要求32B对齐，使用32×40的UB布局	35	+constexpr uint32_t TILE_PAD = 8; // MTE搬运要求32B对齐，使用32×40的UB布局
36	-constexpr int TILE_PAD_STRIDE = TILE_DIM + TILE_PAD;	36	+constexpr uint32_t TILE_PAD_STRIDE = TILE_DIM + TILE_PAD;
37	constexpr uint32_t THREADS_2048 = 2048;	37	constexpr uint32_t THREADS_2048 = 2048;
38	constexpr uint32_t THREADS_PER_BLOCK = THREADS_2048;	38	constexpr uint32_t THREADS_PER_BLOCK = THREADS_2048;
39	-constexpr int TILES_PER_BLOCK = 2; // 2048线程对应2个32×32 tile	39	+constexpr uint32_t TILES_PER_BLOCK = 2; // 2048线程对应2个32×32 tile
40	constexpr uint32_t TILE_ELEMENTS = TILE_DIM * TILE_DIM;	40	constexpr uint32_t TILE_ELEMENTS = TILE_DIM * TILE_DIM;
41	constexpr uint32_t TILE_ROW_BYTES = TILE_DIM * sizeof(float);	41	constexpr uint32_t TILE_ROW_BYTES = TILE_DIM * sizeof(float);
42	constexpr uint32_t TILE_PAD_ELEMENTS = TILE_DIM * TILE_PAD_STRIDE;	42	constexpr uint32_t TILE_PAD_ELEMENTS = TILE_DIM * TILE_PAD_STRIDE;
	@@ -280,10 +280,9 @@ __global__ __vector__ void transpose_ub_2tile_loop_kernel(
280	uint32_t tiles_x = width / TILE_DIM;	280	uint32_t tiles_x = width / TILE_DIM;
281	uint32_t loop_start = block_idx * TILES_PER_BLOCK;	281	uint32_t loop_start = block_idx * TILES_PER_BLOCK;
282	uint32_t loop_step = block_num * TILES_PER_BLOCK;	282	uint32_t loop_step = block_num * TILES_PER_BLOCK;
283	- uint32_t loop_end = total_tiles;
284		283
285	// Case 2固定Thread Block数，每个Thread Block按gridDim.x步长循环处理多个tile组。	284	// Case 2固定Thread Block数，每个Thread Block按gridDim.x步长循环处理多个tile组。
286	- for (uint32_t tile_base = loop_start; tile_base < loop_end; tile_base += loop_step) {	285	+ for (uint32_t tile_base = loop_start; tile_base < total_tiles; tile_base += loop_step) {
287	asc_lock(PIPE_MTE2, SINGLE_BUFFER_MUTEX);	286	asc_lock(PIPE_MTE2, SINGLE_BUFFER_MUTEX);
288	copy_gm_2tile_to_ub(&in_tile[0][0][0], input, width, tiles_x, tile_base, total_tiles);	287	copy_gm_2tile_to_ub(&in_tile[0][0][0], input, width, tiles_x, tile_base, total_tiles);
289	asc_unlock(PIPE_MTE2, SINGLE_BUFFER_MUTEX);	288	asc_unlock(PIPE_MTE2, SINGLE_BUFFER_MUTEX);
	@@ -460,7 +459,7 @@ std::vector<float> run_case(const std::vector<float>& input_host, const std::vec
460	return output_host;	459	return output_host;
461	}	460	}
462		461
463	-int main(int argc, char** argv)	462	+int32_t main(int32_t argc, char** argv)
464	{	463	{
465	const uint32_t kernel_kind = SCENARIO_NUM;	464	const uint32_t kernel_kind = SCENARIO_NUM;
466	uint32_t matrix_size = 1024;	465	uint32_t matrix_size = 1024;
	@@ -478,7 +477,7 @@ int main(int argc, char** argv)
478	printf("Matrix shape: %u x %u\n", matrix_size, matrix_size);	477	printf("Matrix shape: %u x %u\n", matrix_size, matrix_size);
479		478
480	if (matrix_size - (matrix_size / TILE_DIM) * TILE_DIM != 0) {	479	if (matrix_size - (matrix_size / TILE_DIM) * TILE_DIM != 0) {
481	- std::fprintf(stderr, "Matrix shape must be multiple of TILE_DIM=%d.\n", TILE_DIM);	480	+ std::fprintf(stderr, "Matrix shape must be multiple of TILE_DIM=%u.\n", TILE_DIM);
482	return EXIT_FAILURE;	481	return EXIT_FAILURE;
483	}	482	}
484		483