[AMD] Reorganize hip-related header files in sgl-kernel (#9320)
This commit is contained in:
@@ -1,7 +1,7 @@
|
||||
// Copyright (c) Microsoft Corporation.
|
||||
// Licensed under the MIT license.
|
||||
#pragma once
|
||||
#if defined(__HIP_PLATFORM_AMD__)
|
||||
#ifdef USE_ROCM
|
||||
#include <hip/hip_fp16.h>
|
||||
#else
|
||||
#include <cuda_bf16.h>
|
||||
|
||||
@@ -25,7 +25,7 @@
|
||||
#include "utils.h"
|
||||
|
||||
#else
|
||||
#include "hip_act_and_mul.cuh"
|
||||
#include "hip/hip_act_and_mul.cuh"
|
||||
#endif
|
||||
|
||||
// Adapted from flashinfer activation
|
||||
|
||||
@@ -69,7 +69,7 @@ __global__ void per_tensor_quant_fp8_kernel(
|
||||
#pragma unroll
|
||||
for (uint32_t j = 0; j < VEC_SIZE; ++j) {
|
||||
float val = fmax(fmin(static_cast<float>(input_vec[j]) * scale_val, FP8_E4M3_MAX), -FP8_E4M3_MAX);
|
||||
#ifndef USE_ROCM
|
||||
#if !defined(USE_ROCM) || defined(HIP_FP8_TYPE_E4M3)
|
||||
output_arr[j] = static_cast<DST_DTYPE>(val);
|
||||
#else
|
||||
output_arr[j] = c10::Float8_e4m3fnuz(
|
||||
@@ -83,7 +83,7 @@ __global__ void per_tensor_quant_fp8_kernel(
|
||||
const int32_t remaining_start = num_vec_elems * VEC_SIZE;
|
||||
for (int32_t idx = remaining_start + gid; idx < num_elements; idx += grid_size) {
|
||||
float val = fmax(-FP8_E4M3_MAX, fmin(static_cast<float>(input[idx]) * scale_val, FP8_E4M3_MAX));
|
||||
#ifndef USE_ROCM
|
||||
#if !defined(USE_ROCM) || defined(HIP_FP8_TYPE_E4M3)
|
||||
output[idx] = static_cast<DST_DTYPE>(val);
|
||||
#else
|
||||
output[idx] = c10::Float8_e4m3fnuz(
|
||||
|
||||
@@ -67,7 +67,7 @@ __global__ void per_token_quant_fp8_kernel(
|
||||
for (uint32_t j = 0; j < kVecSize; ++j) {
|
||||
float val = static_cast<float>(input_vec[j]) * scale_inv;
|
||||
val = fmaxf(fminf(val, FP8_E4M3_MAX), -FP8_E4M3_MAX);
|
||||
#ifndef USE_ROCM
|
||||
#if !defined(USE_ROCM) || defined(HIP_FP8_TYPE_E4M3)
|
||||
output_arr[j] = static_cast<DST_DTYPE>(val);
|
||||
#else
|
||||
output_arr[j] = c10::Float8_e4m3fnuz(
|
||||
@@ -143,7 +143,7 @@ __global__ void per_token_quant_fp8_small_batch_kernel(
|
||||
#pragma unroll
|
||||
for (uint32_t j = 0; j < kVecSize; ++j) {
|
||||
float val = fmaxf(fminf(static_cast<float>(input_vec[j]) * scale_inv, FP8_E4M3_MAX), -FP8_E4M3_MAX);
|
||||
#ifndef USE_ROCM
|
||||
#if !defined(USE_ROCM) || defined(HIP_FP8_TYPE_E4M3)
|
||||
output_arr[j] = static_cast<DST_DTYPE>(val);
|
||||
#else
|
||||
output_arr[j] = c10::Float8_e4m3fnuz(
|
||||
|
||||
@@ -21,8 +21,6 @@ limitations under the License.
|
||||
|
||||
#include "utils.h"
|
||||
|
||||
#define WARP_SIZE 32
|
||||
|
||||
#define VEC_SIZE 4
|
||||
using Vec = int4;
|
||||
|
||||
|
||||
Reference in New Issue
Block a user