issue/312: causal softmax算子支持bf16

src/infiniop/ops/causal_softmax/maca/causal_softmax_maca.maca

@@ -38,6 +38,12 @@ infiniStatus_t launchKernel(void *y, const void *x, infiniDtype_t dtype,
                                               batch_size, seq_len, total_seq_len,
                                               y_stride_b, y_stride_i,
                                               x_stride_b, x_stride_i);
+    } else if (dtype == INFINI_DTYPE_BF16) {
+        causalSoftmax<BLOCK_SIZE, __hpcc_bfloat16, float>
+            <<<grid, BLOCK_SIZE, 0, stream>>>((__hpcc_bfloat16 *)y, (const __hpcc_bfloat16 *)x,
+                                              batch_size, seq_len, total_seq_len,
+                                              y_stride_b, y_stride_i,
+                                              x_stride_b, x_stride_i);
     } else if (dtype == INFINI_DTYPE_F32) {
         causalSoftmax<BLOCK_SIZE, float, float>
             <<<grid, BLOCK_SIZE, 0, stream>>>((float *)y, (const float *)x,