simt-c conversion: trilu-ops-math-AtomGit

已合并

simt-c conversion: trilu #2691

renshuojian创建于 5月12日从rsj007/ops-math_3118： master 合入到 cann/ops-math： master

simt-c conversion: trilu #2691

已合并

renshuojian创建于 5月12日

从rsj007/ops-math_3118 master合入到cann/ops-math master

共 1 个文件变更+6-5

Mexperimental/math/trilu/op_kernel/arch35/trilu_simt.h+6-5

	@@ -15,6 +15,7 @@
15	#include "kernel_tiling/kernel_tiling.h"	15	#include "kernel_tiling/kernel_tiling.h"
16	#include "trilu_tiling_data.h"	16	#include "trilu_tiling_data.h"
17	#include "trilu_tiling_key.h"	17	#include "trilu_tiling_key.h"
		18	+#include "simt_api/asc_simt.h"
18		19
19	namespace NsTrilu {	20	namespace NsTrilu {
20		21
	@@ -27,11 +28,11 @@ inline void OpTriluSimt(int64_t totalElements, int64_t diagonal, int32_t upper,
27	{	28	{
28	int64_t matrixSize = h * w;	29	int64_t matrixSize = h * w;
29	for (uint64_t index = static_cast<uint64_t>(	30	for (uint64_t index = static_cast<uint64_t>(
30	- AscendC::Simt::GetBlockIdx() * AscendC::Simt::GetThreadNum()	31	+ blockIdx.x * blockDim.x
31	- + AscendC::Simt::GetThreadIdx());	32	+ + threadIdx.x);
32	index < static_cast<uint64_t>(totalElements);	33	index < static_cast<uint64_t>(totalElements);
33	index += static_cast<uint64_t>(	34	index += static_cast<uint64_t>(
34	- AscendC::Simt::GetThreadNum() * AscendC::Simt::GetBlockNum())) {	35	+ blockDim.x * gridDim.x)) {
35	int64_t matrixOffset = static_cast<int64_t>(index) % matrixSize;	36	int64_t matrixOffset = static_cast<int64_t>(index) % matrixSize;
36	int64_t row = matrixOffset / w;	37	int64_t row = matrixOffset / w;
37	int64_t col = matrixOffset % w;	38	int64_t col = matrixOffset % w;
	@@ -47,8 +48,8 @@ __aicore__ inline void Process(GM_ADDR x, GM_ADDR y, const TriluTilingData* tili
47	int64_t totalElements = tilingData->totalElements;	48	int64_t totalElements = tilingData->totalElements;
48	__gm__ T* x_gm = (__gm__ T*)x;	49	__gm__ T* x_gm = (__gm__ T*)x;
49	__gm__ T* y_gm = (__gm__ T*)y;	50	__gm__ T* y_gm = (__gm__ T*)y;
50	- AscendC::Simt::VF_CALL<OpTriluSimt<T>>(	51	+ asc_vf_call<OpTriluSimt<T>>(
51	- AscendC::Simt::Dim3(512), totalElements,	52	+ dim3(512), totalElements,
52	tilingData->diagonal, tilingData->upper,	53	tilingData->diagonal, tilingData->upper,
53	tilingData->h, tilingData->w,	54	tilingData->h, tilingData->w,
54	x_gm, y_gm);	55	x_gm, y_gm);