AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Merge 'master' into 'amd-master'

Przeglądaj źródła 
Change-Id: Ie46e5c47b00abc773e5f5c03188852aaee32aa74
This commit is contained in:
Jenkins
2019-01-11 04:09:51 -06:00
rodzic 7f9d78f7a1 eb618d2140
commit a4cbc04802
53 zmienionych plików z 236 dodań i 81 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
README.md
+1 -1
Wyświetl plik
@@ -35,7 +35,7 @@ HIP releases are typically of two types. The tag naming convention is different
- [HIP Programming Guide](docs/markdown/hip_programming_guide.md)
- [HIP Profiling ](docs/markdown/hip_profiling.md)
- [HIP Debugging](docs/markdown/hip_debugging.md)
- [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenCL)
- [hipify-clang](hipify-clang/README.md)
- [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
- [Release Notes](RELEASE.md)
hipify-clang/README.md
+2 -2
Wyświetl plik
@@ -320,9 +320,9 @@ For example:
```shell
./hipify-clang \
square.cu \
-- \
--cuda-path=/usr/local/cuda-8.0 \
-isystem /usr/local/cuda-8.0/samples/common/inc
-- \
-I /usr/local/cuda-8.0/samples/common/inc
```
`hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
hipify-clang/src/ArgParse.cpp
+10
Wyświetl plik
@@ -39,6 +39,11 @@ cl::opt<std::string> TemporaryDir("temp-dir",
cl::value_desc("directory"),
cl::cat(ToolTemplateCategory));
cl::opt<std::string> CudaPath("cuda-path",
cl::desc("CUDA installation path"),
cl::value_desc("directory"),
cl::cat(ToolTemplateCategory));
cl::opt <bool> SaveTemps("save-temps",
cl::desc("Save temporary files"),
cl::value_desc("save-temps"),
@@ -79,4 +84,9 @@ cl::opt<bool> Examine("examine",
cl::value_desc("examine"),
cl::cat(ToolTemplateCategory));
cl::opt<bool> DashDash("-",
cl::desc("Separator between hipify-clang and clang options;\ndon't specify if there are no clang options."),
cl::value_desc("--"),
cl::cat(ToolTemplateCategory));
cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage);
hipify-clang/src/ArgParse.h
+2
Wyświetl plik
@@ -32,6 +32,7 @@ extern cl::OptionCategory ToolTemplateCategory;
extern cl::opt<std::string> OutputFilename;
extern cl::opt<std::string> OutputDir;
extern cl::opt<std::string> TemporaryDir;
extern cl::opt<std::string> CudaPath;
extern cl::opt<bool> Inplace;
extern cl::opt<bool> SaveTemps;
extern cl::opt<bool> NoBackup;
@@ -41,3 +42,4 @@ extern cl::opt<std::string> OutputStatsFilename;
extern cl::opt<bool> Examine;
extern cl::extrahelp CommonHelp;
extern cl::opt<bool> TranslateToRoc;
extern cl::opt<bool> DashDash;
hipify-clang/src/main.cpp
+11
Wyświetl plik
@@ -68,6 +68,13 @@ std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC,
}
int main(int argc, const char **argv) {
std::vector<const char*> new_argv(argv, argv + argc);
if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) {
new_argv.push_back("--");
new_argv.push_back(nullptr);
argv = new_argv.data();
argc++;
}
llcompat::PrintStackTraceOnErrorSignal();
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
@@ -174,6 +181,10 @@ int main(int argc, const char **argv) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
if (!CudaPath.empty()) {
std::string sCudaPath = "--cuda-path=" + CudaPath;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
// Includes for clang's CUDA wrappers for using by packaged hipify-clang
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
include/hip/hcc_detail/driver_types.h
+3 -3
Wyświetl plik
@@ -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;
include/hip/hcc_detail/hip_runtime_api.h
+43 -21
Wyświetl plik
@@ -41,7 +41,12 @@ THE SOFTWARE.
#include <hip/hcc_detail/hip_texture_types.h>
#include <hip/hcc_detail/hip_surface_types.h>
#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.
include/hip/hcc_detail/hip_vector_types.h
+107
Wyświetl plik
@@ -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 <mmintrin.h>
#include <xmmintrin.h>
#include <emmintrin.h>
#include <immintrin.h>
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
include/hip/hcc_detail/surface_functions.h
+1 -1
Wyświetl plik
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include <hip/hcc_detail/hip_surface_types.h>
#define __SURFACE_FUNCTIONS_DECL__ static __inline__ __device__
#define __SURFACE_FUNCTIONS_DECL__ static inline __device__
template <class T>
__SURFACE_FUNCTIONS_DECL__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y,
int boundaryMode = hipBoundaryModeZero) {
include/hip/hcc_detail/texture_functions.h
+1 -1
Wyświetl plik
@@ -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)
src/program_state.cpp
+1 -1
Wyświetl plik
@@ -650,7 +650,7 @@ extern "C" void __attribute__((constructor)) __startup_kernel_loader_init() {
int hip_startup_loader=0;
if (std::getenv("HIP_STARTUP_LOADER"))
hip_startup_loader = atoi(std::getenv("HIP_STARTUP_LOADER"));
if (hip_startup_loader) functions();
if (hip_startup_loader) functions(true);
}
extern "C" void __attribute__((destructor)) __startup_kernel_loader_fini() {
tests/hipify-clang/lit.cfg
+7 -5
Wyświetl plik
@@ -64,18 +64,20 @@ if obj_root is not None:
config.environment['PATH'] = path
hipify_path = obj_root
clang_args = "-v --cuda-path='%s'"
clang_arguments = "-v"
if sys.platform in ['win32']:
run_test_ext = ".bat"
hipify_path += "/" + config.build_type
clang_args += " -isystem'%s'/common/inc"
clang_arguments += " -isystem'%s'/common/inc"
else:
run_test_ext = ".sh"
clang_args += " -isystem'%s'/samples/common/inc"
clang_arguments += " -isystem'%s'/samples/common/inc"
clang_arguments += " -I'%s'/include"
clang_args += " -I'%s'/include"
hipify_arguments = "--cuda-path='%s'"
config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root, config.cuda_dnn_root)))
config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root)))
config.substitutions.append(("%hipify_args", hipify_arguments % (config.cuda_root)))
config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"'))
config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"'))
tests/hipify-clang/run_test.bat
+5 -4
Wyświetl plik
@@ -2,17 +2,18 @@
setlocal
for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i
if not defined FILE_CHECK (echo Error: FileCheck.exe not found in PATH. && exit /b 1)
if not defined FILE_CHECK (echo Error: FileCheck.exe not found in PATH. && exit /b 1)
set HIPIFY=%1
set IN_FILE=%2
set TMP_FILE=%3
set CUDA_ROOT=%4
set all_args=%*
call set clang_args=%%all_args:*%4=%%
set clang_args=%4%clang_args%
call set clang_args=%%all_args:*%5=%%
set clang_args=%5%clang_args%
%HIPIFY% -o=%TMP_FILE% %IN_FILE% -- %clang_args%
%HIPIFY% -o=%TMP_FILE% %IN_FILE% %CUDA_ROOT% -- %clang_args%
if errorlevel 1 (echo Error: hipify-clang.exe failed with exit code: %errorlevel% && exit /b %errorlevel%)
findstr /v /r /c:"[ ]*//[ ]*[CHECK*|RUN]" %TMP_FILE% | %FILE_CHECK% %IN_FILE%
tests/hipify-clang/run_test.sh
+3 -3
Wyświetl plik
@@ -9,9 +9,9 @@ set -o errexit
HIPIFY=$1
IN_FILE=$2
TMP_FILE=$3
shift 3
CUDA_ROOT=$4
shift 4
# Remaining args are the ones to forward to clang proper.
$HIPIFY -o=$TMP_FILE $IN_FILE -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE
$HIPIFY -o=$TMP_FILE $IN_FILE $CUDA_ROOT -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE
tests/hipify-clang/unit_tests/headers/headers_test_01.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK-NOT: #include <cuda_runtime.h>
tests/hipify-clang/unit_tests/headers/headers_test_02.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include "hip/hip_runtime.h"
// CHECK-NOT: #include "cuda_runtime.h"
tests/hipify-clang/unit_tests/headers/headers_test_03.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #pragma once
// CHECK-NEXT: #include <hip/hip_runtime.h>
tests/hipify-clang/unit_tests/headers/headers_test_04.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK-NEXT: #include <stdio.h>
tests/hipify-clang/unit_tests/headers/headers_test_05.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #pragma once
// CHECK-NEXT: #include <hip/hip_runtime.h>
tests/hipify-clang/unit_tests/headers/headers_test_06.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hipblas.h>
// CHECK-NOT: #include <cublas_v2.h>
tests/hipify-clang/unit_tests/headers/headers_test_07.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include "hipblas.h"
// CHECK-NOT: #include "cublas.h"
tests/hipify-clang/unit_tests/headers/headers_test_08.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK-NOT: #include <cuda_runtime.h>
tests/hipify-clang/unit_tests/headers/headers_test_09.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK: #include <memory>
tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK: #include "hip/hip_complex.h"
tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <iomanip>
#include <iostream>
tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <stdio.h>
tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <cuda.h>
tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Copyright (c) 2017 Advanced Micro Devices, Inc. All rights reserved.
//
tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Copyright (c) 2017 Advanced Micro Devices, Inc. All rights reserved.
//
tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Taken from: http://docs.nvidia.com/cuda/curand/device-api-overview.html#poisson-api-example
/*
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
// CHECK: #include <hip/hip_runtime.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/samples/allocators.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#pragma once
tests/hipify-clang/unit_tests/samples/axpy.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <iostream>
tests/hipify-clang/unit_tests/samples/coalescing.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// To measure effects of memory coalescing. Coalescing.cu
// B. Wilkinson Jan 30, 2011
tests/hipify-clang/unit_tests/samples/cudaRegister.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Taken from Jonathan Hui blog https://jhui.github.io/2017/03/06/CUDA
tests/hipify-clang/unit_tests/samples/intro.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
#include <stdio.h>
#include <stdlib.h>
tests/hipify-clang/unit_tests/samples/square.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
tests/hipify-clang/unit_tests/samples/static_shared_memory.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Taken from Jonathan Hui blog https://jhui.github.io/2017/03/06/CUDA
tests/hipify-clang/unit_tests/samples/vec_add.cu
+1 -1
Wyświetl plik
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Kernel definition
__global__ void vecAdd(float* A, float* B, float* C)