From 3cc0a3fcaf1a6d42a84d4f0864fa0125e0dbe0a2 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Tue, 22 Feb 2022 15:16:13 +0000 Subject: [PATCH] SWDEV-321656 - Address incorrect warp size with hiprtc Change-Id: I97051d4e2cf27cb073c39663c62e43d5c19b1a12 --- .../hip/amd_detail/amd_device_functions.h | 421 +---------------- .../hip/amd_detail/amd_warp_functions.h | 444 ++++++++++++++++++ hipamd/src/hiprtc/CMakeLists.txt | 3 +- hipamd/src/hiprtc/cmake/HIPRTC.cmake | 2 + 4 files changed, 452 insertions(+), 418 deletions(-) create mode 100644 hipamd/include/hip/amd_detail/amd_warp_functions.h diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h index eb66b86697..8383a2d720 100644 --- a/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -386,423 +386,10 @@ __device__ static inline int __hip_move_dpp_N(int src) { bound_ctrl); } -static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE; +#if !defined(__HIPCC_RTC__) +#include "amd_warp_functions.h" +#endif -__device__ -inline -int __shfl(int var, int src_lane, int width = warpSize) { - int self = __lane_id(); - int index = src_lane + (self & ~(width-1)); - return __builtin_amdgcn_ds_bpermute(index<<2, var); -} -__device__ -inline -unsigned int __shfl(unsigned int var, int src_lane, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.u = var; - tmp.i = __shfl(tmp.i, src_lane, width); - return tmp.u; -} -__device__ -inline -float __shfl(float var, int src_lane, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.f = var; - tmp.i = __shfl(tmp.i, src_lane, width); - return tmp.f; -} -__device__ -inline -double __shfl(double var, int src_lane, int width = warpSize) { - static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl(tmp[0], src_lane, width); - tmp[1] = __shfl(tmp[1], src_lane, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -long __shfl(long var, int src_lane, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl(tmp[0], src_lane, width); - tmp[1] = __shfl(tmp[1], src_lane, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(long) == sizeof(int), ""); - return static_cast(__shfl(static_cast(var), src_lane, width)); - #endif -} -__device__ -inline -unsigned long __shfl(unsigned long var, int src_lane, int width = warpSize) { - #ifndef _MSC_VER - static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); - - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl(tmp[0], src_lane, width); - tmp[1] = __shfl(tmp[1], src_lane, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); - return static_cast(__shfl(static_cast(var), src_lane, width)); - #endif -} -__device__ -inline -long long __shfl(long long var, int src_lane, int width = warpSize) -{ - static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl(tmp[0], src_lane, width); - tmp[1] = __shfl(tmp[1], src_lane, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -unsigned long long __shfl(unsigned long long var, int src_lane, int width = warpSize) { - static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); - - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl(tmp[0], src_lane, width); - tmp[1] = __shfl(tmp[1], src_lane, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} - -__device__ -inline -int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) { - int self = __lane_id(); - int index = self - lane_delta; - index = (index < (self & ~(width-1)))?self:index; - return __builtin_amdgcn_ds_bpermute(index<<2, var); -} -__device__ -inline -unsigned int __shfl_up(unsigned int var, unsigned int lane_delta, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.u = var; - tmp.i = __shfl_up(tmp.i, lane_delta, width); - return tmp.u; -} -__device__ -inline -float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.f = var; - tmp.i = __shfl_up(tmp.i, lane_delta, width); - return tmp.f; -} -__device__ -inline -double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) { - static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_up(tmp[0], lane_delta, width); - tmp[1] = __shfl_up(tmp[1], lane_delta, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -long __shfl_up(long var, unsigned int lane_delta, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_up(tmp[0], lane_delta, width); - tmp[1] = __shfl_up(tmp[1], lane_delta, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(long) == sizeof(int), ""); - return static_cast(__shfl_up(static_cast(var), lane_delta, width)); - #endif -} - -__device__ -inline -unsigned long __shfl_up(unsigned long var, unsigned int lane_delta, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); - - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_up(tmp[0], lane_delta, width); - tmp[1] = __shfl_up(tmp[1], lane_delta, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); - return static_cast(__shfl_up(static_cast(var), lane_delta, width)); - #endif -} - -__device__ -inline -long long __shfl_up(long long var, unsigned int lane_delta, int width = warpSize) -{ - static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(uint64_t), ""); - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_up(tmp[0], lane_delta, width); - tmp[1] = __shfl_up(tmp[1], lane_delta, width); - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} - -__device__ -inline -unsigned long long __shfl_up(unsigned long long var, unsigned int lane_delta, int width = warpSize) -{ - static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_up(tmp[0], lane_delta, width); - tmp[1] = __shfl_up(tmp[1], lane_delta, width); - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} - -__device__ -inline -int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) { - int self = __lane_id(); - int index = self + lane_delta; - index = (int)((self&(width-1))+lane_delta) >= width?self:index; - return __builtin_amdgcn_ds_bpermute(index<<2, var); -} -__device__ -inline -unsigned int __shfl_down(unsigned int var, unsigned int lane_delta, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.u = var; - tmp.i = __shfl_down(tmp.i, lane_delta, width); - return tmp.u; -} -__device__ -inline -float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.f = var; - tmp.i = __shfl_down(tmp.i, lane_delta, width); - return tmp.f; -} -__device__ -inline -double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) { - static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_down(tmp[0], lane_delta, width); - tmp[1] = __shfl_down(tmp[1], lane_delta, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -long __shfl_down(long var, unsigned int lane_delta, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_down(tmp[0], lane_delta, width); - tmp[1] = __shfl_down(tmp[1], lane_delta, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(long) == sizeof(int), ""); - return static_cast(__shfl_down(static_cast(var), lane_delta, width)); - #endif -} -__device__ -inline -unsigned long __shfl_down(unsigned long var, unsigned int lane_delta, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); - - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_down(tmp[0], lane_delta, width); - tmp[1] = __shfl_down(tmp[1], lane_delta, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); - return static_cast(__shfl_down(static_cast(var), lane_delta, width)); - #endif -} -__device__ -inline -long long __shfl_down(long long var, unsigned int lane_delta, int width = warpSize) -{ - static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(uint64_t), ""); - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_down(tmp[0], lane_delta, width); - tmp[1] = __shfl_down(tmp[1], lane_delta, width); - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -unsigned long long __shfl_down(unsigned long long var, unsigned int lane_delta, int width = warpSize) -{ - static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_down(tmp[0], lane_delta, width); - tmp[1] = __shfl_down(tmp[1], lane_delta, width); - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} - -__device__ -inline -int __shfl_xor(int var, int lane_mask, int width = warpSize) { - int self = __lane_id(); - int index = self^lane_mask; - index = index >= ((self+width)&~(width-1))?self:index; - return __builtin_amdgcn_ds_bpermute(index<<2, var); -} -__device__ -inline -unsigned int __shfl_xor(unsigned int var, int lane_mask, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.u = var; - tmp.i = __shfl_xor(tmp.i, lane_mask, width); - return tmp.u; -} -__device__ -inline -float __shfl_xor(float var, int lane_mask, int width = warpSize) { - union { int i; unsigned u; float f; } tmp; tmp.f = var; - tmp.i = __shfl_xor(tmp.i, lane_mask, width); - return tmp.f; -} -__device__ -inline -double __shfl_xor(double var, int lane_mask, int width = warpSize) { - static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_xor(tmp[0], lane_mask, width); - tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -long __shfl_xor(long var, int lane_mask, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(uint64_t), ""); - - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_xor(tmp[0], lane_mask, width); - tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(long) == sizeof(int), ""); - return static_cast(__shfl_xor(static_cast(var), lane_mask, width)); - #endif -} -__device__ -inline -unsigned long __shfl_xor(unsigned long var, int lane_mask, int width = warpSize) -{ - #ifndef _MSC_VER - static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); - - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_xor(tmp[0], lane_mask, width); - tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; - #else - static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); - return static_cast(__shfl_xor(static_cast(var), lane_mask, width)); - #endif -} -__device__ -inline -long long __shfl_xor(long long var, int lane_mask, int width = warpSize) -{ - static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(uint64_t), ""); - int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_xor(tmp[0], lane_mask, width); - tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} -__device__ -inline -unsigned long long __shfl_xor(unsigned long long var, int lane_mask, int width = warpSize) -{ - static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); - unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); - tmp[0] = __shfl_xor(tmp[0], lane_mask, width); - tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); - unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - return tmp1; -} #define MASK1 0x00ff00ff #define MASK2 0xff00ff00 diff --git a/hipamd/include/hip/amd_detail/amd_warp_functions.h b/hipamd/include/hip/amd_detail/amd_warp_functions.h new file mode 100644 index 0000000000..ebdc86374d --- /dev/null +++ b/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -0,0 +1,444 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_WARP_FUNCTIONS_H +#define HIP_INCLUDE_HIP_AMD_DETAIL_WARP_FUNCTIONS_H + +static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE; + +__device__ +inline +int __shfl(int var, int src_lane, int width = warpSize) { + int self = __lane_id(); + int index = src_lane + (self & ~(width-1)); + return __builtin_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl(unsigned int var, int src_lane, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.u = var; + tmp.i = __shfl(tmp.i, src_lane, width); + return tmp.u; +} +__device__ +inline +float __shfl(float var, int src_lane, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.f = var; + tmp.i = __shfl(tmp.i, src_lane, width); + return tmp.f; +} +__device__ +inline +double __shfl(double var, int src_lane, int width = warpSize) { + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +long __shfl(long var, int src_lane, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl(static_cast(var), src_lane, width)); + #endif +} +__device__ +inline +unsigned long __shfl(unsigned long var, int src_lane, int width = warpSize) { + #ifndef _MSC_VER + static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); + + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); + return static_cast(__shfl(static_cast(var), src_lane, width)); + #endif +} +__device__ +inline +long long __shfl(long long var, int src_lane, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +unsigned long long __shfl(unsigned long long var, int src_lane, int width = warpSize) { + static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); + + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} + +__device__ +inline +int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) { + int self = __lane_id(); + int index = self - lane_delta; + index = (index < (self & ~(width-1)))?self:index; + return __builtin_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl_up(unsigned int var, unsigned int lane_delta, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.u = var; + tmp.i = __shfl_up(tmp.i, lane_delta, width); + return tmp.u; +} +__device__ +inline +float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.f = var; + tmp.i = __shfl_up(tmp.i, lane_delta, width); + return tmp.f; +} +__device__ +inline +double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) { + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +long __shfl_up(long var, unsigned int lane_delta, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl_up(static_cast(var), lane_delta, width)); + #endif +} + +__device__ +inline +unsigned long __shfl_up(unsigned long var, unsigned int lane_delta, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); + + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); + return static_cast(__shfl_up(static_cast(var), lane_delta, width)); + #endif +} + +__device__ +inline +long long __shfl_up(long long var, unsigned int lane_delta, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} + +__device__ +inline +unsigned long long __shfl_up(unsigned long long var, unsigned int lane_delta, int width = warpSize) +{ + static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} + +__device__ +inline +int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) { + int self = __lane_id(); + int index = self + lane_delta; + index = (int)((self&(width-1))+lane_delta) >= width?self:index; + return __builtin_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl_down(unsigned int var, unsigned int lane_delta, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.u = var; + tmp.i = __shfl_down(tmp.i, lane_delta, width); + return tmp.u; +} +__device__ +inline +float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.f = var; + tmp.i = __shfl_down(tmp.i, lane_delta, width); + return tmp.f; +} +__device__ +inline +double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) { + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +long __shfl_down(long var, unsigned int lane_delta, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl_down(static_cast(var), lane_delta, width)); + #endif +} +__device__ +inline +unsigned long __shfl_down(unsigned long var, unsigned int lane_delta, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); + + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); + return static_cast(__shfl_down(static_cast(var), lane_delta, width)); + #endif +} +__device__ +inline +long long __shfl_down(long long var, unsigned int lane_delta, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +unsigned long long __shfl_down(unsigned long long var, unsigned int lane_delta, int width = warpSize) +{ + static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} + +__device__ +inline +int __shfl_xor(int var, int lane_mask, int width = warpSize) { + int self = __lane_id(); + int index = self^lane_mask; + index = index >= ((self+width)&~(width-1))?self:index; + return __builtin_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl_xor(unsigned int var, int lane_mask, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.u = var; + tmp.i = __shfl_xor(tmp.i, lane_mask, width); + return tmp.u; +} +__device__ +inline +float __shfl_xor(float var, int lane_mask, int width = warpSize) { + union { int i; unsigned u; float f; } tmp; tmp.f = var; + tmp.i = __shfl_xor(tmp.i, lane_mask, width); + return tmp.f; +} +__device__ +inline +double __shfl_xor(double var, int lane_mask, int width = warpSize) { + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +long __shfl_xor(long var, int lane_mask, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(long) == 2 * sizeof(int), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(long) == sizeof(int), ""); + return static_cast(__shfl_xor(static_cast(var), lane_mask, width)); + #endif +} +__device__ +inline +unsigned long __shfl_xor(unsigned long var, int lane_mask, int width = warpSize) +{ + #ifndef _MSC_VER + static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); + + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; + #else + static_assert(sizeof(unsigned long) == sizeof(unsigned int), ""); + return static_cast(__shfl_xor(static_cast(var), lane_mask, width)); + #endif +} +__device__ +inline +long long __shfl_xor(long long var, int lane_mask, int width = warpSize) +{ + static_assert(sizeof(long long) == 2 * sizeof(int), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} +__device__ +inline +unsigned long long __shfl_xor(unsigned long long var, int lane_mask, int width = warpSize) +{ + static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); + unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; +} + +#endif diff --git a/hipamd/src/hiprtc/CMakeLists.txt b/hipamd/src/hiprtc/CMakeLists.txt index f72758ccd9..c37f6127cc 100644 --- a/hipamd/src/hiprtc/CMakeLists.txt +++ b/hipamd/src/hiprtc/CMakeLists.txt @@ -115,6 +115,7 @@ set(HIPRTC_GEN_HEADER "${HIPRTC_GEN_DIR}/hipRTC_header.h") set(HIPRTC_GEN_MCIN "${HIPRTC_GEN_DIR}/hipRTC_header.mcin") set(HIPRTC_GEN_PREPROCESSED "${HIPRTC_GEN_DIR}/hipRTC") set(HIPRTC_GEN_OBJ "${HIPRTC_GEN_DIR}/hipRTC_header${CMAKE_CXX_OUTPUT_EXTENSION}") +set(HIPRTC_WARP_FUNCS "${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_functions.h") # Generate required HIPRTC files. FILE(MAKE_DIRECTORY ${HIPRTC_GEN_DIR}) @@ -127,7 +128,7 @@ generate_hiprtc_mcin("${HIPRTC_GEN_MCIN}" "${HIPRTC_GEN_PREPROCESSED}") add_custom_command( OUTPUT ${HIPRTC_GEN_PREPROCESSED} COMMAND $ -O3 --rocm-path=${PROJECT_SOURCE_DIR}/include/.. -std=c++17 -nogpulib --hip-version=3.6 -isystem ${HIP_COMMON_INCLUDE_DIR} -isystem ${PROJECT_SOURCE_DIR}/include -isystem ${PROJECT_BINARY_DIR}/include -isystem ${CMAKE_CURRENT_SOURCE_DIR}/include --cuda-device-only -D__HIPCC_RTC__ -x hip ${HIPRTC_GEN_HEADER} -E -o ${HIPRTC_GEN_PREPROCESSED} - COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE} + COMMAND ${CMAKE_COMMAND} -DHIPRTC_ADD_MACROS=1 -DHIPRTC_WARP_HEADER_FILE=${HIPRTC_WARP_FUNCS} -DHIPRTC_PREPROCESSED_FILE=${HIPRTC_GEN_PREPROCESSED} -P ${HIPRTC_CMAKE} DEPENDS clang ${HIPRTC_GEN_HEADER}) add_custom_command( OUTPUT ${HIPRTC_GEN_OBJ} diff --git a/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/hipamd/src/hiprtc/cmake/HIPRTC.cmake index 0443588a4b..209a484e3c 100644 --- a/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -56,6 +56,8 @@ if(HIPRTC_ADD_MACROS) message(STATUS "Appending hiprtc macros to ${HIPRTC_PREPROCESSED_FILE}.") get_hiprtc_macros(HIPRTC_DEFINES) FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_DEFINES}") + FILE(READ "${HIPRTC_WARP_HEADER_FILE}" HIPRTC_WARP_HEADER) + FILE(APPEND ${HIPRTC_PREPROCESSED_FILE} "${HIPRTC_WARP_HEADER}") endif() macro(generate_hiprtc_header HiprtcHeader)