AI/rocm-systems
2
0
Derivar 0
Código Questões Pedidos de integração Operações Pacotes Planeamentos Lançamentos Wiki Trabalho

SWDEV-395996 - Add HIPRTC support for missing hip headers

Ver a proveniência 
hip headers such as hip_math_constants.h, library_types.h,
hip_common.h, channel_descriptor.h, device_functions.h, hip_complex.h,
hip_texture_types.h, math_functions.h, surface_types.h
are added in HIPRTC

Change-Id: I4a4c198449ceb609c3ff55e00b43056c1f085431
Este cometimento está contido em:
Satyanvesh Dittakavi
2023-04-26 18:13:36 +00:00
ascendente e1d34cb24f
cometimento d7d0f11318
17 ficheiros modificados com 271 adições e 1481 eliminações
Descarregar ficheiro patch Descarregar ficheiro diff Expandir todos os ficheiros Colapsar todos os ficheiros
hipamd/include/hip/amd_detail/amd_channel_descriptor.h
+12 -1
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -23,9 +23,16 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
#define HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wold-style-cast"
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#include <hip/driver_types.h>
#include <hip/amd_detail/amd_hip_vector_types.h>
#endif
#ifdef __cplusplus
@@ -353,4 +360,8 @@ struct hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
#endif /* __cplusplus */
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif /* !HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H */
hipamd/include/hip/amd_detail/amd_device_functions.h
+25 -6
Ver ficheiro
@@ -23,18 +23,32 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_FUNCTIONS_H
#define HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_FUNCTIONS_H
#include "host_defines.h"
#include "math_fwd.h"
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wsign-conversion"
#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
#pragma clang diagnostic ignored "-Wold-style-cast"
#pragma clang diagnostic ignored "-Wshorten-64-to-32"
#pragma clang diagnostic ignored "-Wimplicit-int-conversion"
#pragma clang diagnostic ignored "-Wimplicit-float-conversion"
#pragma clang diagnostic ignored "-Wmissing-noreturn"
#pragma clang diagnostic ignored "-Wimplicit-fallthrough"
#pragma clang diagnostic ignored "-Wunneeded-internal-declaration"
#pragma clang diagnostic ignored "-Wshift-count-overflow"
#endif
#if !defined(__HIPCC_RTC__)
#include "host_defines.h"
#include "math_fwd.h"
#include <hip/hip_runtime_api.h>
#include <stddef.h>
#endif // !defined(__HIPCC_RTC__)
#include <hip/hip_vector_types.h>
#include <hip/amd_detail/device_library_decls.h>
#endif // !defined(__HIPCC_RTC__)
#if __HIP_CLANG_ONLY__
#if defined(__clang__) && defined(__HIP__)
extern "C" __device__ int printf(const char *fmt, ...);
#else
template <typename... All>
@@ -639,7 +653,7 @@ __device__ static inline float __ull2float_rz(unsigned long long int x) {
return __ocml_cvtrtz_f32_u64(x);
}
#if __HIP_CLANG_ONLY__
#if defined(__clang__) && defined(__HIP__)
// Clock functions
__device__ long long int __clock64();
@@ -1105,4 +1119,9 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) {
return __hip_hc_memset(ptr, val8, size);
}
#endif // !__OPENMP_AMDGCN__
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif
hipamd/include/hip/amd_detail/amd_hip_atomic.h
+14
Ver ficheiro
@@ -22,7 +22,17 @@ THE SOFTWARE.
#pragma once
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wc++98-compat"
#pragma clang diagnostic ignored "-Wold-style-cast"
#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
#pragma clang diagnostic ignored "-Wsign-conversion"
#endif
#if !defined(__HIPCC_RTC__)
#include "amd_device_functions.h"
#endif
#if __has_builtin(__hip_atomic_compare_exchange_strong)
@@ -1625,4 +1635,8 @@ unsigned long long atomicXor(
return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
}
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif // __hip_atomic_compare_exchange_strong
hipamd/include/hip/amd_detail/amd_hip_complex.h
+19 -1
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -23,7 +23,21 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wfloat-equal"
#pragma clang diagnostic ignored "-Wc++98-compat"
#pragma clang diagnostic ignored "-Wold-style-cast"
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#pragma clang diagnostic ignored "-Wimplicit-int-float-conversion"
#pragma clang diagnostic ignored "-Wimplicit-float-conversion"
#pragma clang diagnostic ignored "-Wdouble-promotion"
#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
#endif
#if !defined(__HIPCC_RTC__)
#include "hip/amd_detail/amd_hip_vector_types.h"
#endif
#if defined(__HIPCC_RTC__)
#define __HOST_DEVICE__ __device__
@@ -311,4 +325,8 @@ __HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDo
return make_hipDoubleComplex(real, imag);
}
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
hipamd/include/hip/amd_detail/amd_hip_runtime.h
+22 -19
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -110,18 +110,19 @@ using ::int64_t;
#if !defined(__HIPCC_RTC__)
#include <hip/hip_runtime_api.h>
#include <hip/amd_detail/amd_hip_atomic.h>
#include <hip/amd_detail/amd_device_functions.h>
#include <hip/amd_detail/amd_surface_functions.h>
#include <hip/amd_detail/texture_fetch_functions.h>
#include <hip/amd_detail/texture_indirect_functions.h>
extern int HIP_TRACE_API;
#endif // !defined(__HIPCC_RTC__)
#ifdef __cplusplus
#include <hip/amd_detail/hip_ldg.h>
#endif
#include <hip/amd_detail/amd_hip_atomic.h>
#include <hip/amd_detail/host_defines.h>
#include <hip/amd_detail/amd_device_functions.h>
#include <hip/amd_detail/amd_surface_functions.h>
#include <hip/amd_detail/texture_fetch_functions.h>
#include <hip/amd_detail/texture_indirect_functions.h>
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
#if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
@@ -244,10 +245,10 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
#include <hip/hip_runtime_api.h>
#endif // !defined(__HIPCC_RTC__)
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(unsigned int);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(unsigned int);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(unsigned int);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(unsigned int);
struct __HIP_BlockIdx {
__device__
std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_group_id(x); }
@@ -282,7 +283,7 @@ typedef struct dim3 {
} dim3;
#endif // !defined(__HIPCC_RTC__)
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(unsigned int);
#ifdef __cplusplus
template <typename F> struct __HIP_Coordinates {
@@ -367,33 +368,35 @@ static constexpr __HIP_Coordinates<__HIP_GridDim> gridDim{};
static constexpr __HIP_Coordinates<__HIP_ThreadIdx> threadIdx{};
#endif // __cplusplus
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(unsigned int);
#define hipThreadIdx_x (__ockl_get_local_id(0))
#define hipThreadIdx_y (__ockl_get_local_id(1))
#define hipThreadIdx_z (__ockl_get_local_id(2))
extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(unsigned int);
#define hipBlockIdx_x (__ockl_get_group_id(0))
#define hipBlockIdx_y (__ockl_get_group_id(1))
#define hipBlockIdx_z (__ockl_get_group_id(2))
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(unsigned int);
#define hipBlockDim_x (__ockl_get_local_size(0))
#define hipBlockDim_y (__ockl_get_local_size(1))
#define hipBlockDim_z (__ockl_get_local_size(2))
extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(unsigned int);
#define hipGridDim_x (__ockl_get_num_groups(0))
#define hipGridDim_y (__ockl_get_num_groups(1))
#define hipGridDim_z (__ockl_get_num_groups(2))
#if !defined(__HIPCC_RTC__)
#include <hip/amd_detail/amd_math_functions.h>
#endif
#if __HIP_HCC_COMPAT_MODE__
// Define HCC work item functions in terms of HIP builtin variables.
#pragma push_macro("__DEFINE_HCC_FUNC")
#define __DEFINE_HCC_FUNC(hc_fun,hip_var) \
inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \
inline __device__ __attribute__((always_inline)) unsigned int hc_get_##hc_fun(unsigned int i) { \
if (i==0) \
return hip_var.x; \
else if(i==1) \
@@ -408,11 +411,11 @@ __DEFINE_HCC_FUNC(group_size, blockDim)
__DEFINE_HCC_FUNC(num_groups, gridDim)
#pragma pop_macro("__DEFINE_HCC_FUNC")
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(uint);
inline __device__ __attribute__((always_inline)) uint
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(unsigned int);
inline __device__ __attribute__((always_inline)) unsigned int
hc_get_workitem_absolute_id(int dim)
{
return (uint)__ockl_get_global_id(dim);
return (unsigned int)__ockl_get_global_id(dim);
}
#endif
hipamd/include/hip/amd_detail/amd_math_functions.h
+13 -1402
Ver ficheiro
A apresentação das diferenças no ficheiro foi suprimida por ser demasiado grande Carregar diff
hipamd/include/hip/amd_detail/amd_surface_functions.h
+27 -1
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2018 - 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2018 - 2023 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
@@ -23,12 +23,33 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
#define HIP_INCLUDE_HIP_AMD_DETAIL_SURFACE_FUNCTIONS_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#pragma clang diagnostic ignored "-Wc++17-extensions"
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wc++98-compat"
#pragma clang diagnostic ignored "-Wold-style-cast"
#pragma clang diagnostic ignored "-Wextra-semi-stmt"
#pragma clang diagnostic ignored "-Wunused-variable"
#pragma clang diagnostic ignored "-Wunused-parameter"
#pragma clang diagnostic ignored "-Wunused-template"
#endif
#if defined(__cplusplus)
#if !defined(__HIPCC_RTC__)
#include <hip/surface_types.h>
#include <hip/hip_vector_types.h>
#include <hip/amd_detail/texture_fetch_functions.h>
#include <hip/amd_detail/ockl_image.h>
#endif
#if defined(__HIPCC_RTC__)
#define __HOST_DEVICE__ __device__
#else
#define __HOST_DEVICE__ __host__ __device__
#endif
#define __HIP_SURFACE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)surfObj;
@@ -232,4 +253,9 @@ static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfa
}
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif
hipamd/include/hip/amd_detail/device_library_decls.h
+14 -1
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -31,7 +31,16 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
#define HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wc++98-compat-pedantic"
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wold-style-cast"
#endif
#if !defined(__HIPCC_RTC__)
#include "hip/amd_detail/host_defines.h"
#endif
typedef unsigned char uchar;
typedef unsigned short ushort;
@@ -128,4 +137,8 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local
#define __CLK_LOCAL_MEM_FENCE 0x01
typedef unsigned __cl_mem_fence_flags;
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif
hipamd/include/hip/amd_detail/math_fwd.h
+15 -2
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -22,13 +22,22 @@ THE SOFTWARE.
#pragma once
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wundef"
#pragma clang diagnostic ignored "-Wreserved-identifier"
#endif
#if !defined(__HIPCC_RTC__)
#include "host_defines.h"
#endif
#if defined(__cplusplus)
extern "C" {
#endif
// DOT FUNCTIONS
#if __HIP_CLANG_ONLY__
#if defined(__clang__) && defined(__HIP__)
__device__
__attribute__((const))
int __ockl_sdot2(
@@ -692,3 +701,7 @@ double __ocml_fma_rtz_f64(double, double, double);
#if defined(__cplusplus)
} // extern "C"
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
hipamd/include/hip/amd_detail/ockl_image.h
+13 -2
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -22,7 +22,14 @@ THE SOFTWARE.
#pragma once
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/hip_vector_types.h>
#endif
extern "C" {
@@ -172,4 +179,8 @@ __device__ int __ockl_image_channel_order_CM(unsigned int ADDRESS_SPACE_CONSTANT
__device__ int __ockl_image_channel_order_CMa(unsigned int ADDRESS_SPACE_CONSTANT* i);
};
}
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
hipamd/include/hip/amd_detail/texture_fetch_functions.h
+18 -3
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -22,13 +22,24 @@ THE SOFTWARE.
#pragma once
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wc++17-extensions"
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wc++98-compat"
#pragma clang diagnostic ignored "-Wold-style-cast"
#pragma clang diagnostic ignored "-Wextra-semi-stmt"
#pragma clang diagnostic ignored "-Wunused-variable"
#pragma clang diagnostic ignored "-Wunused-parameter"
#pragma clang diagnostic ignored "-Wunused-template"
#endif
#if defined(__cplusplus)
#if !defined(__HIPCC_RTC__)
#include <hip/hip_vector_types.h>
#include <hip/hip_texture_types.h>
#include <hip/amd_detail/ockl_image.h>
#if !defined(__HIPCC_RTC__)
#include <type_traits>
#endif // !defined(__HIPCC_RTC__)
@@ -490,3 +501,7 @@ static __forceinline__ __device__ __hip_img_chk__ __hip_tex2dgather_ret_t<T, rea
}
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
hipamd/include/hip/amd_detail/texture_indirect_functions.h
+18 -4
Ver ficheiro
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2023 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
@@ -22,14 +22,24 @@ THE SOFTWARE.
#pragma once
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wc++17-extensions"
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wc++98-compat"
#pragma clang diagnostic ignored "-Wold-style-cast"
#pragma clang diagnostic ignored "-Wunused-variable"
#pragma clang diagnostic ignored "-Wunused-parameter"
#pragma clang diagnostic ignored "-Wunused-template"
#endif
#if defined(__cplusplus)
#if !defined(__HIPCC_RTC__)
#include <hip/hip_vector_types.h>
#include <hip/hip_texture_types.h>
#include <hip/amd_detail/texture_fetch_functions.h>
#include <hip/amd_detail/ockl_image.h>
#if !defined(__HIPCC_RTC__)
#include <type_traits>
#endif // !defined(__HIPCC_RTC__)
@@ -208,7 +218,7 @@ static __device__ __hip_img_chk__ T tex2Dgather(hipTextureObject_t textureObject
return __hipMapFrom<T>(tmp);
break;
}
};
}
return {};
}
@@ -477,3 +487,7 @@ static __device__ __hip_img_chk__ void texCubemapLayeredGrad(T *ptr, hipTextureO
}
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif