Merge pull request #1 from aravindhbalaji1985/aravindhbalaji1985-patch-1

aravindhbalaji1985 · web-flow · commit 08f998122219 · 2025-04-02T12:41:31.000-07:00
Update gpu_prim.h
diff --git a/tensorflow/core/kernels/gpu_prim.h b/tensorflow/core/kernels/gpu_prim.h
@@ -44,10 +44,9 @@ __device__ __forceinline__ void ThreadStoreVolatilePtr<Eigen::half>(
       Eigen::numext::bit_cast<uint16_t>(val);
 }
 
-template <>
-__device__ __forceinline__ Eigen::half ThreadLoadVolatilePointer<Eigen::half>(
+__device__ __forceinline__ Eigen::half ThreadLoadVolatilePointer(
     Eigen::half *ptr, Int2Type<true> /*is_primitive*/) {
-  uint16_t result = *reinterpret_cast<volatile uint16_t *>(ptr);
+  const uint16_t result = *reinterpret_cast<volatile const uint16_t *>(ptr);
   return Eigen::numext::bit_cast<Eigen::half>(result);
 }
 
@@ -59,10 +58,8 @@ __device__ __forceinline__ void ThreadStoreVolatilePtr<Eigen::bfloat16>(
       Eigen::numext::bit_cast<uint16_t>(val);
 }
 
-template <>
-__device__ __forceinline__ Eigen::bfloat16
-ThreadLoadVolatilePointer<Eigen::bfloat16>(Eigen::bfloat16 *ptr,
-                                           Int2Type<true> /*is_primitive*/) {
+__device__ __forceinline__ Eigen::bfloat16 ThreadLoadVolatilePointer(
+  Eigen::bfloat16 *ptr, Int2Type<true> /*is_primitive*/) {
   uint16_t result = *reinterpret_cast<volatile uint16_t *>(ptr);
   return Eigen::numext::bit_cast<Eigen::bfloat16>(result);
 }

Original file line number	Diff line number	Diff line change
`@@ -44,10 +44,9 @@ __device__ __forceinline__ void ThreadStoreVolatilePtr<Eigen::half>(`
`44`	`44`	`Eigen::numext::bit_cast<uint16_t>(val);`
`45`	`45`	`}`
`46`	`46`
`47`		`-template <>`
`48`		`-__device__ __forceinline__ Eigen::half ThreadLoadVolatilePointer<Eigen::half>(`
	`47`	`+__device__ __forceinline__ Eigen::half ThreadLoadVolatilePointer(`
`49`	`48`	`Eigen::half ptr, Int2Type<true> /is_primitive*/) {`
`50`		`- uint16_t result = reinterpret_cast<volatile uint16_t >(ptr);`
	`49`	`+ const uint16_t result = reinterpret_cast<volatile const uint16_t >(ptr);`
`51`	`50`	`return Eigen::numext::bit_cast<Eigen::half>(result);`
`52`	`51`	`}`
`53`	`52`
`@@ -59,10 +58,8 @@ __device__ __forceinline__ void ThreadStoreVolatilePtr<Eigen::bfloat16>(`
`59`	`58`	`Eigen::numext::bit_cast<uint16_t>(val);`
`60`	`59`	`}`
`61`	`60`
`62`		`-template <>`
`63`		`-__device__ __forceinline__ Eigen::bfloat16`
`64`		`-ThreadLoadVolatilePointer<Eigen::bfloat16>(Eigen::bfloat16 *ptr,`
`65`		`- Int2Type<true> /is_primitive/) {`
	`61`	`+__device__ __forceinline__ Eigen::bfloat16 ThreadLoadVolatilePointer(`
	`62`	`+ Eigen::bfloat16 ptr, Int2Type<true> /is_primitive*/) {`
`66`	`63`	`uint16_t result = reinterpret_cast<volatile uint16_t >(ptr);`
`67`	`64`	`return Eigen::numext::bit_cast<Eigen::bfloat16>(result);`
`68`	`65`	`}`