CANN ops-math BallQuery 算子技術(shù)指南PointNet 球查詢算子的原理、參數(shù)與源碼實(shí)現(xiàn)解析【免費(fèi)下載鏈接】ops-math本項(xiàng)目是CANN提供的數(shù)學(xué)類基礎(chǔ)計(jì)算算子庫(kù)實(shí)現(xiàn)網(wǎng)絡(luò)在NPU上加速計(jì)算。項(xiàng)目地址: https://gitcode.com/cann/ops-math導(dǎo)讀BallQuery球查詢是 PointNet 系列點(diǎn)云網(wǎng)絡(luò)中的經(jīng)典算子用于在給定半徑范圍內(nèi)為每個(gè)查詢中心點(diǎn)收集鄰域點(diǎn)的索引。本文以 conversion/ball_query/README.md 為主線完整梳理該算子在 CANN ops-math 中的產(chǎn)品支持情況、數(shù)學(xué)定義、參數(shù)約束與調(diào)用方式并深入到算子定義、InferShape、Tiling 與 SIMT Kernel 的實(shí)現(xiàn)源碼幫助讀者既能在實(shí)踐中正確使用該算子也能理解它在 NPU 上的并行計(jì)算原理。一、算子概述BallQuery 在點(diǎn)云網(wǎng)絡(luò)中的角色Ball Query 是 PointNet 中用于局部區(qū)域特征聚合的核心操作。與 KNNK 近鄰相比Ball Query 通過(guò)固定半徑的球形鄰域替代 K 近鄰從而保證鄰域采樣在空間上是各向同性isotropic的不受點(diǎn)云密度分布影響因而在點(diǎn)云分割、分類等任務(wù)中被廣泛采用。在 CANN ops-math 倉(cāng)庫(kù)中該算子位于 conversion/ball_query 目錄屬于 conversion轉(zhuǎn)換/采樣類算子集合。它的基本職責(zé)是對(duì)每個(gè)查詢中心點(diǎn)center_xyz[m,b]在點(diǎn)集xyz[b]中按k 0..N-1的自然順序查找滿足距離判定條件的點(diǎn)并收集前sample_num個(gè)點(diǎn)的索引到輸出idx中。產(chǎn)品支持情況該算子當(dāng)前支持以下產(chǎn)品的 NPU 運(yùn)行環(huán)境產(chǎn)品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 訓(xùn)練系列產(chǎn)品 / Atlas A3 推理系列產(chǎn)品√Atlas A2 訓(xùn)練系列產(chǎn)品 / Atlas A2 推理系列產(chǎn)品√Atlas 200I/500 A2 推理產(chǎn)品×Atlas 推理系列產(chǎn)品√Atlas 訓(xùn)練系列產(chǎn)品√從 算子定義 可以看到源碼中通過(guò)this-AICore().AddConfig(ascend950, aicoreConfig)顯式注冊(cè)了 ascend950 平臺(tái)的 AI Core 配置與文檔中Ascend 950PR/950DT 支持的聲明相互印證同時(shí)算子 Kernel 與 Tiling 均位于arch35目錄op_kernel/arch35對(duì)應(yīng) 950 系列所采用的架構(gòu)。二、功能與計(jì)算公式2.1 數(shù)據(jù)布局約定BallQuery 的輸入輸出采用如下非對(duì)稱的 3D 布局xyzshape 為(B, 3, N)坐標(biāo)位于中間維。其中B為 batch 數(shù)N為每個(gè) batch 內(nèi)的點(diǎn)云點(diǎn)數(shù)第二維固定為 3 表示 (x, y, z) 坐標(biāo)。center_xyzshape 為(M, B, 3)坐標(biāo)位于最后一維。M即 npoint為每個(gè) batch 內(nèi)的查詢中心點(diǎn)數(shù)量其第 1 維B必須與xyz的第 0 維B一致。idx輸出shape 為(M, B, sample_num)dtype 固定為 INT32。注意兩個(gè)輸入張量的坐標(biāo)維位置不同xyz 坐標(biāo)在中間維center_xyz 坐標(biāo)在末維這是該算子的一個(gè)重要記憶點(diǎn)在編寫推理腳本或構(gòu)圖代碼時(shí)極易混淆。2.2 距離計(jì)算與命中判定對(duì)每個(gè)(m, b)取中心點(diǎn) $(c_x, c_y, c_z) center_xyz[m, b, :]$計(jì)算其與xyz[b]中每個(gè)點(diǎn) $(x_k, y_k, z_k) (xyz[b, 0, k], xyz[b, 1, k], xyz[b, 2, k])$ 的距離平方$$ d2 (c_x - x_k)^2 (c_y - y_k)^2 (c_z - z_k)^2 $$距離判定條件為$$ d2 0 \quad 或 \quad min_radius^2 \le d2 max_radius^2 $$即命中點(diǎn)要么與中心點(diǎn)完全重合距離為 0典型場(chǎng)景為查詢中心點(diǎn)本身取自點(diǎn)云要么落在以中心點(diǎn)為圓心、min_radius為內(nèi)半徑、max_radius為外半徑的球形環(huán)帶內(nèi)。當(dāng)min_radius 0時(shí)即退化為經(jīng)典的實(shí)心球鄰域查詢。2.3 輸出收集與填充規(guī)則命中點(diǎn)按照k的自然順序0 到 N-1依次寫入idx首個(gè)命中點(diǎn)的索引記為first_num當(dāng)命中點(diǎn)數(shù)量不足sample_num時(shí)剩余位置用first_num填充當(dāng)沒(méi)有任何命中點(diǎn)時(shí)剩余位置統(tǒng)一填充 0該行為在 Kernel 實(shí)現(xiàn)中體現(xiàn)見(jiàn)下文源碼分析。三、參數(shù)說(shuō)明參數(shù)名輸入/輸出/屬性描述數(shù)據(jù)類型數(shù)據(jù)格式xyz輸入所有點(diǎn)的 xyz 坐標(biāo)shape 為(B,3,N)中間維必須為 3坐標(biāo)在中間維FLOAT、FLOAT16NDcenter_xyz輸入查詢中心點(diǎn)坐標(biāo)shape 為(M,B,3)最后一維必須為 3第 1 維 B 必須與 xyz 的第 0 維一致與 xyz 保持一致NDidx輸出采樣點(diǎn)在 xyz 中的索引shape 為(M,B,sample_num)INT32NDmin_radius屬性最小查詢半徑環(huán)形內(nèi)邊界必須大于等于 0Float-max_radius屬性最大查詢半徑球外邊界必須大于 min_radiusFloat-sample_num屬性每球最大采樣點(diǎn)數(shù)必須大于等于 1Int-3.1 算子定義層的參數(shù)約束在 算子 IR 定義 中三個(gè)屬性均被聲明為必填REQUIRED_ATTRREG_OP(BallQuery) .INPUT(xyz, TensorType({DT_FLOAT16, DT_FLOAT})) .INPUT(center_xyz, TensorType({DT_FLOAT16, DT_FLOAT})) .OUTPUT(idx, TensorType({DT_INT32})) .REQUIRED_ATTR(min_radius, Float) .REQUIRED_ATTR(max_radius, Float) .REQUIRED_ATTR(sample_num, Int) .OP_END_FACTORY_REG(BallQuery)而在 ball_query_def.cpp 中算子注冊(cè)了如下關(guān)鍵配置輸入xyz、center_xyz均聲明AutoContiguous()即框架會(huì)對(duì)非連續(xù)輸入自動(dòng)做連續(xù)化處理對(duì)應(yīng) README 中非連續(xù) Tensor 自動(dòng)連續(xù)化的說(shuō)明AICore 配置開(kāi)啟DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)、DynamicCompileStaticFlag(true)說(shuō)明該算子支持動(dòng)態(tài) shape 與靜態(tài)編譯相結(jié)合的調(diào)度方式ExtendCfgInfo(opFile.value, ball_query_apt)指定了算子內(nèi)核入口文件為 ball_query_apt.cpp。四、約束說(shuō)明以下是使用 BallQuery 時(shí)必須遵守的約束條件其中多數(shù)約束在 InferShape 與 Tiling 源碼中都有對(duì)應(yīng)的校驗(yàn)邏輯見(jiàn)第五節(jié)xyz為 3 維 Tensorshape 為(B,3,N)中間維必須為 3center_xyz為 3 維 Tensorshape 為(M,B,3)最后一維必須為 3xyz的第 0 維 B 與center_xyz的第 1 維 B 必須一致xyz的第 0 維 B、第 2 維 N以及center_xyz的第 0 維 Mnpoint取值均不得超過(guò) INT32_MAX2147483647xyz與center_xyz的 dtype 必須相同且為 FLOAT 或 FLOAT16輸出idx的 dtype 固定為 INT32shape 為(M,B,sample_num)min_radius必須為有限值非 NaN、非 Inf且大于等于 0max_radius必須為有限值非 NaN、非 Inf且大于 0并必須大于min_radiussample_num必須大于等于 1 且不超過(guò) INT32_MAX輸入 Tensor 需為連續(xù)內(nèi)存布局contiguous即數(shù)據(jù)在內(nèi)存中按行主序緊密排列不含 stride 跳步或間隙概念詳見(jiàn)非連續(xù)的Tensor。若傳入非連續(xù) Tensor如轉(zhuǎn)置、切片得到的視圖框架會(huì)自動(dòng)連續(xù)化處理產(chǎn)生額外拷貝開(kāi)銷算子僅支持 3D 輸入不支持標(biāo)量、1D 或 8D 場(chǎng)景。五、源碼級(jí)實(shí)現(xiàn)解析5.1 InferShape輸出形狀推導(dǎo)與校驗(yàn)ball_query_infershape.cpp 實(shí)現(xiàn)了形狀推導(dǎo)函數(shù)InferShape4BallQuery其核心邏輯與 README 約束一一對(duì)應(yīng)坐標(biāo)維校驗(yàn)對(duì)xyz檢查dimNum 3 dim(1) 3對(duì)center_xyz檢查dimNum 3 dim(2) 3。當(dāng)輸入為 unknownRank動(dòng)態(tài)秩時(shí)會(huì)跳過(guò)維數(shù)校驗(yàn)Batch 一致性校驗(yàn)在非 unknownRank 場(chǎng)景下要求xyz.dim(0) center_xyz.dim(1)若任一維為動(dòng)態(tài)值 -1 則跳過(guò)等值校驗(yàn)dtype 校驗(yàn)要求xyz與center_xyz的 dtype 相同且取值必須為DT_FLOAT或DT_FLOAT16屬性校驗(yàn)min_radius必須isfinite且 0max_radius必須isfinite且 0且 min_radiussample_num必須 1輸出推導(dǎo)idx.shape (center_xyz.dim(0), center_xyz.dim(1), sample_num)即(M, B, sample_num)。若center_xyz為 unknownRank則 M、B 位置寫 -1rank 與 sample_num 仍已知。倉(cāng)庫(kù)中的 InferShape 單元測(cè)試 覆蓋了 FP32 靜態(tài) shapexyz(5,3,15)、center_xyz(20,5,3)、sample_num4期望輸出idx(20,5,4)與 FP16 動(dòng)態(tài) shape期望輸出idx(-1,2,16)等典型場(chǎng)景可作為理解輸出推導(dǎo)規(guī)則的參考用例。5.2 Tiling多核切分策略ball_query_tiling.cpp 實(shí)現(xiàn)了 host 側(cè)的切分邏輯核心步驟包括平臺(tái)信息獲取通過(guò)platform_ascendc::PlatformAscendC獲取 AIV 核數(shù)coreNum與 UB 內(nèi)存大小輸入與屬性校驗(yàn)與 InferShape 相同規(guī)格的 shape/dtype/屬性合法性檢查并將 B、N、npoint 收窄為 int32核數(shù)切分以查詢點(diǎn)總數(shù)為并行單元totalQueryPoints B * npointperCoreElements ceil(totalQueryPoints / coreNum)并設(shè)置每核最小處理量PER_CORE_MIN 64同時(shí)向上對(duì)齊到WARP_SIZE 32最終needCoreNum ceil(totalQueryPoints / perCoreElements)通過(guò)context-SetBlockDim()下發(fā)TilingData 填充向 BallQueryTilingData 寫入B、N、npoint、sampleNum以及預(yù)計(jì)算的半徑平方minRadius2 min_radius2、maxRadius2 max_radius2。將半徑平方預(yù)計(jì)算放到 host 側(cè)可避免 kernel 端重復(fù)開(kāi)方與乘法運(yùn)算TilingKey 設(shè)置通過(guò) ball_query_tiling_key.h 中的BALL_QUERY_SCH_MODE_DEFAULT宏選擇默認(rèn)調(diào)度模式并聲明為ASCENDC_TPL_AIV_ONLY僅使用 AIV 向量/標(biāo)量核。5.3 SIMT Kernel并行查詢與索引收集ball_query_simt.h 是算子的核心 SIMT 內(nèi)核實(shí)現(xiàn)采用Grid-Stride 循環(huán)每個(gè)線程以stride blockDim.x * gridDim.x為間隔處理多個(gè)查詢點(diǎn)適合查詢點(diǎn)數(shù)遠(yuǎn)大于線程數(shù)的場(chǎng)景。其處理流程為索引分解由一維查詢點(diǎn)索引qIdx反解(m, b)qIdx m*B b其中除法采用Simt::UintDiv的 magic/shift 快速除法由 Process 階段通過(guò)GetUintDivMagicAndShift預(yù)計(jì)算GM 基址計(jì)算分別定位xyz[b, 0, 0]、center_xyz[m, b, :]、idx[m, b, :]三個(gè)地址坐標(biāo)讀取與精度提升讀取中心點(diǎn)坐標(biāo)后通過(guò)CoordToFloat統(tǒng)一提升為 float32——對(duì) half 類型調(diào)用__half2float避免 FP16 計(jì)算距離平方時(shí)的溢出與截?cái)嗑嚯x判定與收集遍歷該 batch 內(nèi)全部 N 個(gè)點(diǎn)按d2 0.0f || (d2 minRadius2 d2 maxRadius2)判定命中命中索引按序?qū)懭雐dxBase[cnt]并記錄首個(gè)命中點(diǎn)firstResult當(dāng)cnt達(dá)到sampleNum時(shí)提前終止內(nèi)層循環(huán)填充與未命中處理命中數(shù)不足時(shí)用firstResult填充剩余槽位完全未命中時(shí)將所有槽位置 0。這與 README 中命中數(shù)不足 sample_num 時(shí)剩余位置用 first_num 填充的語(yǔ)義一致并補(bǔ)充了零命中時(shí)填 0的實(shí)現(xiàn)細(xì)節(jié)。Kernel 入口 ball_query_apt.cpp 通過(guò)REGISTER_TILING_DEFAULT/GET_TILING_DATA_WITH_STRUCT讀取 TilingData并按schMode BALL_QUERY_SCH_MODE_DEFAULT分派到NsBallQuery::Process。六、調(diào)用方式GE 圖模式6.1 調(diào)用總覽調(diào)用方式調(diào)用樣例說(shuō)明GE 圖模式test_geir_ball_query通過(guò)算子IR構(gòu)圖方式調(diào)用 BallQuery 算子參見(jiàn)算子調(diào)用完成算子編譯和驗(yàn)證。當(dāng)前倉(cāng)庫(kù)為 BallQuery 提供的官方調(diào)用樣例是GEGraph Engine圖模式即通過(guò)ge::op::BallQuery(ball1)在計(jì)算圖中創(chuàng)建算子節(jié)點(diǎn)再經(jīng)AddGraphRunGraph完成構(gòu)圖、編譯與執(zhí)行。6.2 構(gòu)圖與輸入輸出聲明示例程序 test_geir_ball_query.cpp 中關(guān)鍵構(gòu)圖邏輯位于CreateOppInGraphauto ball1 op::BallQuery(ball1); // 輸入xyz(B,3,N) 與 center_xyz(M,B,3) ADD_INPUT_MODE(1, xyz, inDtype, xyzShape, xyzDShape, mode); ADD_INPUT_MODE(2, center_xyz, inDtype, centerXyzShape, centerXyzDShape, mode); // 輸出idx(M,B,sample_num)dtype 固定 INT32 std::vectorint64_t idxShape {centerXyzShape[0], centerXyzShape[1], sampleNum}; ADD_OUTPUT_MODE(1, idx, DT_INT32, idxShape, mode); // 屬性設(shè)置min_radius - max_radius - sample_num ball1.set_attr_min_radius(minRadius); ball1.set_attr_max_radius(maxRadius); ball1.set_attr_sample_num(sampleNum);該示例同時(shí)演示了靜態(tài)S與動(dòng)態(tài)D兩種構(gòu)圖模式。特別值得注意的是 D 模式下動(dòng)態(tài) shape 的設(shè)置方式由于 InferShape 會(huì)校驗(yàn)坐標(biāo)維因此不能把所有維都置為 -1而應(yīng)僅將動(dòng)態(tài)維置 -1、保留坐標(biāo)維為已知值 3即xyz使用 D shape{-1, 3, -1}center_xyz使用 D shape{-1, -1, 3}。這一細(xì)節(jié)體現(xiàn)了 BallQuery 坐標(biāo)維布局約束對(duì)動(dòng)態(tài) shape 構(gòu)圖的影響是實(shí)際開(kāi)發(fā)中容易踩坑的點(diǎn)。6.3 執(zhí)行與驗(yàn)證矩陣示例程序的測(cè)試矩陣體現(xiàn)了算子的完整適用邊界dtype 矩陣DT_FLOATFP32與DT_FLOAT16FP16兩類輸入輸出固定 INT32shape 場(chǎng)景矩陣共 5 類——常規(guī) 3Dxyz(2,3,1024)、center_xyz(512,2,3)、sample_num16、小規(guī)模 3D、空點(diǎn)云N0、空查詢M0、多 batchB4運(yùn)行模式每個(gè) (dtype, shape) 組合都同時(shí)跑 S靜態(tài)與 D動(dòng)態(tài)兩種模式最終輸出逐 case 的 Build / RunGraph / OutputExists 報(bào)告與通過(guò)率統(tǒng)計(jì)。程序啟動(dòng)時(shí)會(huì)通過(guò)全局選項(xiàng)初始化 GE 引擎包括ge.exec.deviceId需通過(guò)npu-smi info查詢狀態(tài)為 OK 的設(shè)備號(hào)、ge.graphRunMode以及ge.exec.precision_mode must_keep_origin_dtype保持原始 dtype 精度。上述 shape 組合可直接作為讀者自行編寫驗(yàn)證腳本的參考模板。七、總結(jié)與使用建議BallQuery 是 CANN ops-math 中面向點(diǎn)云處理的鄰域采樣算子其核心設(shè)計(jì)可以概括為四點(diǎn)布局特殊xyz坐標(biāo)在中間維、center_xyz坐標(biāo)在末維構(gòu)圖與推理時(shí)必須嚴(yán)格區(qū)分判定規(guī)則d2 0或min_radius2 ≤ d2 max_radius2才計(jì)入采樣天然支持環(huán)形空心球鄰域半徑平方預(yù)計(jì)算host 側(cè) Tiling 階段完成min_radius2、max_radius2計(jì)算Kernel 僅做乘加與比較降低片上計(jì)算開(kāi)銷并行充分以B × npoint個(gè)查詢點(diǎn)為并行單元Grid-Stride 循環(huán)配合 magic-division 索引反解適配 AIV 多核場(chǎng)景。實(shí)際使用時(shí)建議先對(duì)照約束說(shuō)明核對(duì) shape/dtype/屬性取值再按算子調(diào)用指南編譯運(yùn)行示例 test_geir_ball_query.cpp若需深度定制可進(jìn)一步閱讀 Tiling 實(shí)現(xiàn) 與 SIMT Kernel 了解切分與并行細(xì)節(jié)或參考 InferShape UT 中的用例組織方式編寫自己的驗(yàn)證用例?！久赓M(fèi)下載鏈接】ops-math本項(xiàng)目是CANN提供的數(shù)學(xué)類基礎(chǔ)計(jì)算算子庫(kù)實(shí)現(xiàn)網(wǎng)絡(luò)在NPU上加速計(jì)算。項(xiàng)目地址: https://gitcode.com/cann/ops-math創(chuàng)作聲明：本文部分內(nèi)容由AI輔助生成（AIGC），僅供參考