From e7a40bf6f6ba404bccabd79e63629d584d7277a9 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Wed, 8 Sep 2021 08:30:16 -0400 Subject: [PATCH] SWDEV-302147 - fix __hip_ds_permutef device function to return float Change-Id: Ie452e90a7f7410401df4a66de468ad850403953c [ROCm/clr commit: 673c2deed1153d230f4acb822a23e73d50c73f0f] --- .../clr/hipamd/include/hip/amd_detail/amd_device_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index 5666c62e02..f9b9570dfc 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -283,9 +283,9 @@ __device__ static inline unsigned __hip_ds_permute(int index, unsigned src) { } __device__ static inline float __hip_ds_permutef(int index, float src) { - union { int i; unsigned u; float f; } tmp; tmp.u = src; + union { int i; unsigned u; float f; } tmp; tmp.f = src; tmp.i = __builtin_amdgcn_ds_permute(index, tmp.i); - return tmp.u; + return tmp.f; } #define __hip_ds_swizzle(src, pattern) __hip_ds_swizzle_N<(pattern)>((src))