From 1ce36d001b0f50f492963bdabc54e7e49bf1b005 Mon Sep 17 00:00:00 2001
From: Matthew Wong <Matthew.Wong2@amd.com>
Date: Wed, 27 Mar 2024 01:02:20 -0500
Subject: [PATCH] Remove redundant inline specifiers; preparing for upstream

---
 csrc/layernorm_kernels.cu | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/csrc/layernorm_kernels.cu b/csrc/layernorm_kernels.cu
index 5c9425fec7ad3..5ee8b2609027e 100644
--- a/csrc/layernorm_kernels.cu
+++ b/csrc/layernorm_kernels.cu
@@ -103,7 +103,7 @@ struct _f16Vec {
   using T2 = typename Converter::packed_hip_type;
   T1 data[width];
 
-  __device__ inline _f16Vec& operator+=(const _f16Vec<scalar_t, width>& other) {
+  __device__ _f16Vec& operator+=(const _f16Vec<scalar_t, width>& other) {
     if constexpr (width % 2 == 0) {
       #pragma unroll
       for (int i = 0; i < width; i += 2) {
@@ -120,7 +120,7 @@ struct _f16Vec {
     return *this;
   }
 
-  __device__ inline _f16Vec& operator*=(const _f16Vec<scalar_t, width>& other) {
+  __device__ _f16Vec& operator*=(const _f16Vec<scalar_t, width>& other) {
     if constexpr (width % 2 == 0) {
       #pragma unroll
       for (int i = 0; i < width; i += 2) {
@@ -137,7 +137,7 @@ struct _f16Vec {
     return *this;
   }
 
-  __device__ inline _f16Vec& operator*=(const float scale) {
+  __device__ _f16Vec& operator*=(const float scale) {
     if constexpr (width % 2 == 0) {
       #pragma unroll
       for (int i = 0; i < width; i += 2) {
@@ -158,7 +158,7 @@ struct _f16Vec {
     return *this;
   }
 
-  __device__ inline float sum_squares() const {
+  __device__ float sum_squares() const {
     float result = 0.0f;
     if constexpr (width % 2 == 0) {
       #pragma unroll