[AMD][Diffusion] support timestep embedding kernel for AMD GPUs (#16766)

sgl-kernel/csrc/common_extension_rocm.cc

@@ -219,6 +219,18 @@ TORCH_LIBRARY_EXPAND(sgl_kernel, m) {
       "                 Tensor!? key, int head_size,"
       "                 Tensor cos_sin_cache, bool is_neox) -> ()");
   m.impl("rotary_embedding", torch::kCUDA, &rotary_embedding);
+  /*
+   * From csrc/sgl_diffusion/elementwise
+   */
+  m.def(
+      "timestep_embedding(Tensor input,"
+      "Tensor output,"
+      "int dim,"
+      "bool flip_sin_to_cos,"
+      "float downscale_freq_shift,"
+      "float scale,"
+      "int max_period) -> Tensor");
+  m.impl("timestep_embedding", torch::kCUDA, &timestep_embedding);
 }
 
 REGISTER_EXTENSION(common_ops)

sgl-kernel/csrc/sgl_diffusion/elementwise/timestep_embedding.cu

@@ -33,7 +33,8 @@ __global__ void timestep_embedding_kernel(
   if (row_idx >= batch_size) {
     return;
   }
-  float t_val = castToFloat(__ldg(&t_ptr[row_idx]));
+  // Use the portable LDG helper (maps to __ldg on CUDA, plain load on ROCm/HIP).
+  float t_val = castToFloat(SGLANG_LDG(&t_ptr[row_idx]));
   float* output_batch_base_ptr = output_ptr + row_idx * dim;
 
   // Calculate half dimension

sgl-kernel/include/hip/hip_math_def.h

@@ -29,7 +29,11 @@ template <typename T>
 __forceinline__ __device__ T shfl_xor_sync(unsigned mask, T var, int laneMask, int width = warpSize);
 
 template <typename srcDtype, typename destDtype>
-__forceinline__ __device__ destDtype cast(srcDtype val);
+__forceinline__ __device__ destDtype cast(srcDtype val) {
+  // Generic fallback used by most scalar types (int/float/double/etc).
+  // Specific types like fp16/bf16 have explicit specializations below.
+  return static_cast<destDtype>(val);
+}
 
 // specialization
 template <>
@@ -43,27 +47,27 @@ __forceinline__ __device__ int shfl_xor_sync(unsigned mask, int var, int laneMas
 }
 
 template <>
-__forceinline__ __device__ float cast(float val) {
+__forceinline__ __device__ float cast<float, float>(float val) {
   return val;
 }
 
 template <>
-__forceinline__ __device__ float cast(__half val) {
+__forceinline__ __device__ float cast<__half, float>(__half val) {
   return __half2float(val);
 }
 
 template <>
-__forceinline__ __device__ float cast(__hip_bfloat16 val) {
+__forceinline__ __device__ float cast<__hip_bfloat16, float>(__hip_bfloat16 val) {
   return __bfloat162float(val);
 }
 
 template <>
-__forceinline__ __device__ __half cast(float fval) {
+__forceinline__ __device__ __half cast<float, __half>(float fval) {
   return __float2half(fval);
 }
 
 template <>
-__forceinline__ __device__ __hip_bfloat16 cast(float fval) {
+__forceinline__ __device__ __hip_bfloat16 cast<float, __hip_bfloat16>(float fval) {
   return __float2bfloat16(fval);
 }
 

sgl-kernel/setup_rocm.py

@@ -54,6 +54,7 @@ sources = [
     "csrc/speculative/eagle_utils.cu",
     "csrc/kvcacheio/transfer.cu",
     "csrc/elementwise/pos_enc.cu",
+    "csrc/sgl_diffusion/elementwise/timestep_embedding.cu",
 ]
 
 cxx_flags = ["-O3"]