issue/312: 沐曦BF16编译问题

src/infiniop/devices/maca/maca_kernel_common.h

@@ -60,4 +60,9 @@ __forceinline__ __device__ __half
 exp_(const __half x) {
     return hexp(x);
 }
+
+__forceinline__ __device__ __hpcc_bfloat16;
+exp_(const __hpcc_bfloat16; x) {
+    return hexp(x);
+}
 #endif

src/infiniop/ops/random_sample/maca/random_sample_kernel.h

@@ -107,6 +107,11 @@ struct CudaTval<fp16_t> {
     using Type = half;
 };
 
+template <>
+struct CudaTval<bf16_t> {
+    using Type = __hpcc_bfloat16;
+};
+
 // ↑↑↑ 通过特化将 fp16_t 转换为 half
 // ↓↓↓ 用于采样过程的小型 kernel
 

src/infiniop/ops/random_sample/maca/random_sample_maca.maca

@@ -34,15 +34,16 @@ infiniStatus_t Descriptor::create(
         workspace_size = workspace_result.take();                       \
     } break
 
-#define CASE_I(CASE, Tidx)                          \
-    case CASE:                                      \
-        switch (info.dt_p) {                        \
-            CASE_P(INFINI_DTYPE_F16, Tidx, half);   \
-            CASE_P(INFINI_DTYPE_F32, Tidx, float);  \
-            CASE_P(INFINI_DTYPE_F64, Tidx, double); \
-        default:                                    \
-            abort();                                \
-        }                                           \
+#define CASE_I(CASE, Tidx)                                    \
+    case CASE:                                                \
+        switch (info.dt_p) {                                  \
+            CASE_P(INFINI_DTYPE_F16, Tidx, half);             \
+            CASE_P(INFINI_DTYPE_BF16, Tidx, __hpcc_bfloat16); \
+            CASE_P(INFINI_DTYPE_F32, Tidx, float);            \
+            CASE_P(INFINI_DTYPE_F64, Tidx, double);           \
+        default:                                              \
+            abort();                                          \
+        }                                                     \
         break
 
     switch (info.dt_i) {