revert silu_and_mul to debug

Edenzzzz · Edenzzzz · commit 9a62dda4d2cd · 2025-11-21T14:14:38.000Z
diff --git a/sgl-kernel/csrc/elementwise/activation.cu b/sgl-kernel/csrc/elementwise/activation.cu
@@ -85,34 +85,20 @@ __device__ __forceinline__ T gelu_tanh(const T& x) {
 void silu_and_mul(at::Tensor& out, at::Tensor& input, bool enable_pdl) {
   int d = input.size(-1) / 2;
   int64_t num_tokens = input.numel() / input.size(-1);
+  dim3 grid(num_tokens);
+
   const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
-  const c10::cuda::OptionalCUDAGuard device_guard(device_of(input));
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
 
   DISPATCH_PYTORCH_DTYPE_TO_CTYPE_FLOAT_FP16(input.scalar_type(), c_type, [&] {
     uint32_t vec_size = 16 / sizeof(c_type);
-#if USE_ROCM
-    dim3 grid(num_tokens);
     dim3 block(std::min(d / vec_size, 1024U));
+#if USE_ROCM
     sgl_hip::activation::act_and_mul_kernel<c_type, silu>
         <<<grid, block, 0, stream>>>(static_cast<c_type*>(out.data_ptr()), static_cast<c_type*>(input.data_ptr()), d);
 #else
-    cudaLaunchConfig_t config;
-    config.gridDim = num_tokens;
-    config.blockDim = std::min(d / vec_size, 1024U);
-    config.dynamicSmemBytes = 0;
-    config.stream = stream;
-    cudaLaunchAttribute attrs[1];
-    attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-    attrs[0].val.programmaticStreamSerializationAllowed = enable_pdl;
-    config.numAttrs = 1;
-    config.attrs = attrs;
-
-    auto kernel = flashinfer::activation::act_and_mul_kernel<c_type, silu>;
-    cudaLaunchKernelEx(
-        &config, kernel, static_cast<c_type*>(out.data_ptr()), static_cast<c_type*>(input.data_ptr()), d);
-
-    cudaError_t err = cudaGetLastError();
-    TORCH_CHECK(err == cudaSuccess, "Failed to launch kernel: ", cudaGetErrorString(err));
+    flashinfer::activation::act_and_mul_kernel<c_type, silu>
+        <<<grid, block, 0, stream>>>(static_cast<c_type*>(out.data_ptr()), static_cast<c_type*>(input.data_ptr()), d);
 #endif
     return true;
   });
