LV.1

hxf0223

128积分0赞

8 帖子 23 回复 0 收藏

TA的动态

TA的帖子

TA的回复

cnrtConvertFloatToHalf转换慢 我的回复：#4jiapeiyuan回复1）你现在用的cnrtConvertFloatToHalf这个接口，不能做量化，它的作用是float32转float16，直接截断成16bit浮点数2）想要量化，要用cnrtCreateQuantizedParam + cnrtCastDataType 这两个接口来做，先设置pos/scale，生成量化参数变量，然后再cast。pos和scale需要你自己设置。以下几个公式可参考：positon = floor(log(max)) - (n - 2)scale = (2^(n-1) - 1) * 2^pos / maxq = round(r * scale / 2^pos)注：log是以2为底的，r表示实数，q表示量化后的数，max是float32里面绝对值最大的数。这是按照绝对值最大值的方式计算量化参数的方法。可参考cnml用户文档里面关于量化的描述，或者参考公网的量化方法，都是一致的展开非常感谢您帮我厘清这些错误，以及提供的量化公式！ 0

cnrtConvertFloatToHalf转换慢 我的回复：#1jiapeiyuan回复你是要把float转成fp16，还是int16？这两个接口不一样，int16的话，可以用cnrtCastDataType这个转成half（fp16）跟“向量乘”，在计算复杂度上不在一个量级，前者的逻辑和操作要复杂的多，所以这样对比不是很合适cnrtConvertFloatToHalf这个接口是在cpu上做的，我们已经做到尽量最优了，如果cpu支持avx，也使用avx指令集加速，所以我们找不到什么优化空间。如果你有别的方法，也可以替换成自己的算法展开您也提到了SIMD指令，SIMD指令的优势批处理的带宽跟并行计算。想问下，能否提供一个 \"float32数组 -> int16数组\" 的转换？因为我现在调用so里面的 \"cnrtConvertFloatToHalf\"，应该不能针对这一块作优化。另外，我不了解是使用什么算法来作量化的（量化这一块也不熟悉）。可以的话，望指点一二！ 0

