replace more asm code

Signed-off-by: Zhenhuan Chen <chenzhh3671@gmail.com>

Author: dc3671

cpp/tensorrt_llm/kernels/fusedLayernormKernels/low_latency_layernorm.cuh

@@ -178,8 +178,8 @@ struct LowLatencyLayerNorm
 #if (defined(__CUDA_ARCH__) && (__CUDACC_VER_MAJOR__ >= 12))
         if constexpr (arch::is_major_v<9> || arch::is_major_v<10>)
         {
-            asm volatile("griddepcontrol.wait;\n");
-            asm volatile("griddepcontrol.launch_dependents;\n");
+            cudaGridDependencySynchronize();
+            cudaTriggerProgrammaticLaunchCompletion();
         }
 #endif
         load_to_register(&param.input[work_id * param.n], data, param.n);

cpp/tensorrt_llm/kernels/fusedLayernormKernels/ws_layernorm.cuh

@@ -211,7 +211,7 @@ struct WarpSpecializedLayerNorm
 
                 if constexpr (FIRST_RUN)
                 {
-                    asm volatile("griddepcontrol.wait;\n");
+                    cudaGridDependencySynchronize();
                 }
 
                 for (int i = 0; i < Traits::M_BLOCK; i++)
@@ -817,7 +817,7 @@ struct WarpSpecializedLayerNorm
                 {
                     scheduler(lane_id, gridDim.x * gridDim.y * gridDim.z, param, shared);
                     // PRE-EXIT after all tiles have been scheduled.
-                    asm volatile("griddepcontrol.launch_dependents;\n");
+                    cudaTriggerProgrammaticLaunchCompletion();
                 }
                 else if (warp_id == 1)
                 {

cpp/tensorrt_llm/kernels/llama4MinLatencyKernels/llama4Bf16Bf16Gemm.cu

@@ -60,7 +60,7 @@ __global__ void llama4_bf16_bf16_gemm_kernel(int num_tokens,
         b_vec[chunk] = reinterpret_cast<aligned_bf16x4 const*>(B)[row * GEMM_K / VEC_SIZE + base_idx];
     }
 
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 
     // Process 5 chunks of 4 elements each
 #pragma unroll

cpp/tensorrt_llm/kernels/llama4MinLatencyKernels/llama4Fp8Bf16GemmAttnScalingPerBlockTemplate.cuh

@@ -100,7 +100,7 @@ __launch_bounds__(BLOCK_SIZE) __global__ void llama4_fp8_bf16_gemm_attn_scaling_
 #endif
 
 #if ENABLE_ACQBULK
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 #endif
 
     // Processing 8 elements each

cpp/tensorrt_llm/kernels/llama4MinLatencyKernels/llama4Fp8Bf16GemmPerBlockTemplate.cuh

@@ -89,7 +89,7 @@ __launch_bounds__(BLOCK_SIZE) __global__ void llama4_fp8_bf16_gemm_per_block_ker
 #endif
 
 #if ENABLE_ACQBULK
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 #endif
 
     // Processing 8 elements each

cpp/tensorrt_llm/kernels/llama4MinLatencyKernels/llama4Fp8Bf16GemmPerWarpTemplate.cuh

@@ -91,7 +91,7 @@ __launch_bounds__(BLOCK_SIZE) __global__ void llama4_fp8_bf16_gemm_per_warp_kern
 #endif
 
 #if ENABLE_ACQBULK
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 #endif
 
     // Processing 8 elements each

cpp/tensorrt_llm/kernels/llama4MinLatencyKernels/llama4Fp8Fp8GemmSwiGLUPerBlockTemplate.cuh

@@ -98,7 +98,7 @@ __launch_bounds__(BLOCK_SIZE) __global__ void llama4_fp8_fp8_gemm_swiglu_per_blo
 #endif
 
 #if ENABLE_ACQBULK
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 #endif
 
     // Processing 8 elements each

cpp/tensorrt_llm/kernels/llama4MinLatencyKernels/llama4MinLatencyMoEOp.cu

@@ -82,7 +82,7 @@ __global__ void llama4_moe_fc13_swiglu_fp8_kernel(int num_tokens,
 
     // Logits depends on the previous kernel, so we cannot prefetch anything.
 #if ENABLE_ACQBULK
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 #endif
 
     // Perform top1 within the current thread, which processes 4 experts.
@@ -242,7 +242,7 @@ __global__ void llama4_moe_fc2_fp8_kernel(int num_tokens,
     scaling_factors_shared[tid] = scaling_factors[tid];
 
 #if ENABLE_ACQBULK
-    asm volatile("griddepcontrol.wait;" ::: "memory");
+    cudaGridDependencySynchronize();
 #endif
 
     // Select the corresponding expert weight.