Merge remote-tracking branch 'refs/remotes/origin/privatestaging' into kernel_synchronization
[ROCm/hip commit: f0974e5867]
This commit is contained in:
bwicakso
@@ -1,6 +1,6 @@
|
||||
#ifndef HCC_ACC_H
|
||||
#define HCC_ACC_H
|
||||
#include "hip_runtime_api.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
|
||||
#if __cplusplus
|
||||
#ifdef __HCC__
|
||||
|
||||
@@ -21,8 +21,8 @@ THE SOFTWARE.
|
||||
#define HIP_HCC_H
|
||||
|
||||
#include <hc.hpp>
|
||||
#include "hcc_detail/hip_util.h"
|
||||
#include "hcc_detail/staging_buffer.h"
|
||||
#include "hip/hcc_detail/hip_util.h"
|
||||
#include "hip/hcc_detail/staging_buffer.h"
|
||||
|
||||
|
||||
#if defined(__HCC__) && (__hcc_workweek__ < 1502)
|
||||
|
||||
@@ -23,36 +23,82 @@ THE SOFTWARE.
|
||||
#if __HCC__
|
||||
#include"hip_vector_types.h"
|
||||
#include"host_defines.h"
|
||||
__device__ char __ldg(const char* );
|
||||
__device__ signed char __ldg(const signed char* );
|
||||
__device__ short __ldg(const short* );
|
||||
__device__ int __ldg(const int* );
|
||||
__device__ long __ldg(const long* );
|
||||
__device__ long long __ldg(const long long* );
|
||||
__device__ char2 __ldg(const char2* );
|
||||
__device__ char4 __ldg(const char4* );
|
||||
__device__ short2 __ldg(const short2* );
|
||||
__device__ short4 __ldg(const short4* );
|
||||
__device__ int2 __ldg(const int2* );
|
||||
__device__ int4 __ldg(const int4* );
|
||||
__device__ longlong2 __ldg(const longlong2* );
|
||||
__device__ unsigned char __ldg(const unsigned char* );
|
||||
__device__ unsigned short __ldg(const unsigned short* );
|
||||
__device__ unsigned int __ldg(const unsigned int* );
|
||||
__device__ unsigned long __ldg(const unsigned long* );
|
||||
__device__ unsigned long long __ldg(const unsigned long long* );
|
||||
__device__ uchar2 __ldg(const uchar2* );
|
||||
__device__ uchar4 __ldg(const uchar4* );
|
||||
__device__ ushort2 __ldg(const ushort2* );
|
||||
__device__ uint2 __ldg(const uint2* );
|
||||
__device__ uint4 __ldg(const uint4* );
|
||||
__device__ ulonglong2 __ldg(const ulonglong2* );
|
||||
__device__ float __ldg(const float* );
|
||||
__device__ double __ldg(const double* );
|
||||
__device__ float2 __ldg(const float2* );
|
||||
__device__ float4 __ldg(const float4* );
|
||||
__device__ double2 __ldg(const double2* );
|
||||
#if __hcc_workweek__ >= 16164
|
||||
#include"hip/hip_vector_types.h"
|
||||
#include"hip/hcc_detail/host_defines.h"
|
||||
|
||||
#endif
|
||||
|
||||
#endif
|
||||
__device__ char __ldg(const char* );
|
||||
__device__ char1 __ldg(const char1* );
|
||||
__device__ char2 __ldg(const char2* );
|
||||
__device__ char3 __ldg(const char3* );
|
||||
__device__ char4 __ldg(const char4* );
|
||||
__device__ signed char __ldg(const signed char* );
|
||||
__device__ unsigned char __ldg(const unsigned char* );
|
||||
|
||||
__device__ short __ldg(const short* );
|
||||
__device__ short1 __ldg(const short1* );
|
||||
__device__ short2 __ldg(const short2* );
|
||||
__device__ short3 __ldg(const short3* );
|
||||
__device__ short4 __ldg(const short4* );
|
||||
__device__ unsigned short __ldg(const unsigned short* );
|
||||
|
||||
__device__ int __ldg(const int* );
|
||||
__device__ int1 __ldg(const int1* );
|
||||
__device__ int2 __ldg(const int2* );
|
||||
__device__ int3 __ldg(const int3* );
|
||||
__device__ int4 __ldg(const int4* );
|
||||
__device__ unsigned int __ldg(const unsigned int* );
|
||||
|
||||
|
||||
__device__ long __ldg(const long* );
|
||||
__device__ long1 __ldg(const long1* );
|
||||
__device__ long2 __ldg(const long2* );
|
||||
__device__ long3 __ldg(const long3* );
|
||||
__device__ long4 __ldg(const long4* );
|
||||
__device__ unsigned long __ldg(const unsigned long* );
|
||||
|
||||
__device__ long long __ldg(const long long* );
|
||||
__device__ longlong1 __ldg(const longlong1* );
|
||||
__device__ longlong2 __ldg(const longlong2* );
|
||||
__device__ longlong3 __ldg(const longlong3* );
|
||||
__device__ longlong4 __ldg(const longlong4* );
|
||||
__device__ unsigned long long __ldg(const unsigned long long* );
|
||||
|
||||
__device__ uchar1 __ldg(const uchar1* );
|
||||
__device__ uchar2 __ldg(const uchar2* );
|
||||
__device__ uchar3 __ldg(const uchar3* );
|
||||
__device__ uchar4 __ldg(const uchar4* );
|
||||
|
||||
__device__ ushort1 __ldg(const ushort1* );
|
||||
__device__ ushort2 __ldg(const ushort2* );
|
||||
__device__ ushort3 __ldg(const ushort3* );
|
||||
__device__ ushort4 __ldg(const ushort4* );
|
||||
|
||||
__device__ uint1 __ldg(const uint1* );
|
||||
__device__ uint2 __ldg(const uint2* );
|
||||
__device__ uint3 __ldg(const uint3* );
|
||||
__device__ uint4 __ldg(const uint4* );
|
||||
|
||||
__device__ ulonglong1 __ldg(const ulonglong1* );
|
||||
__device__ ulonglong2 __ldg(const ulonglong2* );
|
||||
__device__ ulonglong3 __ldg(const ulonglong3* );
|
||||
__device__ ulonglong4 __ldg(const ulonglong4* );
|
||||
|
||||
__device__ float __ldg(const float* );
|
||||
__device__ float1 __ldg(const float1* );
|
||||
__device__ float2 __ldg(const float2* );
|
||||
__device__ float3 __ldg(const float3* );
|
||||
__device__ float4 __ldg(const float4* );
|
||||
|
||||
__device__ double __ldg(const double* );
|
||||
__device__ double1 __ldg(const double1* );
|
||||
__device__ double2 __ldg(const double2* );
|
||||
__device__ double3 __ldg(const double3* );
|
||||
__device__ double4 __ldg(const double4* );
|
||||
|
||||
#endif // __hcc_workweek__
|
||||
|
||||
#endif // __HCC__
|
||||
|
||||
#endif // HIP_LDG_H
|
||||
|
||||
@@ -40,8 +40,8 @@ THE SOFTWARE.
|
||||
|
||||
#define CUDA_SUCCESS hipSuccess
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
//#include "hcc_detail/hip_hcc.h"
|
||||
#include <hip/hip_runtime_api.h>
|
||||
//#include "hip/hcc_detail/hip_hcc.h"
|
||||
//---
|
||||
// Remainder of this file only compiles with HCC
|
||||
#ifdef __HCC__
|
||||
@@ -55,10 +55,10 @@ extern int HIP_TRACE_API;
|
||||
//typedef grid_launch_parm hipLaunchParm ;
|
||||
#define hipLaunchParm grid_launch_parm
|
||||
#ifdef __cplusplus
|
||||
#include <hcc_detail/hip_texture.h>
|
||||
#include <hcc_detail/hip_ldg.h>
|
||||
#include <hip/hcc_detail/hip_texture.h>
|
||||
#include <hip/hcc_detail/hip_ldg.h>
|
||||
#endif
|
||||
#include <hcc_detail/host_defines.h>
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
|
||||
#if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
|
||||
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
|
||||
|
||||
@@ -30,11 +30,11 @@ THE SOFTWARE.
|
||||
#include <stdint.h>
|
||||
#include <stddef.h>
|
||||
|
||||
#include <hcc_detail/host_defines.h>
|
||||
#include <hip_runtime_api.h>
|
||||
//#include "hip_hcc.h"
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
//#include "hip/hip_hcc.h"
|
||||
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 16074)
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 16155)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
|
||||
@@ -32,7 +32,7 @@ THE SOFTWARE.
|
||||
|
||||
#include <limits.h>
|
||||
|
||||
#include <hip_runtime.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
//----
|
||||
//Texture - TODO - likely need to move this to a separate file only included with kernel compilation.
|
||||
|
||||
@@ -115,7 +115,74 @@ typedef hc::short_vector::double2 double2;
|
||||
typedef hc::short_vector::double3 double3;
|
||||
typedef hc::short_vector::double4 double4;
|
||||
|
||||
#if __HCC__
|
||||
#include"hip/hcc_detail/host_defines.h"
|
||||
#define __HIP_DEVICE__ __device__ __host__
|
||||
#else
|
||||
#define __HIP_DEVICE__
|
||||
#endif
|
||||
|
||||
__HIP_DEVICE__ char1 make_char1(signed char );
|
||||
__HIP_DEVICE__ char2 make_char2(signed char, signed char );
|
||||
__HIP_DEVICE__ char3 make_char3(signed char, signed char, signed char );
|
||||
__HIP_DEVICE__ char4 make_char4(signed char, signed char, signed char, signed char );
|
||||
|
||||
__HIP_DEVICE__ short1 make_short1(short );
|
||||
__HIP_DEVICE__ short2 make_short2(short, short );
|
||||
__HIP_DEVICE__ short3 make_short3(short, short, short );
|
||||
__HIP_DEVICE__ short4 make_short4(short, short, short, short );
|
||||
|
||||
__HIP_DEVICE__ int1 make_int1(int );
|
||||
__HIP_DEVICE__ int2 make_int2(int, int );
|
||||
__HIP_DEVICE__ int3 make_int3(int, int, int );
|
||||
__HIP_DEVICE__ int4 make_int4(int, int, int, int );
|
||||
|
||||
__HIP_DEVICE__ long1 make_long1(long );
|
||||
__HIP_DEVICE__ long2 make_long2(long, long );
|
||||
__HIP_DEVICE__ long3 make_long3(long, long, long );
|
||||
__HIP_DEVICE__ long4 make_long4(long, long, long, long );
|
||||
|
||||
__HIP_DEVICE__ longlong1 make_longlong1(long long );
|
||||
__HIP_DEVICE__ longlong2 make_longlong2(long long, long long );
|
||||
__HIP_DEVICE__ longlong3 make_longlong3(long long, long long, long long );
|
||||
__HIP_DEVICE__ longlong4 make_longlong4(long long, long long, long long, long long );
|
||||
|
||||
__HIP_DEVICE__ uchar1 make_uchar1(unsigned char );
|
||||
__HIP_DEVICE__ uchar2 make_uchar2(unsigned char, unsigned char );
|
||||
__HIP_DEVICE__ uchar3 make_uchar3(unsigned char, unsigned char, unsigned char );
|
||||
__HIP_DEVICE__ uchar4 make_uchar4(unsigned char, unsigned char, unsigned char, unsigned char );
|
||||
|
||||
__HIP_DEVICE__ ushort1 make_ushort1(unsigned short );
|
||||
__HIP_DEVICE__ ushort2 make_ushort2(unsigned short, unsigned short );
|
||||
__HIP_DEVICE__ ushort3 make_ushort3(unsigned short, unsigned short, unsigned short );
|
||||
__HIP_DEVICE__ ushort4 make_ushort4(unsigned short, unsigned short, unsigned short, unsigned short );
|
||||
|
||||
__HIP_DEVICE__ uint1 make_uint1(unsigned int );
|
||||
__HIP_DEVICE__ uint2 make_uint2(unsigned int, unsigned int );
|
||||
__HIP_DEVICE__ uint3 make_uint3(unsigned int, unsigned int, unsigned int );
|
||||
__HIP_DEVICE__ uint4 make_uint4(unsigned int, unsigned int, unsigned int, unsigned int );
|
||||
|
||||
__HIP_DEVICE__ ulong1 make_ulong1(unsigned long );
|
||||
__HIP_DEVICE__ ulong2 make_ulong2(unsigned long, unsigned long );
|
||||
__HIP_DEVICE__ ulong3 make_ulong3(unsigned long, unsigned long, unsigned long );
|
||||
__HIP_DEVICE__ ulong4 make_ulong4(unsigned long, unsigned long, unsigned long, unsigned long );
|
||||
|
||||
__HIP_DEVICE__ ulonglong1 make_ulonglong1(unsigned long long );
|
||||
__HIP_DEVICE__ ulonglong2 make_ulonglong2(unsigned long long, unsigned long long);
|
||||
__HIP_DEVICE__ ulonglong3 make_ulonglong3(unsigned long long, unsigned long long, unsigned long long);
|
||||
__HIP_DEVICE__ ulonglong4 make_ulonglong4(unsigned long long, unsigned long long, unsigned long long, unsigned long long );
|
||||
|
||||
__HIP_DEVICE__ float1 make_float1(float );
|
||||
__HIP_DEVICE__ float2 make_float2(float, float );
|
||||
__HIP_DEVICE__ float3 make_float3(float, float, float );
|
||||
__HIP_DEVICE__ float4 make_float4(float, float, float, float );
|
||||
|
||||
__HIP_DEVICE__ double1 make_double1(double );
|
||||
__HIP_DEVICE__ double2 make_double2(double, double );
|
||||
__HIP_DEVICE__ double3 make_double3(double, double, double );
|
||||
__HIP_DEVICE__ double4 make_double4(double, double, double, double );
|
||||
|
||||
/*
|
||||
///---
|
||||
// Inline functions for creating vector types from basic types
|
||||
#define ONE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x) { VT t; t.x = x; return t; };
|
||||
@@ -198,7 +265,7 @@ ONE_COMPONENT_ACCESS (double, double1);
|
||||
TWO_COMPONENT_ACCESS (double, double2);
|
||||
THREE_COMPONENT_ACCESS(double, double3);
|
||||
FOUR_COMPONENT_ACCESS (double, double4);
|
||||
|
||||
*/
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
@@ -48,17 +48,17 @@ THE SOFTWARE.
|
||||
#endif
|
||||
|
||||
|
||||
#include <hip_common.h>
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hcc_detail/hip_runtime.h>
|
||||
#include <hip/hcc_detail/hip_runtime.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include <nvcc_detail/hip_runtime.h>
|
||||
#include <hip/nvcc_detail/hip_runtime.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
#include <hip_vector_types.h>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
|
||||
|
||||
@@ -30,7 +30,7 @@ THE SOFTWARE.
|
||||
|
||||
|
||||
#include <string.h> // for getDeviceProp
|
||||
#include <hip_common.h>
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
typedef struct {
|
||||
// 32-bit Atomics
|
||||
@@ -202,9 +202,9 @@ typedef enum hipDeviceAttribute_t {
|
||||
*/
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include "hcc_detail/hip_runtime_api.h"
|
||||
#include "hip/hcc_detail/hip_runtime_api.h"
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include "nvcc_detail/hip_runtime_api.h"
|
||||
#include "hip/nvcc_detail/hip_runtime_api.h"
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
@@ -23,12 +23,12 @@ THE SOFTWARE.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip_common.h>
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#if __cplusplus
|
||||
#include <hcc_detail/hip_vector_types.h>
|
||||
#include <hip/hcc_detail/hip_vector_types.h>
|
||||
#endif
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include <vector_types.h>
|
||||
|
||||
@@ -23,7 +23,7 @@ THE SOFTWARE.
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
|
||||
|
||||
Reference in New Issue
Block a user