NRAM中定义数组过大时，运行出错（不是编译链接出错） 我的回复：请问该如何计算所需的内存空。更新代码如下：#define SHR_BUFF32_MAX (1024*1024/4) #define BLOCK_SIZE (8192*1)//16384 #define CORE_DIM 4 #define STAGE 2 __mlu_func__ void kernel_row_mul_addb(bang::pipeline& pipe, float* __restrict__ nC, float* __restrict__ nA, float* __restrict__ nB, int count) { __nram__ float temp_buff[BLOCK_SIZE]; bang::vector::mul(pipe, temp_buff, nA, nB, count); bang::vector::add(pipe, nC, nC, temp_buff, count); } __mlu_func__ void kernel_row_mul_add(float* __restrict__ gA, float* __restrict__ sB, float* __restrict__ gC, int colNum, int rowNum, int rowA1, int rowA2, int rowNumB) { __nram__ float nbuffa[BLOCK_SIZE]; __nram__ float nbuffc[BLOCK_SIZE]; __nram__ float nbuffb[BLOCK_SIZE*STAGE]; bang::pipeline pipe0(0); bang::pipeline pipe1(1); bang::pipeline pipe4(2); bang::pipeline pipe5(3); for ( int row = rowA1; row < rowA2; row++ ) { size_t cp_sizeb = colNum * sizeof(float); //__bang_write_zero(nbuffc, BLOCK_SIZE); bang::memcpy_async(pipe5, nbuffc, gC+row*colNum, cp_sizeb, GDRAM2NRAM); pipe5.wait_copy_dram_to_nram(); bang::memcpy_async(pipe4, nbuffa, gA+row*colNum, cp_sizeb, GDRAM2NRAM); pipe4.wait_copy_dram_to_nram(); for (int rowb = 0; rowb < rowNumB; rowb += 2) { bang::memcpy_async(pipe0, nbuffb, sB + rowb*colNum, cp_sizeb, SRAM2NRAM); bang::memcpy_async(pipe1, nbuffb+BLOCK_SIZE, sB + (rowb+1)*colNum, cp_sizeb, SRAM2NRAM); pipe0.wait_copy_sram_to_nram(); kernel_row_mul_addb(pipe0, nbuffc, nbuffa, nbuffb, colNum); pipe1.wait_copy_sram_to_nram(); kernel_row_mul_addb(pipe1, nbuffc, nbuffa, nbuffb+BLOCK_SIZE, colNum); } // store result line buffer to C pipe5.wait_copy_nram_to_dram(); bang::memcpy_async(pipe5, gC+row*colNum, nbuffc, cp_sizeb, NRAM2GDRAM); } pipe5.wait_copy_nram_to_dram(); } __mlu_global__ void kernelMM(float* __restrict__ gA, float* __restrict__ gB, float* __restrict__ gC, int rowNum, int colNum) { __mlu_shared__ float shr_buffb[SHR_BUFF32_MAX]; int max_shr_lines = SHR_BUFF32_MAX / colNum; int rows_per_core = rowNum / taskDim; int start_row = rows_per_core * taskId; int end_row = start_row + rows_per_core; if (start_row >= rowNum) return; if (end_row > rowNum) end_row = rowNum; for (int i = 0; i < rowNum; i += max_shr_lines) { int cp_row_num = max_shr_lines; if ((cp_row_num+i) > rowNum) cp_row_num = rowNum-i; int cp_bytes = cp_row_num * colNum * sizeof(float); if (taskId == 0) { // printf("rowid %d, cp_row_num %d \\n", i, cp_row_num); __memcpy(shr_buffb, gB + i * colNum, cp_bytes, GDRAM2SRAM); } __sync_all(); cnscclBroadcast((void*)&shr_buffb, cp_bytes/sizeof(float), cnscclFloat, 0); __sync_all(); kernel_row_mul_add(gA, shr_buffb, gC, colNum, rowNum, start_row, end_row, max_shr_lines); } } int matrixMul(float* ptrC, float* ptrA, float* ptrB, int rowNum, int colNum) { float* deva, *devb, *devc; constexpr int loops = 1; cnrtQueue_t queue; CNRT_CHECK( cnrtCreateQueue(&queue) ); const size_t mem_size = rowNum * colNum * sizeof(float); CNRT_CHECK( cnrtMalloc((void**)(&deva), mem_size) ); CNRT_CHECK( cnrtMalloc((void**)(&devb), mem_size) ); CNRT_CHECK( cnrtMalloc((void**)(&devc), mem_size) ); CNRT_CHECK( cnrtMemcpy(deva, ptrA, mem_size, CNRT_MEM_TRANS_DIR_HOST2DEV) ); CNRT_CHECK( cnrtMemcpy(devb, ptrB, mem_size, CNRT_MEM_TRANS_DIR_HOST2DEV) ); CNRT_CHECK( cnrtMemcpy(devc, ptrC, mem_size, CNRT_MEM_TRANS_DIR_HOST2DEV) ); /*CNRT_CHECK( cnrtMemcpyAsync(deva, ptrA, mem_size, queue, CNRT_MEM_TRANS_DIR_HOST2DEV) ); CNRT_CHECK( cnrtMemcpyAsync(devb, ptrB, mem_size, queue, CNRT_MEM_TRANS_DIR_HOST2DEV) ); CNRT_CHECK( cnrtMemcpyAsync(devc, ptrC, mem_size, queue, CNRT_MEM_TRANS_DIR_HOST2DEV) );*/ // Create cnrt event handles. cnrtNotifier_t start, stop; CNRT_CHECK( cnrtCreateNotifier(&start) ); CNRT_CHECK( cnrtCreateNotifier(&stop) ); // Record timer start CNRT_CHECK( cnrtPlaceNotifier(start, queue) ); cnrtDim3_t dim = {16, 1, 1}; cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION4; for (int i = 0; i < loops; i++) { kernelMM<<<dim, ktype, queue>>>(deva, devb, devc, rowNum, colNum); } CNRT_CHECK( cnrtSyncQueue(queue) ); // Record timer stop CNRT_CHECK( cnrtPlaceNotifier(stop, queue) ); CNRT_CHECK( cnrtMemcpy(ptrC, devc, mem_size, CNRT_MEM_TRANS_DIR_DEV2HOST) ); float usecTotal = 0.0f; CNRT_CHECK(cnrtNotifierDuration(start, stop, &usecTotal)); printf("time spent executing by the MLU: %.3f ms\\n", usecTotal / (1000*loops) ); CNRT_CHECK( cnrtFree((void*)deva) ); CNRT_CHECK( cnrtFree((void*)devb) ); CNRT_CHECK( cnrtFree((void*)devc) ); CNRT_CHECK( cnrtDestroyQueue(queue) ); return 0; } int bang_dev_init() { cnrtDev_t dev; CNRT_CHECK(cnrtInit(0)); CNRT_CHECK(cnrtGetDeviceHandle(&dev, 0)); CNRT_CHECK(cnrtSetCurrentDevice(dev)); return 0; } int bang_dev_deinit() { cnrtDestroy(); return 0; } 0

