[INFO] Initializing environment for https://gitcode.com/pre-commit-clang/mirrors-clang-format. [INFO] Installing environment for https://gitcode.com/pre-commit-clang/mirrors-clang-format. [INFO] Once installed this environment will be reused. [INFO] This may take a few minutes... clang-format.............................................................Failed - hook id: clang-format - files were modified by this hook Check....................................................................Passed All changes made by hooks: diff --git a/blas/axpy/arch22/caxpy_host.cpp b/blas/axpy/arch22/caxpy_host.cpp index d11f1da..e1abdeb 100644 --- a/blas/axpy/arch22/caxpy_host.cpp +++ b/blas/axpy/arch22/caxpy_host.cpp @@ -39,8 +39,7 @@ constexpr uint32_t FLOATS_PER_COMPLEX = 2; constexpr uint32_t K_FACTOR_4 = 4; constexpr uint32_t DEFAULT_VECTOR_NUM = 40; static_assert( - DEFAULT_VECTOR_NUM == CAXPY_MAX_VECTOR_CORES, - "Host block cap must match the tiling struct's per-block array size"); + DEFAULT_VECTOR_NUM == CAXPY_MAX_VECTOR_CORES, "Host block cap must match the tiling struct's per-block array size"); constexpr uint32_t MAX_DATA_COUNT = 38 * 1024 / sizeof(float); constexpr uint32_t STRIDED_TILE_COMPLEX_COUNT = 512; constexpr uint32_t BLOCKED_FLOATS_PER_COMPLEX = 8; @@ -170,10 +169,9 @@ void GenDenseSwapOffsets(uint32_t* offsets) void GenXNormalizeOffsets(uint32_t* offsets, uint32_t xIncrement) { for (uint32_t i = 0; i < STRIDED_TILE_COMPLEX_COUNT; ++i) { - const uint32_t sourceByteOffset = - UseCaxpySharedXSpanPipeline(static_cast(xIncrement)) ? - i * xIncrement * FLOATS_PER_COMPLEX * sizeof(float) : - i * BLOCKED_FLOATS_PER_COMPLEX * sizeof(float); + const uint32_t sourceByteOffset = UseCaxpySharedXSpanPipeline(static_cast(xIncrement)) ? + i * xIncrement * FLOATS_PER_COMPLEX * sizeof(float) : + i * BLOCKED_FLOATS_PER_COMPLEX * sizeof(float); for (uint32_t lane = 0; lane < BLOCKED_FLOATS_PER_COMPLEX; ++lane) { offsets[i * BLOCKED_FLOATS_PER_COMPLEX + lane] = lane == 1 ? sourceByteOffset + sizeof(float) : sourceByteOffset; @@ -217,9 +215,8 @@ aclblasStatus_t GetCaxpyMaskCache(aclblasHandle_t handle, uint64_t absIncx, uint if (h->caxpy_mask_cache == nullptr) { void* buffer = nullptr; const aclError aclRet = aclrtMalloc(&buffer, MASK_DATA_COUNT * sizeof(uint32_t), ACL_MEM_MALLOC_HUGE_FIRST); - CHECK_RET( - aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); - return ACLBLAS_STATUS_ALLOC_FAILED); + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", aclRet); + return ACLBLAS_STATUS_ALLOC_FAILED); h->caxpy_mask_cache = buffer; } if (h->caxpy_mask_cache_incx != absIncx || h->caxpy_mask_cache_incy != absIncy) { @@ -232,9 +229,8 @@ aclblasStatus_t GetCaxpyMaskCache(aclblasHandle_t handle, uint64_t absIncx, uint const aclError aclRet = aclrtMemcpy( h->caxpy_mask_cache, MASK_DATA_COUNT * sizeof(uint32_t), maskHost.data(), MASK_DATA_COUNT * sizeof(uint32_t), ACL_MEMCPY_HOST_TO_DEVICE); - CHECK_RET( - aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); - return ACLBLAS_STATUS_INTERNAL_ERROR); + CHECK_RET(aclRet == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", aclRet); + return ACLBLAS_STATUS_INTERNAL_ERROR); h->caxpy_mask_cache_incx = absIncx; h->caxpy_mask_cache_incy = absIncy; } @@ -301,10 +297,9 @@ aclblasStatus_t BuildCaxpyExecution( const CaxpyKernelVariant kernelVariant = SelectCaxpyKernelVariant(static_cast(n), incx, incy, useStridedKernel, numBlocks); auto* h = handle; - CHECK_RET( - sizeof(CaxpyTilingData) <= GetEffectiveWorkspaceSize(h), - LOG_PRINT("workspace need %zu, available %zu\n", sizeof(CaxpyTilingData), GetEffectiveWorkspaceSize(h)); - return ACLBLAS_STATUS_EXECUTION_FAILED); + CHECK_RET(sizeof(CaxpyTilingData) <= GetEffectiveWorkspaceSize(h), + LOG_PRINT("workspace need %zu, available %zu\n", sizeof(CaxpyTilingData), GetEffectiveWorkspaceSize(h)); + return ACLBLAS_STATUS_EXECUTION_FAILED); uint8_t* maskDevice = nullptr; if (kernelVariant != CaxpyKernelVariant::STRIDED_SCALAR) { const aclblasStatus_t maskStatus = GetCaxpyMaskCache(handle, absIncx, absIncy, &maskDevice); @@ -347,15 +342,13 @@ aclblasStatus_t AllocatePackedBuffers(const CaxpyExecution& execution, CaxpyPack const size_t bytes = static_cast(execution.n) * sizeof(aclblasComplex); if (execution.packX) { const aclError result = aclrtMalloc(reinterpret_cast(&buffers.x), bytes, ACL_MEM_MALLOC_HUGE_FIRST); - CHECK_RET( - result == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", result); - return ACLBLAS_STATUS_ALLOC_FAILED); + CHECK_RET(result == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", result); + return ACLBLAS_STATUS_ALLOC_FAILED); } if (execution.packY) { const aclError result = aclrtMalloc(reinterpret_cast(&buffers.y), bytes, ACL_MEM_MALLOC_HUGE_FIRST); - CHECK_RET( - result == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", result); FreePackedBuffers(buffers); - return ACLBLAS_STATUS_ALLOC_FAILED); + CHECK_RET(result == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", result); + FreePackedBuffers(buffers); return ACLBLAS_STATUS_ALLOC_FAILED); } return ACLBLAS_STATUS_SUCCESS; } @@ -370,9 +363,8 @@ aclblasStatus_t CopyStridedVector( uint8_t* source = unpack ? packed + packedOffset : strided + stridedOffset; const aclError result = aclrtMemcpyAsync( destination, sizeof(aclblasComplex), source, sizeof(aclblasComplex), ACL_MEMCPY_DEVICE_TO_DEVICE, stream); - CHECK_RET( - result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpyAsync failed. ERROR: %d\n", result); - return ACLBLAS_STATUS_INTERNAL_ERROR); + CHECK_RET(result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpyAsync failed. ERROR: %d\n", result); + return ACLBLAS_STATUS_INTERNAL_ERROR); } return ACLBLAS_STATUS_SUCCESS; } @@ -400,9 +392,8 @@ aclblasStatus_t UploadTiling(const CaxpyExecution& execution, const CaxpyTilingD { const aclError result = aclrtMemcpyAsync( execution.tilingDevice, sizeof(tiling), &tiling, sizeof(tiling), ACL_MEMCPY_HOST_TO_DEVICE, execution.stream); - CHECK_RET( - result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpyAsync failed. ERROR: %d\n", result); - return ACLBLAS_STATUS_INTERNAL_ERROR); + CHECK_RET(result == ACL_SUCCESS, LOG_PRINT("aclrtMemcpyAsync failed. ERROR: %d\n", result); + return ACLBLAS_STATUS_INTERNAL_ERROR); return ACLBLAS_STATUS_SUCCESS; } @@ -413,8 +404,8 @@ aclblasStatus_t UploadTiling(const CaxpyExecution& execution, const CaxpyTilingD CaxpyTilingData BuildStridedTiling( const CaxpyExecution& execution, const aclblasComplex& alpha, uint32_t count, uint32_t blocks) { - CaxpyTilingData tiling = CalTilingData( - count, blocks, alpha.real, alpha.imag, execution.incx, execution.incy, 0, count); + CaxpyTilingData tiling = + CalTilingData(count, blocks, alpha.real, alpha.imag, execution.incx, execution.incy, 0, count); tiling.activeBlocks = blocks; // Only the shared pipeline's run_shared_pipeline/shared_tile_bounds read // this field; the disjoint and scalar kernel entry points ignore it and @@ -514,9 +505,8 @@ aclblasStatus_t FinishCaxpy(const CaxpyExecution& execution, aclblasComplex* y, CHECK_RET(status == ACLBLAS_STATUS_SUCCESS, FreePackedBuffers(buffers); return status); } const aclError result = aclrtSynchronizeStream(execution.stream); - CHECK_RET( - result == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", result); - FreePackedBuffers(buffers); return ACLBLAS_STATUS_INTERNAL_ERROR); + CHECK_RET(result == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", result); + FreePackedBuffers(buffers); return ACLBLAS_STATUS_INTERNAL_ERROR); FreePackedBuffers(buffers); return ACLBLAS_STATUS_SUCCESS; } diff --git a/blas/axpy/arch22/caxpy_kernel.cpp b/blas/axpy/arch22/caxpy_kernel.cpp index ae1574d..bea4d2a 100644 --- a/blas/axpy/arch22/caxpy_kernel.cpp +++ b/blas/axpy/arch22/caxpy_kernel.cpp @@ -52,8 +52,8 @@ constexpr uint32_t CAXPY_DENSE_UB_COEFFICIENT_COUNT = 8; // even though the dense and strided kernels are otherwise independent code // paths. See caxpy_host.cpp's DENSE_SWAP_OFFSET_START for the authoritative // host-side computation this mirrors. -constexpr uint32_t CAXPY_DENSE_SWAP_SKIP_X_OFFSET_WORDS = 512 * 8 + 64; // == CAXPY_X_OFFSET_COUNT -constexpr uint32_t CAXPY_DENSE_SWAP_SKIP_SPAN_OFFSET_WORDS = 2 * (512 * 2 + 64); // == 2 * CAXPY_SPAN_OFFSET_COUNT +constexpr uint32_t CAXPY_DENSE_SWAP_SKIP_X_OFFSET_WORDS = 512 * 8 + 64; // == CAXPY_X_OFFSET_COUNT +constexpr uint32_t CAXPY_DENSE_SWAP_SKIP_SPAN_OFFSET_WORDS = 2 * (512 * 2 + 64); // == 2 * CAXPY_SPAN_OFFSET_COUNT constexpr uint32_t CAXPY_DENSE_SWAP_OFFSET_START = CAXPY_DENSE_TILE_FLOATS + CAXPY_DENSE_SWAP_SKIP_X_OFFSET_WORDS + CAXPY_DENSE_SWAP_SKIP_SPAN_OFFSET_WORDS; static_assert( @@ -177,8 +177,8 @@ __aicore__ __inline__ __attribute__((always_inline)) void copy_complex_ub2gm_sig } __aicore__ __inline__ __attribute__((always_inline)) void caxpy_compute_planar( - __ubuf__ float* resultReal, __ubuf__ float* resultImag, __ubuf__ float* sourceReal, - __ubuf__ float* sourceImag, float alphaReal, float alphaImag, uint32_t repeat) + __ubuf__ float* resultReal, __ubuf__ float* resultImag, __ubuf__ float* sourceReal, __ubuf__ float* sourceImag, + float alphaReal, float alphaImag, uint32_t repeat) { // Keep every vector address 32-byte aligned. An interleaved implementation using // source + 1 is mathematically equivalent, but faults on C220 vector instructions. @@ -190,7 +190,6 @@ __aicore__ __inline__ __attribute__((always_inline)) void caxpy_compute_planar( pipe_barrier(PIPE_V); } - __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_load( __gm__ float* gmX, __ubuf__ float* ubX, uint32_t copyLength, uint32_t eventId) { @@ -199,16 +198,15 @@ __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_load( } __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_compute( - __ubuf__ float* ubX, __ubuf__ float* ubOut, __ubuf__ uint32_t* ubOffset, - __ubuf__ float* ubCoefficient, float alphaReal, uint32_t computeLength, uint32_t eventId) + __ubuf__ float* ubX, __ubuf__ float* ubOut, __ubuf__ uint32_t* ubOffset, __ubuf__ float* ubCoefficient, + float alphaReal, uint32_t computeLength, uint32_t eventId) { const uint32_t repeat = (computeLength + 63) / 64; wait_flag(PIPE_MTE2, PIPE_V, eventId); // Turn [real, imag] into [imag, real]. The alternating coefficient then // applies [-alphaImag, +alphaImag] without splitting into planar buffers. - vgather( - reinterpret_cast<__ubuf__ uint32_t*>(ubOut), ubOffset, (uintptr_t)ubX, 8, repeat); + vgather(reinterpret_cast<__ubuf__ uint32_t*>(ubOut), ubOffset, (uintptr_t)ubX, 8, repeat); pipe_barrier(PIPE_V); vmuls(ubX, ubX, alphaReal, repeat, 1, 1, 8, 8); vmul(ubOut, ubOut, ubCoefficient, repeat, 1, 1, 0, 8, 8, 0); @@ -239,8 +237,8 @@ __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_prepare( } __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_load_next_if_any( - __gm__ float* gmX, __ubuf__ float* ubX[2], uint32_t offset, uint32_t calNum, uint32_t tileFloats, - uint32_t tile, uint32_t tileCount, uint32_t current) + __gm__ float* gmX, __ubuf__ float* ubX[2], uint32_t offset, uint32_t calNum, uint32_t tileFloats, uint32_t tile, + uint32_t tileCount, uint32_t current) { if (tile + 1 >= tileCount) return; @@ -253,8 +251,8 @@ __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_load_next_ } __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_run( - __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag, - uint32_t offset, uint32_t calNum) + __gm__ float* gmX, __gm__ uint32_t* gmAug, __gm__ float* gmY, float alphaReal, float alphaImag, uint32_t offset, + uint32_t calNum) { constexpr uint32_t tileFloats = CAXPY_DENSE_TILE_FLOATS; __ubuf__ float* ubOut[2] = { @@ -281,13 +279,11 @@ __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_run( const uint32_t current = tile & 1; const uint32_t currentEvent = current == 0 ? EVENT_ID0 : EVENT_ID1; const uint32_t currentOffset = tile * tileFloats; - const uint32_t currentCount = - calNum - currentOffset > tileFloats ? tileFloats : calNum - currentOffset; + const uint32_t currentCount = calNum - currentOffset > tileFloats ? tileFloats : calNum - currentOffset; caxpy_dense_load_next_if_any(gmX, ubX, offset, calNum, tileFloats, tile, tileCount, current); - caxpy_dense_compute( - ubX[current], ubOut[current], ubOffset, ubCoefficient, alphaReal, tileFloats, currentEvent); + caxpy_dense_compute(ubX[current], ubOut[current], ubOffset, ubCoefficient, alphaReal, tileFloats, currentEvent); caxpy_dense_store(gmY + offset + currentOffset, ubOut[current], currentCount, currentEvent); } @@ -297,7 +293,6 @@ __aicore__ __inline__ __attribute__((always_inline)) void caxpy_dense_run( set_atomic_none(); } - constexpr uint32_t CAXPY_MAX_DATA_COUNT = 38 * 1024 / sizeof(float); constexpr uint32_t CAXPY_UB_COMPLEX_CAPACITY = CAXPY_MAX_DATA_COUNT / 2; // CAXPY_STRIDED_TILE_COUNT is defined once in caxpy_dispatch_policy.h (shared @@ -490,8 +485,8 @@ __aicore__ __inline__ __attribute__((always_inline)) CaxpyStridedUb disjoint_ub_ } __aicore__ __inline__ __attribute__((always_inline)) void disjoint_load_tile( - CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, - int64_t incx, uint32_t eventId) + CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, int64_t incx, + uint32_t eventId) { copy_complex_gm2ub_signed(ub.x, ub.xDense, gmX, logicalOffset, totalCount, count, incx); set_flag(PIPE_MTE2, PIPE_V, eventId); @@ -506,14 +501,13 @@ __aicore__ __inline__ __attribute__((always_inline)) void disjoint_compute_tile( } __aicore__ __inline__ __attribute__((always_inline)) void strided_store_tile( - CaxpyStridedUb& ub, __gm__ float* gmY, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, - int64_t incy, uint32_t eventId) + CaxpyStridedUb& ub, __gm__ float* gmY, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, int64_t incy, + uint32_t eventId) { __ubuf__ float* source = ub.work; if (incy == 1) { const uint32_t repeat = (count * 2 + 63) / 64; - vgather( - reinterpret_cast<__ubuf__ uint32_t*>(ub.workDense), ub.storeOffset, (uintptr_t)ub.work, 8, repeat); + vgather(reinterpret_cast<__ubuf__ uint32_t*>(ub.workDense), ub.storeOffset, (uintptr_t)ub.work, 8, repeat); pipe_barrier(PIPE_V); source = ub.workDense; } @@ -528,14 +522,12 @@ __aicore__ __inline__ __attribute__((always_inline)) void disjoint_preload_offse { copy_vec_gm2ub_uint32(ping.computeOffset, gmAug, CAXPY_COMPUTE_OFFSET_COUNT); if (incy == 1) - copy_vec_gm2ub_uint32( - ping.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT); + copy_vec_gm2ub_uint32(ping.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT); } __aicore__ __inline__ __attribute__((always_inline)) void disjoint_load_next_if_any( - CaxpyStridedUb& ping, CaxpyStridedUb& pong, __gm__ float* gmX, uint32_t logicalOffset, - uint32_t totalComplexCount, uint32_t complexCount, uint32_t tileCapacity, int64_t incx, uint32_t tile, - uint32_t tileCount, uint32_t current) + CaxpyStridedUb& ping, CaxpyStridedUb& pong, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalComplexCount, + uint32_t complexCount, uint32_t tileCapacity, int64_t incx, uint32_t tile, uint32_t tileCount, uint32_t current) { if (tile + 1 >= tileCount) return; @@ -543,11 +535,9 @@ __aicore__ __inline__ __attribute__((always_inline)) void disjoint_load_next_if_ const uint32_t nextEvent = next == 0 ? EVENT_ID0 : EVENT_ID1; CaxpyStridedUb& nextUb = next == 0 ? ping : pong; const uint32_t nextOffset = (tile + 1) * tileCapacity; - const uint32_t nextCount = - complexCount - nextOffset > tileCapacity ? tileCapacity : complexCount - nextOffset; + const uint32_t nextCount = complexCount - nextOffset > tileCapacity ? tileCapacity : complexCount - nextOffset; wait_flag(PIPE_MTE3, PIPE_MTE2, nextEvent); - disjoint_load_tile( - nextUb, gmX, logicalOffset + nextOffset, totalComplexCount, nextCount, incx, nextEvent); + disjoint_load_tile(nextUb, gmX, logicalOffset + nextOffset, totalComplexCount, nextCount, incx, nextEvent); } __aicore__ __inline__ __attribute__((always_inline)) void run_disjoint_pipeline( @@ -603,8 +593,8 @@ __aicore__ __inline__ __attribute__((always_inline)) CaxpyStridedUb shared_ub_la } __aicore__ __inline__ __attribute__((always_inline)) void shared_load_tile( - CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, - int64_t incx, uint32_t waitEvent, uint32_t readyEvent, bool waitForRaw) + CaxpyStridedUb& ub, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalCount, uint32_t count, int64_t incx, + uint32_t waitEvent, uint32_t readyEvent, bool waitForRaw) { // MTE3 owns each work buffer after a store. Hand it back through MTE2, // which is the event direction used by the stable dense pipeline. @@ -666,14 +656,12 @@ __aicore__ __inline__ __attribute__((always_inline)) void shared_preload_offsets { strided_preload_offsets(ping, gmAug, true, false); if (incy == 1) - copy_vec_gm2ub_uint32( - ping.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT); + copy_vec_gm2ub_uint32(ping.storeOffset, gmAug + CAXPY_STORE_OFFSET_START, CAXPY_SPAN_OFFSET_COUNT); } __aicore__ __inline__ __attribute__((always_inline)) void shared_load_next_if_any( - CaxpyStridedUb& ping, CaxpyStridedUb& pong, __gm__ float* gmX, uint32_t logicalOffset, - uint32_t totalComplexCount, uint32_t complexCount, uint32_t tileCapacity, int64_t incx, uint32_t tile, - uint32_t tileCount, uint32_t current) + CaxpyStridedUb& ping, CaxpyStridedUb& pong, __gm__ float* gmX, uint32_t logicalOffset, uint32_t totalComplexCount, + uint32_t complexCount, uint32_t tileCapacity, int64_t incx, uint32_t tile, uint32_t tileCount, uint32_t current) { if (tile + 1 >= tileCount) return; @@ -786,37 +774,45 @@ extern "C" __global__ __aicore__ __vector__ void caxpy( } #if __DAV_C220_VEC__ -#define DEFINE_CAXPY_STRIDED_KERNEL(KERNEL, PIPELINE) \ -extern "C" __global__ __aicore__ __vector__ void KERNEL( \ - __gm__ float* __restrict__ x, __gm__ uint32_t* __restrict__ aug, __gm__ float* __restrict__ y, \ - __gm__ float* __restrict__ tilingGm) \ -{ \ - const uint32_t vecIdx = AscendC::GetBlockIdx(); \ - set_mask_norm(); set_vector_mask((uint64_t)-1, (uint64_t)-1); \ - auto tiling = reinterpret_cast<__gm__ CaxpyTilingData*>(tilingGm); \ - const uint32_t numBlocks = tiling->activeBlocks; \ - const float alphaReal = tiling->alphaReal; \ - const float alphaImag = tiling->alphaImag; \ - const uint32_t totalN = tiling->totalN; \ - const int64_t incx = tiling->incx; \ - const int64_t incy = tiling->incy; \ - const uint32_t tileCapacity = tiling->sharedTileCapacity; \ - const uint32_t waveCapacity = numBlocks * CAXPY_STRIDED_MAX_PER_BLOCK_PER_WAVE; \ - for (uint32_t logicalBase = 0; logicalBase < totalN; logicalBase += waveCapacity) { \ - const uint32_t waveN = totalN - logicalBase > waveCapacity ? waveCapacity : totalN - logicalBase; \ - const uint32_t waveBlocks = numBlocks < waveN ? numBlocks : waveN; \ - if (vecIdx >= waveBlocks) continue; \ - const uint32_t rowNum = waveN / waveBlocks; const uint32_t remain = waveN % waveBlocks; \ - const uint32_t myCount = rowNum + (vecIdx < remain ? 1 : 0); \ - const uint32_t myStart = logicalBase + vecIdx * rowNum + (vecIdx < remain ? vecIdx : remain); \ - if (incy > 0) PIPELINE(x, aug, y, alphaReal, alphaImag, myStart, totalN, myCount, incx, incy, tileCapacity); \ - else { const uint32_t mirroredStart = totalN - myStart - myCount; PIPELINE(x, aug, y, alphaReal, alphaImag, mirroredStart, totalN, myCount, -incx, -incy, tileCapacity); } \ - } \ -} +#define DEFINE_CAXPY_STRIDED_KERNEL(KERNEL, PIPELINE) \ + extern "C" __global__ __aicore__ __vector__ void KERNEL( \ + __gm__ float* __restrict__ x, __gm__ uint32_t* __restrict__ aug, __gm__ float* __restrict__ y, \ + __gm__ float* __restrict__ tilingGm) \ + { \ + const uint32_t vecIdx = AscendC::GetBlockIdx(); \ + set_mask_norm(); \ + set_vector_mask((uint64_t)-1, (uint64_t)-1); \ + auto tiling = reinterpret_cast<__gm__ CaxpyTilingData*>(tilingGm); \ + const uint32_t numBlocks = tiling->activeBlocks; \ + const float alphaReal = tiling->alphaReal; \ + const float alphaImag = tiling->alphaImag; \ + const uint32_t totalN = tiling->totalN; \ + const int64_t incx = tiling->incx; \ + const int64_t incy = tiling->incy; \ + const uint32_t tileCapacity = tiling->sharedTileCapacity; \ + const uint32_t waveCapacity = numBlocks * CAXPY_STRIDED_MAX_PER_BLOCK_PER_WAVE; \ + for (uint32_t logicalBase = 0; logicalBase < totalN; logicalBase += waveCapacity) { \ + const uint32_t waveN = totalN - logicalBase > waveCapacity ? waveCapacity : totalN - logicalBase; \ + const uint32_t waveBlocks = numBlocks < waveN ? numBlocks : waveN; \ + if (vecIdx >= waveBlocks) \ + continue; \ + const uint32_t rowNum = waveN / waveBlocks; \ + const uint32_t remain = waveN % waveBlocks; \ + const uint32_t myCount = rowNum + (vecIdx < remain ? 1 : 0); \ + const uint32_t myStart = logicalBase + vecIdx * rowNum + (vecIdx < remain ? vecIdx : remain); \ + if (incy > 0) \ + PIPELINE(x, aug, y, alphaReal, alphaImag, myStart, totalN, myCount, incx, incy, tileCapacity); \ + else { \ + const uint32_t mirroredStart = totalN - myStart - myCount; \ + PIPELINE(x, aug, y, alphaReal, alphaImag, mirroredStart, totalN, myCount, -incx, -incy, tileCapacity); \ + } \ + } \ + } #else -#define DEFINE_CAXPY_STRIDED_KERNEL(KERNEL, PIPELINE) \ -extern "C" __global__ __aicore__ __vector__ void KERNEL( \ - __gm__ float*, __gm__ uint32_t*, __gm__ float*, __gm__ float*) {} +#define DEFINE_CAXPY_STRIDED_KERNEL(KERNEL, PIPELINE) \ + extern "C" __global__ __aicore__ __vector__ void KERNEL( \ + __gm__ float*, __gm__ uint32_t*, __gm__ float*, __gm__ float*) \ + {} #endif DEFINE_CAXPY_STRIDED_KERNEL(caxpy_strided_shared, caxpy_strided_shared_run) DEFINE_CAXPY_STRIDED_KERNEL(caxpy_strided_disjoint, caxpy_strided_disjoint_run) @@ -837,13 +833,11 @@ void caxpy_strided_shared_kernel_do( void caxpy_strided_disjoint_kernel_do( GM_ADDR x, GM_ADDR maskBuf, GM_ADDR y, GM_ADDR workSpace, GM_ADDR tilingGm, uint32_t numBlocks, void* stream) { - caxpy_strided_disjoint<<>>( - (float*)x, (uint32_t*)maskBuf, (float*)y, (float*)tilingGm); + caxpy_strided_disjoint<<>>((float*)x, (uint32_t*)maskBuf, (float*)y, (float*)tilingGm); } #if __DAV_C220_VEC__ -__aicore__ __inline__ __attribute__((always_inline)) void scalar_copy_complex_in( - __ubuf__ float* dst, __gm__ float* src) +__aicore__ __inline__ __attribute__((always_inline)) void scalar_copy_complex_in(__ubuf__ float* dst, __gm__ float* src) { copy_gm_to_ubuf_align_b32(dst, src, 0, 1, 2 * sizeof(float), 0, 0, 0, 0); } @@ -866,10 +860,8 @@ __aicore__ __inline__ __attribute__((always_inline)) void scalar_compute_all( set_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0); for (uint32_t local = 0; local < complexCount; ++local) { const uint32_t logical = logicalOffset + local; - const uint32_t xPhysical = - incx > 0 ? logical * absIncx : (totalComplexCount - 1 - logical) * absIncx; - const uint32_t yPhysical = - incy > 0 ? logical * absIncy : (totalComplexCount - 1 - logical) * absIncy; + const uint32_t xPhysical = incx > 0 ? logical * absIncx : (totalComplexCount - 1 - logical) * absIncx; + const uint32_t yPhysical = incy > 0 ? logical * absIncy : (totalComplexCount - 1 - logical) * absIncy; wait_flag(PIPE_MTE3, PIPE_MTE2, EVENT_ID0); scalar_copy_complex_in(ubX, gmX + xPhysical * 2); @@ -915,8 +907,7 @@ extern "C" __global__ __aicore__ __vector__ void caxpy_strided_scalar( #endif } -void caxpy_strided_scalar_kernel_do( - GM_ADDR x, GM_ADDR y, GM_ADDR tilingGm, uint32_t numBlocks, void* stream) +void caxpy_strided_scalar_kernel_do(GM_ADDR x, GM_ADDR y, GM_ADDR tilingGm, uint32_t numBlocks, void* stream) { caxpy_strided_scalar<<>>((float*)x, (float*)y, (float*)tilingGm); } diff --git a/blas/axpy/arch22/caxpy_tiling_data.h b/blas/axpy/arch22/caxpy_tiling_data.h index a047d1d..ce69bf3 100644 --- a/blas/axpy/arch22/caxpy_tiling_data.h +++ b/blas/axpy/arch22/caxpy_tiling_data.h @@ -36,8 +36,8 @@ struct CaxpyTilingData { }; float alphaReal; float alphaImag; - uint32_t startOffset[CAXPY_MAX_VECTOR_CORES]; // per-block start, in floats - uint32_t calNum[CAXPY_MAX_VECTOR_CORES]; // per-block element count, in floats + uint32_t startOffset[CAXPY_MAX_VECTOR_CORES]; // per-block start, in floats + uint32_t calNum[CAXPY_MAX_VECTOR_CORES]; // per-block element count, in floats uint32_t totalN; int64_t incx; int64_t incy; diff --git a/blas/common/helper/aclblas_auxiliary.cpp b/blas/common/helper/aclblas_auxiliary.cpp index 636e01c..ff0084e 100644 --- a/blas/common/helper/aclblas_auxiliary.cpp +++ b/blas/common/helper/aclblas_auxiliary.cpp @@ -57,8 +57,7 @@ aclblasStatus_t aclblasDestroy(aclblasHandle_t handle) const aclblasStatus_t syncStatus = SynchronizeHandleStream(h); if (syncStatus != ACLBLAS_STATUS_SUCCESS) { - OP_LOGE("aclblasDestroy", - "stream synchronization failed before handle destruction."); + OP_LOGE("aclblasDestroy", "stream synchronization failed before handle destruction."); return ACLBLAS_STATUS_EXECUTION_FAILED; } diff --git a/blas/common/helper/aclblas_handle_internal.h b/blas/common/helper/aclblas_handle_internal.h index f11b59d..758fac0 100644 --- a/blas/common/helper/aclblas_handle_internal.h +++ b/blas/common/helper/aclblas_handle_internal.h @@ -94,10 +94,14 @@ inline size_t GetEffectiveWorkspaceSize(const _aclblas_handle* h) inline bool CheckEffectiveWorkspaceSize(const _aclblas_handle* h, size_t workSize) { size_t availableBytes = GetEffectiveWorkspaceSize(h); - OP_CHECK_IF(availableBytes < workSize, OP_LOGE("aclblasHandle", - "workspace required %zu bytes, but only %zu bytes available. " - "Please call aclblasSetWorkspace with size >= %zu bytes", - workSize, availableBytes, workSize), return false); + OP_CHECK_IF( + availableBytes < workSize, + OP_LOGE( + "aclblasHandle", + "workspace required %zu bytes, but only %zu bytes available. " + "Please call aclblasSetWorkspace with size >= %zu bytes", + workSize, availableBytes, workSize), + return false); return true; } @@ -157,8 +161,7 @@ inline aclblasStatus_t ResetToDefaultWorkspace(_aclblas_handle* h) h->workspace_size = 0; h->workspace_owner = AclblasWorkspaceOwner::Library; - const size_t allocSize = - h->library_workspace_size > 0 ? h->library_workspace_size : ACLBLAS_DEFAULT_WORKSPACE_SIZE; + const size_t allocSize = h->library_workspace_size > 0 ? h->library_workspace_size : ACLBLAS_DEFAULT_WORKSPACE_SIZE; return AllocateLibraryWorkspace(h, allocSize); } @@ -175,9 +178,9 @@ inline aclblasStatus_t EnsureDefaultWorkspace(_aclblas_handle* h, size_t require } if (h->workspace_owner == AclblasWorkspaceOwner::User) { if (requiredSize > h->workspace_size) { - OP_LOGE("aclblasHandle", - "user workspace too small: required=%zu, available=%zu", - requiredSize, h->workspace_size); + OP_LOGE( + "aclblasHandle", "user workspace too small: required=%zu, available=%zu", requiredSize, + h->workspace_size); return ACLBLAS_STATUS_ALLOC_FAILED; } return ACLBLAS_STATUS_SUCCESS; @@ -188,9 +191,9 @@ inline aclblasStatus_t EnsureDefaultWorkspace(_aclblas_handle* h, size_t require } if (requiredSize > ACLBLAS_MAX_WORKSPACE_SIZE) { - OP_LOGE("aclblasHandle", - "workspace required %zu bytes exceeds maximum limit %zu bytes", - requiredSize, ACLBLAS_MAX_WORKSPACE_SIZE); + OP_LOGE( + "aclblasHandle", "workspace required %zu bytes exceeds maximum limit %zu bytes", requiredSize, + ACLBLAS_MAX_WORKSPACE_SIZE); return ACLBLAS_STATUS_ALLOC_FAILED; } @@ -200,12 +203,11 @@ inline aclblasStatus_t EnsureDefaultWorkspace(_aclblas_handle* h, size_t require } const size_t doubledSize = h->workspace_size > 0 ? h->workspace_size * 2 : ACLBLAS_DEFAULT_WORKSPACE_SIZE; - const size_t newSize = - std::min(std::max(requiredSize, doubledSize), ACLBLAS_MAX_WORKSPACE_SIZE); + const size_t newSize = std::min(std::max(requiredSize, doubledSize), ACLBLAS_MAX_WORKSPACE_SIZE); - OP_LOGW("aclblasHandle", - "library workspace (%zu bytes) insufficient, expanding to %zu bytes", - h->workspace_size, newSize); + OP_LOGW( + "aclblasHandle", "library workspace (%zu bytes) insufficient, expanding to %zu bytes", h->workspace_size, + newSize); FreeLibraryWorkspace(h); diff --git a/test/axpy/caxpy/arch22/caxpy_dispatch_policy_test.cpp b/test/axpy/caxpy/arch22/caxpy_dispatch_policy_test.cpp index ea701d0..01d7bc1 100644 --- a/test/axpy/caxpy/arch22/caxpy_dispatch_policy_test.cpp +++ b/test/axpy/caxpy/arch22/caxpy_dispatch_policy_test.cpp @@ -30,7 +30,6 @@ struct PolicyCase { CaxpyKernelVariant expected; }; - constexpr std::array CASES{{ {1, 1, CaxpyStrideClass::DENSE}, {2, 1, CaxpyStrideClass::X_ONLY_POSITIVE}, @@ -90,7 +89,6 @@ constexpr std::array POLICY_CASES{{ {640, 5, 5, false, 40, CaxpyKernelVariant::DENSE_PIPELINED_ATOMIC}, }}; - } // namespace int main() @@ -101,8 +99,8 @@ int main() if (actual == test.expected) continue; std::cerr << "classification mismatch: incx=" << test.incx << " incy=" << test.incy - << " expected=" << static_cast(test.expected) - << " actual=" << static_cast(actual) << '\n'; + << " expected=" << static_cast(test.expected) << " actual=" << static_cast(actual) + << '\n'; passed = false; } for (const PolicyCase& test : POLICY_CASES) { @@ -111,8 +109,7 @@ int main() if (actual == test.expected) continue; std::cerr << "policy mismatch: n=" << test.n << " incx=" << test.incx << " incy=" << test.incy - << " direct_strided=" << test.directStrided - << " expected=" << static_cast(test.expected) + << " direct_strided=" << test.directStrided << " expected=" << static_cast(test.expected) << " actual=" << static_cast(actual) << '\n'; passed = false; } diff --git a/test/axpy/caxpy/arch22/caxpy_npu_wrapper.h b/test/axpy/caxpy/arch22/caxpy_npu_wrapper.h index 7f49230..f952f89 100644 --- a/test/axpy/caxpy/arch22/caxpy_npu_wrapper.h +++ b/test/axpy/caxpy/arch22/caxpy_npu_wrapper.h @@ -17,17 +17,16 @@ #include "cann_ops_blas.h" #include "fill.h" -static inline bool CaxpyNeedPassThrough(aclblasHandle_t handle, int n) -{ - return handle == nullptr || n <= 0; -} +static inline bool CaxpyNeedPassThrough(aclblasHandle_t handle, int n) { return handle == nullptr || n <= 0; } static inline aclError CaxpyAllocCopyH2D(void*& dPtr, const void* hPtr, size_t bytes) { dPtr = nullptr; - if (hPtr == nullptr || bytes == 0) return ACL_SUCCESS; + if (hPtr == nullptr || bytes == 0) + return ACL_SUCCESS; aclError ret = aclrtMalloc(&dPtr, bytes, ACL_MEM_MALLOC_HUGE_FIRST); - if (ret != ACL_SUCCESS) return ret; + if (ret != ACL_SUCCESS) + return ret; ret = aclrtMemcpy(dPtr, bytes, hPtr, bytes, ACL_MEMCPY_HOST_TO_DEVICE); if (ret != ACL_SUCCESS) { aclrtFree(dPtr); @@ -38,22 +37,25 @@ static inline aclError CaxpyAllocCopyH2D(void*& dPtr, const void* hPtr, size_t b static inline void CaxpyFreeAll(void* dX, void* dY) { - if (dX) aclrtFree(dX); - if (dY) aclrtFree(dY); + if (dX) + aclrtFree(dX); + if (dY) + aclrtFree(dY); } // Physical storage span (in complex elements) covered by a length-n vector // with the given (possibly negative) increment. static inline size_t CaxpySpanElements(int n, int increment) { - if (n <= 0) return 0; + if (n <= 0) + return 0; const int64_t absIncrement = increment >= 0 ? static_cast(increment) : -static_cast(increment); return static_cast((static_cast(n) - 1) * absIncrement + 1); } inline aclblasStatus_t aclblasCaxpy_npu( - aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx, - aclblasComplex* y, int incy) + aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx, aclblasComplex* y, + int incy) { if (CaxpyNeedPassThrough(handle, n)) { return aclblasCaxpy(handle, n, alpha, x, incx, y, incy); diff --git a/test/axpy/caxpy/arch22/caxpy_test.cpp b/test/axpy/caxpy/arch22/caxpy_test.cpp index 4903259..7433d59 100644 --- a/test/axpy/caxpy/arch22/caxpy_test.cpp +++ b/test/axpy/caxpy/arch22/caxpy_test.cpp @@ -49,19 +49,21 @@ std::vector MakeGuardedVector(int n, int increment, int32_t salt } // namespace // ── Test fixture ────────────────────────────────────────────────────────────── -class CaxpyArch22Test : public BlasTest { }; +class CaxpyArch22Test : public BlasTest {}; // ── TEST_F: null handle / null alpha (not expressible via CSV) ─────────────── // aclblasCaxpy's ValidateCaxpyArguments treats handle==nullptr and // alpha==nullptr identically: ACLBLAS_STATUS_INVALID_VALUE (unlike some other // operators that return ACLBLAS_STATUS_HANDLE_IS_NULLPTR for a null handle). -TEST_F(CaxpyArch22Test, NullHandle) { +TEST_F(CaxpyArch22Test, NullHandle) +{ aclblasComplex alpha{1.0F, 0.0F}; aclblasStatus_t ret = aclblasCaxpy_npu(nullptr, 1, &alpha, nullptr, 1, nullptr, 1); EXPECT_EQ(static_cast(ret), static_cast(ACLBLAS_STATUS_INVALID_VALUE)); } -TEST_F(CaxpyArch22Test, NullAlpha) { +TEST_F(CaxpyArch22Test, NullAlpha) +{ aclblasComplex x{1.0F, 0.0F}; aclblasComplex y{1.0F, 0.0F}; aclblasStatus_t ret = aclblasCaxpy_npu(CaxpyArch22Test::handle_, 1, nullptr, &x, 1, &y, 1); @@ -70,12 +72,12 @@ TEST_F(CaxpyArch22Test, NullAlpha) { // ── CSV parameterised test suite ───────────────────────────────────────────── INSTANTIATE_TEST_SUITE_P( - Caxpy, CaxpyArch22Test, - ::testing::ValuesIn(GetCasesFromCsv(ReplaceFileExtension2Csv(__FILE__))), + Caxpy, CaxpyArch22Test, ::testing::ValuesIn(GetCasesFromCsv(ReplaceFileExtension2Csv(__FILE__))), PrintCaseInfoString); // ── TEST_P: 5-step CSV-driven flow ─────────────────────────────────────────── -TEST_P(CaxpyArch22Test, CsvDriven) { +TEST_P(CaxpyArch22Test, CsvDriven) +{ const auto& p = GetParam(); const aclblasComplex alpha{p.alphaReal, p.alphaImag}; @@ -95,16 +97,21 @@ TEST_P(CaxpyArch22Test, CsvDriven) { // Step 3: Verify expected return code EXPECT_EQ(static_cast(ret), static_cast(p.expectResult)); - if (p.expectResult != ACLBLAS_STATUS_SUCCESS) return; - if (p.n == 0) return; + if (p.expectResult != ACLBLAS_STATUS_SUCCESS) + return; + if (p.n == 0) + return; // Step 4: Compute golden on CPU over the same guarded storage, so any // stray write into the guard region is caught by the full-buffer compare. - std::vector xGolden = p.nullx == 0 ? MakeGuardedVector(p.n, p.incx, 3) : std::vector{}; - std::vector yGolden = p.nully == 0 ? MakeGuardedVector(p.n, p.incy, -5) : std::vector{}; + std::vector xGolden = + p.nullx == 0 ? MakeGuardedVector(p.n, p.incx, 3) : std::vector{}; + std::vector yGolden = + p.nully == 0 ? MakeGuardedVector(p.n, p.incy, -5) : std::vector{}; const aclblasComplex* xGoldenPtr = xGolden.empty() ? nullptr : xGolden.data() + GUARD_COUNT; aclblasComplex* yGoldenPtr = yGolden.empty() ? nullptr : yGolden.data() + GUARD_COUNT; - aclblasStatus_t cpuRet = aclblasCaxpy_cpu(CaxpyArch22Test::handle_, p.n, &alpha, xGoldenPtr, p.incx, yGoldenPtr, p.incy); + aclblasStatus_t cpuRet = + aclblasCaxpy_cpu(CaxpyArch22Test::handle_, p.n, &alpha, xGoldenPtr, p.incx, yGoldenPtr, p.incy); EXPECT_EQ(static_cast(cpuRet), static_cast(ACLBLAS_STATUS_SUCCESS)); // Step 5: Precision verification, real/imag components split (MERE/MARE). diff --git a/test/axpy/caxpy/caxpy_golden.h b/test/axpy/caxpy/caxpy_golden.h index fd588c8..85cc47f 100644 --- a/test/axpy/caxpy/caxpy_golden.h +++ b/test/axpy/caxpy/caxpy_golden.h @@ -40,17 +40,16 @@ inline aclblasComplex CaxpyComplexMac(aclblasComplex alpha, aclblasComplex x, ac // convention: the last logical element is physically first. inline int64_t CaxpyStorageIndex(int64_t logicalIndex, int64_t n, int64_t increment) { - const uint64_t absIncrement = - increment > 0 ? static_cast(increment) : static_cast(-increment); - return increment > 0 ? logicalIndex * static_cast(absIncrement) - : (n - 1 - logicalIndex) * static_cast(absIncrement); + const uint64_t absIncrement = increment > 0 ? static_cast(increment) : static_cast(-increment); + return increment > 0 ? logicalIndex * static_cast(absIncrement) : + (n - 1 - logicalIndex) * static_cast(absIncrement); } // y = alpha * x + y, complex, arbitrary (possibly negative) incx/incy. // x and y are the caller-owned base storage pointers (not offset by any guard region). inline aclblasStatus_t aclblasCaxpy_cpu( - aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx, - aclblasComplex* y, int incy) + aclblasHandle_t handle, int n, const aclblasComplex* alpha, const aclblasComplex* x, int incx, aclblasComplex* y, + int incy) { aclblasStatus_t st = CaxpyValidateParams(handle, n, alpha, x, incx, y, incy); if (st != ACLBLAS_STATUS_SUCCESS) { diff --git a/test/axpy/caxpy/caxpy_param.h b/test/axpy/caxpy/caxpy_param.h index 9afa9e1..b6f2d66 100644 --- a/test/axpy/caxpy/caxpy_param.h +++ b/test/axpy/caxpy/caxpy_param.h @@ -25,12 +25,12 @@ struct CaxpyParam : public BlasTestParamBase { CaxpyParam(const csv_map& csv) : BlasTestParamBase(csv) { - n = parseInt(ReadMap(csv, "n", "0")); - incx = parseInt(ReadMap(csv, "incx", "1")); - incy = parseInt(ReadMap(csv, "incy", "1")); + n = parseInt(ReadMap(csv, "n", "0")); + incx = parseInt(ReadMap(csv, "incx", "1")); + incy = parseInt(ReadMap(csv, "incy", "1")); alphaReal = parseFloat(ReadMap(csv, "alpha_real", "1.0")); alphaImag = parseFloat(ReadMap(csv, "alpha_imag", "0.0")); - nullx = parseInt(ReadMap(csv, "nullx", "0")); - nully = parseInt(ReadMap(csv, "nully", "0")); + nullx = parseInt(ReadMap(csv, "nullx", "0")); + nully = parseInt(ReadMap(csv, "nully", "0")); } };