HashTable MTE Queue Optimization Sample for SIMD and SIMT Hybrid Programming【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkitOverviewThis sample compares two value write paths for HashTableinsert_or_assign: direct value writes by a SIMT Warp, and value writes by an MTE task queue in a single kernel. The sample implements a fixed-capacity C-style HashTable. It shows this optimization path and keeps only the semantics and data structures required for the comparison.Optimization paths:CaseSCENARIO_NUMImplementationDescriptionCase 00Direct value writes by a SIMT WarpAfter a Warp completes hash, probe, and CAS operations, the threads in the Warp write one value vector in shards. This case is the baseline implementation.Case 11Value movement by an MTE task queueA task-id ring queue is used in a single kernel. SIMT continuously produces UB tasks, and MTE moves values in batches by task. The MTE movement forms a pipeline with subsequent SIMT probing.The following figure shows the data processing pipelines of the two cases.Supported Products and CANN Software VersionProductCANN Software VersionAscend 950PR/Ascend 950DT CANN 9.1.0Directory Structure├── simd_simt_hash_table_mte_queue │ ├── CMakeLists.txt // Compilation project file │ ├── figures // Image resources in the README │ ├── hash_table_mte_queue.asc // Ascend C sample implementation │ ├── README.md // Sample description document │ └── scripts │ └── run.sh // Compiles and uses msOpProf to collect performance data of the two casesSample DescriptionSample function:This sample implements a fixed-capacity HashTable. It supports insert-or-update semantics similar to C17std::unordered_map::insert_or_assign: when a key exists, the sample updates the value; when a key does not exist, the sample inserts the key and value. The sample handles hash collisions by linear probing.The HashTable consists of two GM storage segments.table[bucket_count]stores bucket metadata. Each bucket stores only a key and a state.table_values[bucket_count, dim]stores value data. Rowbucket_idxstores the value vector of the key that corresponds totable[bucket_idx]. The key and value are not stored contiguously in the same structure. They are associated by the samebucket_idx. Each value vector has a length ofdim, anddimis a runtime parameter.The key field of a bucket also stores occupation information:Key field valueMeaningEMPTY_KEYThe bucket is empty and can be locked by CAS in the insertion flow.LOCKED_KEYThe bucket is being inserted or updated. Other threads must wait until the key is published and then check the key again.Real keyThe bucket is published and visible to find operations.During insertion or update, the kernel first changes the key of the target bucket fromEMPTY_KEYor a real key toLOCKED_KEYby CAS. Then it writes the value and publishes the real key. The find operation matches only real keys, so it cannot observe an intermediate state in which the key is visible but the value is not completely written.The HashTable insertion flow is as follows:$$ bucket hash(key) (capacity - 1) $$$$ bucket (bucket 1) (capacity - 1), \quad \text{when collision occurs} $$Sample specifications:Sample Type (OpType)HashTable insert_or_assignSample Inputnameshapedata typeformatkeys[128 * 1024]int64NDvalues[128 * 1024, dim]float32NDSample Outputtable keys[256 * 1024]int64NDtable values[256 * 1024, dim]float32NDKernel NameSCENARIO_NUM0insert_or_assign_warp_store_kernelSCENARIO_NUM1insert_or_assign_mte_task_queue_kernelThe sample parameters are as follows:ParameterValueHASH_TABLE_CAPACITY256 * 1024KEY_NUM128 * 1024Load factor0.5TEST_DIMS16, 32, 64, 128, 256BLOCK_DIM64THREAD_COUNT512WARP_SIZE32PROFILE_REPEAT_TIMESThe default value is 30.run.shsets it to 1 when collecting performance data.Sample ImplementationThe performance comparison usesTask Duration(μs)in the msOpProf output.Case 0: Direct Value Writes by a SIMT WarpImplementation: For the implementation, refer to theinsert_or_assign_warp_store_vf()function.In this implementation, each Warp processes one key. Lane 0 performs hash, linear probing, and CAS bucket locking. Threads in the Warp obtain the target bucket index byasc_shfland write the value vector in shards byWARP_SIZE. SIMT threads directly access GM to complete value movement.The write sequence is as follows: lock the target bucket by CAS, write the value by threads in the Warp, runasc_threadfence(), and publish the real key from lane 0.Key code:__simt_vf__ inline void insert_or_assign_warp_store_vf(__gm__ Bucket* table, __gm__ float* table_values, __gm__ const int64_t* keys, __gm__ const float* values, uint32_t key_num, uint32_t capacity, uint32_t dim, uint32_t block_index, uint32_t num_blocks) { const uint32_t lane_id static_castuint32_t(threadIdx.x) % WARP_SIZE; ... if (lane_id 0) { int64_t old_key asc_atomic_cas(key_addr, EMPTY_KEY, LOCKED_KEY); if (old_key LOCKED_KEY) { old_key wait_until_unlocked(key_addr); } ... } ... for (uint32_t j lane_id; j dim; j WARP_SIZE) { table_values[dst_base j] values[src_base j]; } asc_threadfence(); if (lane_id 0) { (void)asc_atomic_exch(table[bucket_idx].key, key); } }Performance data:CaseImplementationdimNumber of CoresTask Duration(μs)0Direct value writes by a SIMT Warp1664218.7240Direct value writes by a SIMT Warp3264209.4060Direct value writes by a SIMT Warp6464259.1070Direct value writes by a SIMT Warp12864364.8050Direct value writes by a SIMT Warp25664555.497Performance data analysis:Asdimincreases, the time consumed by element-by-element value writes from SIMT threads increases significantly. Whendim256, value movement becomes the main overhead of this path.Principle:HashTable insertion includes discrete control flow such as hash, linear probing, CAS, and collision handling. SIMT is suitable for this control flow. A value vector is stored in contiguous memory. When SIMT threads access GM element by element, they are usually less efficient than MTE for contiguous movement.Optimization direction:In Case 0, the same group of SIMT threads handles both control flow and value writes. Hash, probe, and CAS operations are discrete control flow and are suitable for SIMT threads. However, value vectors are stored at contiguous GM addresses. Asdimincreases, threads in a Warp must perform more element-by-element GM writes. The SIMT threads are occupied by value movement, and probing for subsequent keys must wait until the current value is written.Case 1 keeps SIMT responsible for hash, probe, and CAS operations, and moves only contiguous value movement to MTE. After SIMT threads lock a bucket, they write the source address, destination address, and key to publish to a UB task queue. Then they continue to produce subsequent tasks. MTE performs GM-UB-GM movement in batches by task queue. In this way, SIMT probing for subsequent keys can run in parallel with previous value movement in a single kernel. This design reduces the impact of contiguous value writes occupying SIMT threads in Case 0.Case 1: Value Movement by an MTE Task QueueImplementation: For the implementation, refer to themte_task_queueclass and theinsert_or_assign_mte_task_queue_kernel()function.In this implementation, SIMT threads still perform hash, probe, and CAS operations, but they no longer write the complete task list to GM. Each Vector Core maintains a UB task-id ring queue. The SIMT producer writes successfully locked tasks to the ring. Each task records the input value source address, the HashTable value destination address, the key to publish, and the key address. The MTE consumer moves values in batches by tasks in the ring. After the value writes are complete, the SIMT drain path publishes the real key.global_id,simt_assign_id,mte_finish_id,key_assign_id, andkey_finish_idtrack task allocation, the MTE-visible range, value movement progress, key publishing claim range, and final publishing progress respectively. The bucket key remainsLOCKED_KEYbefore the value is written. The real key is published only after the complete value vector is written back to GM.The following figure shows the visibility timing of a single task in Case 1. The SIMT path advancesglobal_id,simt_assign_id,key_assign_id, andkey_finish_id. The MTE path advancesmte_finish_id.Key code:class mte_task_queue { public: __simt_callee__ inline void assign_warp_task( uint64_t src_addr, uint64_t dst_addr, int64_t key, uint64_t key_addr, bool need_assign) { uint32_t write_mask asc_ballot(need_assign ? 1 : 0); uint32_t warp_write_count __popc(write_mask); uint32_t task_start_id (laneid() 0) ? apply_id(warp_write_count) : 0; task_start_id asc_shfl(task_start_id, 0); drain(task_start_id warp_write_count); if (need_assign) { uint32_t lane_offset __popc(write_mask lanemask_lt()); fill_task_info(task_start_id lane_offset, src_addr, dst_addr, key, key_addr); } asc_threadfence(); if (laneid() 0) { assign_task(task_start_id, task_start_id warp_write_count); } } __aicore__ inline void run_task() const { while (true) { uint32_t simt_assign_id *reinterpret_cast__ubuf__ volatile uint32_t*(simt_assign_id_); uint32_t proc_count min(max_proc_batch, simt_assign_id - mte_finish_id); ... asc_copy_gm2ub_align(value_local_, src_addr, 1, value_bytes, 0, 0, true, values_cache_mode_, 0, 0); asc_sync_notify(PIPE_MTE2, PIPE_MTE3, EVENT_ID0); asc_sync_wait(PIPE_MTE2, PIPE_MTE3, EVENT_ID0); asc_copy_ub2gm_align(dst_addr, value_local_, 1, value_bytes, table_values_cache_mode, 0, 0); asc_sync_notify(PIPE_MTE3, PIPE_S, EVENT_ID0); asc_sync_wait(PIPE_MTE3, PIPE_S, EVENT_ID0); *reinterpret_cast__ubuf__ volatile uint32_t*(mte_finish_id_) mte_finish_id proc_count; } } };Performance data:CaseImplementationdimNumber of CoresTask Duration(μs)1Value movement by an MTE task queue1664129.1491Value movement by an MTE task queue3264112.5261Value movement by an MTE task queue6464111.7611Value movement by an MTE task queue12864130.7601Value movement by an MTE task queue25664201.567Performance data analysis:Case 1 moves the value vector with a length ofdimout of SIMT threads. SIMT producer and MTE consumer run in parallel in a single kernel.Compared with Case 0, Case 1 accelerates all fivedimconfigurations. The highest speedup is about 2.79x whendim128.Performance Comparison SummaryPerformance Comparison on Ascend 950PRThe performance data of eachdimconfiguration is as follows:dimCase 0 Direct value writes by a SIMT Warp (μs)Case 1 MTE task queue (μs)Case 1 Speedup16218.724129.1491.69x32209.406112.5261.86x64259.107111.7612.32x128364.805130.7602.79x256555.497201.5672.76xOptimization SummaryOptimizationCore PrincipleSample ImplementationReplace SIMT element-by-element value writes with MTEMove value vector writes with a length ofdimout of SIMT threads. MTE performs GM-UB-GM batch movement to reduce the overhead of direct GM access by SIMT threads.In Case 0, Warp threads write values in shards. In Case 1, MTE consumes tasks and completes value movement.Single-kernel task queue pipeline parallelismUse a UB task-id ring queue in a single kernel to connect SIMT probing and MTE movement. This lets probing for subsequent keys run in parallel with previous value movement.Case 1 usesmte_task_queueto organize a producer and consumer pipeline. SIMT probing and MTE movement overlap in the same kernel.Compilation and RunningIn the sample root directory, perform the following steps to compile and run the sample.Switch the case.Specify the case to compile by-DSCENARIO_NUMNduring CMake compilation. The cases are as follows:0: Direct value writes by a SIMT Warp1: Value movement by an MTE task queueCommand:cmake .. -DCMAKE_ASC_ARCHITECTURESdav-3510 -DSCENARIO_NUM1Configure environment variables.Configure environment variables based on the installation method of the CANN development kit package in the current environment.source ${install_path}/cann/set_env.shDescription:${install_path}is the installation directory of the CANN package. If no installation directory is specified, the default installation directory is/usr/local/Ascend.Run the sample.Run the following commands in the sample directory.SCENARIO_NUM1 # Select the execution scenario. mkdir -p build cd build cmake -DSCENARIO_NUM$SCENARIO_NUM -DCMAKE_ASC_ARCHITECTURESdav-3510 .. make -j ./hash_table_mte_queue 128You can also directly run the comparison script. The script compiles the two cases and collects performance data for fivedimconfigurations:16/32/64/128/256../scripts/run.shCompilation optionsOptionValueDescriptionCMAKE_ASC_ARCHITECTURESdav-3510NPU architecture, which corresponds to Ascend 950PR/Ascend 950DT.SCENARIO_NUM0,1Case number: 0 direct value writes by a SIMT Warp; 1 value movement by an MTE task queue.PROFILE_REPEAT_TIMESPositive integerNumber of times thatinsert_or_assignis repeatedly run in a single process. The comparison script sets this option to 1 so that msOpProf collects one complete insert path.Execution resultThe execution result is as follows, which indicates that the HashTable insertion result and value verification passed.Verification PASSEDPerformance AnalysisUse msOpProf to obtain performance data. The comparison script first runs the sample and checksVerification PASSED. After confirming that the result is correct, the script collects target kernel performance data by themsprof opcommand.The Case 0 test method is as follows:cmake -DSCENARIO_NUM0 -DPROFILE_REPEAT_TIMES1 -DCMAKE_ASC_ARCHITECTURESdav-3510 .. make -j ./hash_table_mte_queue 128 msprof op ./hash_table_mte_queue 128The Case 1 test method is as follows:cmake -DSCENARIO_NUM1 -DPROFILE_REPEAT_TIMES1 -DCMAKE_ASC_ARCHITECTURESdav-3510 .. make -j ./hash_table_mte_queue 128 msprof op ./hash_table_mte_queue 128A directory with theOPPROF_prefix is generated in the current directory to store performance analysis data of the current kernel.OPPROF_xxxx_XXXXXXXX ├── ArithmeticUtilization.csv ├── L2Cache.csv ├── Memory.csv ├── MemoryL0.csv ├── MemoryUB.csv ├── OpBasicInfo.csv ├── PipeUtilization.csv └── ResourceConflictRatio.csvView detailed performance analysis results:# View basic information such as Task Duration. cat ./OPPROF_*/OpBasicInfo.csv # View Vector/MTE pipeline time and ratio. cat ./OPPROF_*/PipeUtilization.csv【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