NRAM中定义数组过大时，运行出错（不是编译链接出错） 我的回复：#3guowenxin回复79行按规定传一下参数试试，另外kernel里面的递归调用最好做一下所需总空间的判断展开你好，传参是帖子写的有问题。另外，没有递归调用。 0

使用异步传输(overlapping)，测试执行时间没有得到改善 我的回复：#2goodchong回复您好，请问您的寒武纪板卡是什么型号的？如果是MLU270的话矩阵乘最好要经过量化后使用__bang_conv计算。您目前的程序使用循环标量累加计算，性能肯定是很难达到满意的效果的。展开您好， 1，上面的代码忘记更新了，乘法操作后来改成了向量操作 __bang_mul。 2，使用bang_conv的话，我这边也正在进行测试，估计精度应该是满足不了；我的问题没说清楚，补充如下：上面的示例，分别使用memcpy，以及async_memcpy（overlapping），发现整个时间是一样的。另外一方面，测试发现时间主要消耗在GDRAMNRAM传输，所以怀疑上面测试代码里面写的async_memcpy没有起到overlapping作用。想请您分析下为什么会这样？ 0

NRAM中定义数组过大时，运行出错（不是编译链接出错） 我的回复：当定义的BLOCK_SIZE 为 16384，并且CPU给定的三个matrix也是16384*16384时，执行如上kernel函数，出现运行错误。请问哪位能给诊断下吗？CNRT: 4.9.1 4cd7a8a CPU time with assign 62030.4 ms 2021-06-29 15:05:06.027058: [cnrtError] [197417] [Card : 0] Error occurred in cnrtSyncQueue during calling driver interface. 2021-06-29 15:05:06.027102: [cnrtError] [197417] [Card : 0] Return value is 479, MLU_ERROR_UNKNOWN, means that error unknown 2021-06-29 15:05:06.027108: [cnrtError] [197417] [Card : 0] mlu unfinished! for more information, please use core dump analysis tools 2021-06-29 15:05:06.027155: [cnrtError] [197417] [Card : 0] [/meda_home/roderick/wscuda/mlu270/bang-mix/matrix_matrix/mm_mul.mlu:122] CNRT error, code=632046(Errors occurred from driver functions that are returned via CNRT. See the detailed error messages on the screen. ) "cnrtSyncQueue(queue)" 2021-06-29 15:05:06.027196: [cnrtError] [197417] [Card : 0] Error occurred in cnrtPlaceNotifier during calling driver interface. 2021-06-29 15:05:06.027204: [cnrtError] [197417] [Card : 0] Return value is 505, MLU_NOTIFIER_ERROR_PLACE_NOTIFY, means that place nofify error 2021-06-29 15:05:06.027211: [cnrtError] [197417] [Card : 0] [/meda_home/roderick/wscuda/mlu270/bang-mix/matrix_matrix/mm_mul.mlu:125] CNRT error, code=632046(Errors occurred from driver functions that are returned via CNRT. See the detailed error messages on the screen. ) "cnrtPlaceNotifier(stop, queue)" 0

使用异步传输(overlapping)，测试执行时间没有得到改善 我的回复：写了一个浮点矩阵乘法函数 C=A*B。使用pipeline::async_memcpy，发现执行时间没有减少。使用512*512的矩阵，我这边测试时间都是591ms（kernel执行时间）。想请问是哪里写的不对吗？ 0

请教如何使用异步传输 我的回复：参照例程 BANG/bang-c++/union4/pipeline_nmem_union4。写了一个测试程序，与测试例程有些差别。测试发现使用异步传输，运行时间跟同步传输一样。请问是哪里写的有问题吗？异步传输相关的代码行号：35~37, 45,46, 73~77, 80. 0

MLU270性能问题：请问这样写的代码，是同时开启16个核心计算吗 我的回复：补充说下：已经去掉了初始化，加载等overhead时间。去掉计算，host<-->device传输时间大约 2ms。 0

MLU270及开发平台是否支持 IEEE 754 我的回复：您好，不好意思，我的疑问还是没有解决。可能这样问比较直观些：1，给定一个matrix<complex<double>>；2，用FFTW3计算这个matrix的FFT结果；3，用MLU270同样对这个matrix计算FFT结果；4，比较两次计算结果coeff，比如e-? 0