Don't initialize TMA output buffer (#3105)

TMA will automatically fill zero, besides, the initialization will race
with TMA itself as there is no sync between initialization and TMA.

Matmul perf after enabling circular buffering:
```markdown
 Time (%)  Total Time (ns)  Instances  Avg (ns)   Med (ns)   Min (ns)  Max (ns)  StdDev (ns)                                                  Name

 --------  ---------------  ---------  ---------  ---------  --------  --------  -----------  ----------------------------------------------------------------------------------------------------
     57.3          3218718          1  3218718.0  3218718.0   3218718   3218718          0.0  <unnamed>::nvfuser_none_f0_c0_r0_g0(<unnamed>::Tensor<<unnamed>::__half, (int)3, (int)3>, <unnamed>…
     12.5           700153          1   700153.0   700153.0    700153    700153          0.0  nvjet_hsh_192x192_64x3_2x1_v_bz_coopB_NTN
```

csrc/device_lower/pass/allocation.cpp

@@ -508,10 +508,10 @@ class AllocationInserter : public kir::ExprMutator {
         init = default_val;
       }
 
-      if (ir_utils::isCpAsyncOp(expr)) {
+      if (ir_utils::isCpAsyncOp(expr) || ir_utils::isCpAsyncBulk(expr)) {
         NVF_CHECK(
             init == nullptr || init->isZero(),
-            "cp.async initialized with non-zero is not supported");
+            "cp.async and cp.async.bulk initialized with non-zero is not supported");
         // cp.async will automatically fill zero when out of bound
         init = nullptr;
       }

tests/cpp/test_matmul.cpp

@@ -3624,9 +3624,8 @@ TEST_F(HopperMatmulTest, HSH_NT_128BSwizzle) {
 
   inlineMost();
 
-  // TODO: looks like this test will hang if I enable this
-  // tv0c->circularBuffer(/*number_of_stages=*/4);
-  // tv1c->circularBuffer(/*number_of_stages=*/4);
+  tv0c->circularBuffer(/*number_of_stages=*/4);
+  tv1c->circularBuffer(/*number_of_stages=*/4);
 
   auto inputs =
       matmulAtInput3DHopperSS(M, N, K, layout, data_type_to_aten(dtype));