From 167bbffc4fbcece2a3619061d477693891d911dd Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 9 Jan 2019 09:32:50 -0800 Subject: [PATCH] Fix compilation issues with MSVC --- hipamd/include/hip/hcc_detail/driver_types.h | 6 +- .../include/hip/hcc_detail/hip_runtime_api.h | 64 +++++++---- .../include/hip/hcc_detail/hip_vector_types.h | 107 ++++++++++++++++++ .../hip/hcc_detail/surface_functions.h | 2 +- .../hip/hcc_detail/texture_functions.h | 2 +- 5 files changed, 155 insertions(+), 26 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/driver_types.h b/hipamd/include/hip/hcc_detail/driver_types.h index 661af64cd0..8e1fec11fa 100644 --- a/hipamd/include/hip/hcc_detail/driver_types.h +++ b/hipamd/include/hip/hcc_detail/driver_types.h @@ -279,7 +279,7 @@ typedef struct hipMemcpy3DParms { size_t srcZ; }hipMemcpy3DParms; -static __inline__ struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, +static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz) { struct hipPitchedPtr s; @@ -291,7 +291,7 @@ static __inline__ struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, siz return s; } -static __inline__ struct hipPos make_hipPos(size_t x, size_t y, size_t z) { +static inline struct hipPos make_hipPos(size_t x, size_t y, size_t z) { struct hipPos p; p.x = x; @@ -301,7 +301,7 @@ static __inline__ struct hipPos make_hipPos(size_t x, size_t y, size_t z) { return p; } -static __inline__ struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) { +static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) { struct hipExtent e; e.width = w; diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 0d526ca25d..b6ae88729a 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -41,7 +41,12 @@ THE SOFTWARE. #include #include +#if defined(_MSC_VER) +#define DEPRECATED(msg) __declspec(deprecated(msg)) +#else // !defined(_MSC_VER) #define DEPRECATED(msg) __attribute__ ((deprecated(msg))) +#endif // !defined(_MSC_VER) + #define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases.For more details please refer https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list" #if defined(__HCC__) && (__hcc_workweek__ < 16155) @@ -1044,8 +1049,8 @@ hipError_t hipMalloc(void** ptr, size_t size); * * @deprecated use hipHostMalloc() instead */ -hipError_t hipMallocHost(void** ptr, size_t size) - __attribute__((deprecated("use hipHostMalloc instead"))); +DEPRECATED("use hipHostMalloc instead") +hipError_t hipMallocHost(void** ptr, size_t size); /** * @brief Allocate device accessible page locked host memory @@ -1075,8 +1080,8 @@ hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags); * * @deprecated use hipHostMalloc() instead */ -hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) - __attribute__((deprecated("use hipHostMalloc instead"))); +DEPRECATED("use hipHostMalloc instead") +hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags); /** * @brief Get Device pointer from Host Pointer allocated through hipHostMalloc @@ -1196,7 +1201,8 @@ hipError_t hipFree(void* ptr); * @deprecated use hipHostFree() instead */ -hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree instead"))); +DEPRECATED("use hipHostFree instead") +hipError_t hipFreeHost(void* ptr); /** * @brief Free memory allocated by the hcc hip host memory allocation API @@ -1919,7 +1925,8 @@ hipError_t hipInit(unsigned int flags); * @see hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, * hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device)DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device); /** * @brief Destroy a HIP context. @@ -1931,7 +1938,8 @@ hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device)DE * @see hipCtxCreate, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent,hipCtxSetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -hipError_t hipCtxDestroy(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxDestroy(hipCtx_t ctx); /** * @brief Pop the current/default context and return the popped context. @@ -1943,7 +1951,8 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxSetCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxPopCurrent(hipCtx_t* ctx) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxPopCurrent(hipCtx_t* ctx); /** * @brief Push the context to be set as current/ default context @@ -1955,7 +1964,8 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) DEPRECATED(DEPRECATED_MSG); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -hipError_t hipCtxPushCurrent(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxPushCurrent(hipCtx_t ctx); /** * @brief Set the passed context as current/default @@ -1967,7 +1977,8 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -hipError_t hipCtxSetCurrent(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxSetCurrent(hipCtx_t ctx); /** * @brief Get the handle of the current/ default context @@ -1979,7 +1990,8 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetCurrent(hipCtx_t* ctx) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxGetCurrent(hipCtx_t* ctx); /** * @brief Get the handle of the device associated with current/default context @@ -1992,7 +2004,8 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) DEPRECATED(DEPRECATED_MSG); * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize */ -hipError_t hipCtxGetDevice(hipDevice_t* device) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxGetDevice(hipDevice_t* device); /** * @brief Returns the approximate HIP api version. @@ -2011,7 +2024,8 @@ hipError_t hipCtxGetDevice(hipDevice_t* device) DEPRECATED(DEPRECATED_MSG); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion); /** * @brief Set Cache configuration for a specific function @@ -2026,7 +2040,8 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) DEPRECATED(DEPRECA * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig); /** * @brief Set L1/Shared cache partition. @@ -2041,7 +2056,8 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) DEPRECATED(DEPRECAT * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig); /** * @brief Set Shared memory bank configuration. @@ -2056,7 +2072,8 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) DEPRECATED(DEPRECATE * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config); /** * @brief Get Shared memory bank configuration. @@ -2071,7 +2088,8 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) DEPRECATED(DEPREC * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig); /** * @brief Blocks until the default context has completed all preceding requested tasks. @@ -2084,7 +2102,8 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) DEPRECATED(DEPR * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxGetDevice */ -hipError_t hipCtxSynchronize(void) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxSynchronize(void); /** * @brief Return flags used for creating default context. @@ -2096,7 +2115,8 @@ hipError_t hipCtxSynchronize(void) DEPRECATED(DEPRECATED_MSG); * @see hipCtxCreate, hipCtxDestroy, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetFlags(unsigned int* flags) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxGetFlags(unsigned int* flags); /** * @brief Enables direct access to memory allocations in a peer context. @@ -2117,7 +2137,8 @@ hipError_t hipCtxGetFlags(unsigned int* flags) DEPRECATED(DEPRECATED_MSG); * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice * @warning PeerToPeer support is experimental. */ -hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags); /** * @brief Disable direct access from current context's virtual address space to memory allocations @@ -2135,7 +2156,8 @@ hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) DEPRECAT * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice * @warning PeerToPeer support is experimental. */ -hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) DEPRECATED(DEPRECATED_MSG); +DEPRECATED(DEPRECATED_MSG) +hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx); /** * @brief Get the state of the primary context. diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 00537e3e61..924498583d 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -34,6 +34,7 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" +#if !defined(_MSC_VER) #if defined(__clang__) #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) #elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax. @@ -769,5 +770,111 @@ DECLOP_MAKE_ONE_COMPONENT(signed long long, longlong1); DECLOP_MAKE_TWO_COMPONENT(signed long long, longlong2); DECLOP_MAKE_THREE_COMPONENT(signed long long, longlong3); DECLOP_MAKE_FOUR_COMPONENT(signed long long, longlong4); +#else // defined(_MSC_VER) +#include +#include +#include +#include +typedef union { char data; } char1; +typedef union { char data[2]; } char2; +typedef union { char data[4]; } char4; +typedef union { char4 data; } char3; +typedef union { __m64 data; } char8; +typedef union { __m128i data; } char16; + +typedef union { unsigned char data; } uchar1; +typedef union { unsigned char data[2]; } uchar2; +typedef union { unsigned char data[4]; } uchar4; +typedef union { uchar4 data; } uchar3; +typedef union { __m64 data; } uchar8; +typedef union { __m128i data; } uchar16; + +typedef union { short data; } short1; +typedef union { short data[2]; } short2; +typedef union { __m64 data; } short4; +typedef union { short4 data; } short3; +typedef union { __m128i data; } short8; +typedef union { __m128i data[2]; } short16; + +typedef union { unsigned short data; } ushort1; +typedef union { unsigned short data[2]; } ushort2; +typedef union { __m64 data; } ushort4; +typedef union { ushort4 data; } ushort3; +typedef union { __m128i data; } ushort8; +typedef union { __m128i data[2]; } ushort16; + +typedef union { int data; } int1; +typedef union { __m64 data; } int2; +typedef union { __m128i data; } int4; +typedef union { int4 data; } int3; +typedef union { __m128i data[2]; } int8; +typedef union { __m128i data[4];} int16; + +typedef union { unsigned int data; } uint1; +typedef union { __m64 data; } uint2; +typedef union { __m128i data; } uint4; +typedef union { uint4 data; } uint3; +typedef union { __m128i data[2]; } uint8; +typedef union { __m128i data[4]; } uint16; + +#if !defined(_WIN64) +typedef union { int data; } long1; +typedef union { __m64 data; } long2; +typedef union { __m128i data; } long4; +typedef union { long4 data; } long3; +typedef union { __m128i data[2]; } long8; +typedef union { __m128i data[4]; } long16; + +typedef union { unsigned int data; } ulong1; +typedef union { __m64 data; } ulong2; +typedef union { __m128i data; } ulong4; +typedef union { ulong4 data; } ulong3; +typedef union { __m128i data[2]; } ulong8; +typedef union { __m128i data[4]; } ulong16; +#else // defined(_WIN64) +typedef union { __m64 data; } long1; +typedef union { __m128i data; } long2; +typedef union { __m128i data[2]; } long4; +typedef union { long4 data; } long3; +typedef union { __m128i data[4]; } long8; +typedef union { __m128i data[8]; } long16; + +typedef union { __m64 data; } ulong1; +typedef union { __m128i data; } ulong2; +typedef union { __m128i data[2]; } ulong4; +typedef union { ulong4 data; } ulong3; +typedef union { __m128i data[4]; } ulong8; +typedef union { __m128i data[8]; } ulong16; +#endif // defined(_WIN64) + +typedef union { __m64 data; } longlong1; +typedef union { __m128i data; } longlong2; +typedef union { __m128i data[2]; } longlong4; +typedef union { longlong4 data; } longlong3; +typedef union { __m128i data[4]; } longlong8; +typedef union { __m128i data[8]; } longlong16; + +typedef union { __m64 data; } ulonglong1; +typedef union { __m128i data; } ulonglong2; +typedef union { __m128i data[2]; } ulonglong4; +typedef union { ulonglong4 data; } ulonglong3; +typedef union { __m128i data[4]; } ulonglong8; +typedef union { __m128i data[8]; } ulonglong16; + +typedef union { float data; } float1; +typedef union { __m64 data; } float2; +typedef union { __m128 data; } float4; +typedef union { float4 data; } float3; +typedef union { __m256 data; } float8; +typedef union { __m256 data[2]; } float16; + +typedef union { double data; } double1; +typedef union { __m128d data; } double2; +typedef union { __m256d data; } double4; +typedef union { double4 data; } double3; +typedef union { __m256d data[2]; } double8; +typedef union { __m256d data[4]; } double16; + +#endif // defined(_MSC_VER) #endif diff --git a/hipamd/include/hip/hcc_detail/surface_functions.h b/hipamd/include/hip/hcc_detail/surface_functions.h index 607f221901..b9cab1f466 100644 --- a/hipamd/include/hip/hcc_detail/surface_functions.h +++ b/hipamd/include/hip/hcc_detail/surface_functions.h @@ -25,7 +25,7 @@ THE SOFTWARE. #include -#define __SURFACE_FUNCTIONS_DECL__ static __inline__ __device__ +#define __SURFACE_FUNCTIONS_DECL__ static inline __device__ template __SURFACE_FUNCTIONS_DECL__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipBoundaryModeZero) { diff --git a/hipamd/include/hip/hcc_detail/texture_functions.h b/hipamd/include/hip/hcc_detail/texture_functions.h index 2a7aeb7357..c293d65558 100644 --- a/hipamd/include/hip/hcc_detail/texture_functions.h +++ b/hipamd/include/hip/hcc_detail/texture_functions.h @@ -47,7 +47,7 @@ union TData { __hip_uint4_vector_value_type u; }; -#define __TEXTURE_FUNCTIONS_DECL__ static __inline__ __device__ +#define __TEXTURE_FUNCTIONS_DECL__ static inline __device__ #if (__hcc_workweek__ >= 18114)