diff --git a/README.md b/README.md index 4f7f1a0123..200b2f7474 100644 --- a/README.md +++ b/README.md @@ -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) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 8ef9883b55..812500ad34 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -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 diff --git a/hipify-clang/src/ArgParse.cpp b/hipify-clang/src/ArgParse.cpp index e5915aedb9..9696cb74aa 100644 --- a/hipify-clang/src/ArgParse.cpp +++ b/hipify-clang/src/ArgParse.cpp @@ -39,6 +39,11 @@ cl::opt TemporaryDir("temp-dir", cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); +cl::opt CudaPath("cuda-path", + cl::desc("CUDA installation path"), + cl::value_desc("directory"), + cl::cat(ToolTemplateCategory)); + cl::opt SaveTemps("save-temps", cl::desc("Save temporary files"), cl::value_desc("save-temps"), @@ -79,4 +84,9 @@ cl::opt Examine("examine", cl::value_desc("examine"), cl::cat(ToolTemplateCategory)); +cl::opt 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); diff --git a/hipify-clang/src/ArgParse.h b/hipify-clang/src/ArgParse.h index 70a6a9b28d..97f172d812 100644 --- a/hipify-clang/src/ArgParse.h +++ b/hipify-clang/src/ArgParse.h @@ -32,6 +32,7 @@ extern cl::OptionCategory ToolTemplateCategory; extern cl::opt OutputFilename; extern cl::opt OutputDir; extern cl::opt TemporaryDir; +extern cl::opt CudaPath; extern cl::opt Inplace; extern cl::opt SaveTemps; extern cl::opt NoBackup; @@ -41,3 +42,4 @@ extern cl::opt OutputStatsFilename; extern cl::opt Examine; extern cl::extrahelp CommonHelp; extern cl::opt TranslateToRoc; +extern cl::opt DashDash; diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index bd54975e59..6df8ae9679 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -68,6 +68,13 @@ std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC, } int main(int argc, const char **argv) { + std::vector 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 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)); diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h index 661af64cd0..8e1fec11fa 100644 --- a/include/hip/hcc_detail/driver_types.h +++ b/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/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 0d526ca25d..b6ae88729a 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 00537e3e61..924498583d 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/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/include/hip/hcc_detail/surface_functions.h b/include/hip/hcc_detail/surface_functions.h index 607f221901..b9cab1f466 100644 --- a/include/hip/hcc_detail/surface_functions.h +++ b/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/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index 2a7aeb7357..c293d65558 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/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) diff --git a/src/program_state.cpp b/src/program_state.cpp index bb906b0ad9..601096b48f 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -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() { diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 76a811031a..32a3863f2a 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -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 + '"')) diff --git a/tests/hipify-clang/run_test.bat b/tests/hipify-clang/run_test.bat index d8c8d74cf0..58563be2c7 100644 --- a/tests/hipify-clang/run_test.bat +++ b/tests/hipify-clang/run_test.bat @@ -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% diff --git a/tests/hipify-clang/run_test.sh b/tests/hipify-clang/run_test.sh index 418df5dd4d..eb3e822549 100755 --- a/tests/hipify-clang/run_test.sh +++ b/tests/hipify-clang/run_test.sh @@ -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 diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_01.cu b/tests/hipify-clang/unit_tests/headers/headers_test_01.cu index 3747c339e8..013d7c17c4 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_01.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_01.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_02.cu b/tests/hipify-clang/unit_tests/headers/headers_test_02.cu index 57308efd59..957fd16559 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_02.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_02.cu @@ -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" diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_03.cu b/tests/hipify-clang/unit_tests/headers/headers_test_03.cu index 5f2e479683..14735172fb 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_03.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_03.cu @@ -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 diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_04.cu b/tests/hipify-clang/unit_tests/headers/headers_test_04.cu index 57667b5a34..10a7daf41c 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_04.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_04.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NEXT: #include diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_05.cu b/tests/hipify-clang/unit_tests/headers/headers_test_05.cu index c9428b62d5..4706044b9d 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_05.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_05.cu @@ -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 diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_06.cu b/tests/hipify-clang/unit_tests/headers/headers_test_06.cu index bce73c42df..1adccd95e4 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_06.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_06.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_07.cu b/tests/hipify-clang/unit_tests/headers/headers_test_07.cu index 4237e1eb72..1effc189b8 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_07.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_07.cu @@ -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" diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_08.cu b/tests/hipify-clang/unit_tests/headers/headers_test_08.cu index ad54871bd8..aca7f194b0 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_08.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_08.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK-NOT: #include diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_09.cu b/tests/hipify-clang/unit_tests/headers/headers_test_09.cu index 13d2b688f9..37e718b5a4 100644 --- a/tests/hipify-clang/unit_tests/headers/headers_test_09.cu +++ b/tests/hipify-clang/unit_tests/headers/headers_test_09.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK: #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu b/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu index 5dcaaaf637..af24737dbc 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu b/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu index 4e3c1cace8..6983140eac 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu b/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu index 32692440d4..ecd8fb7eee 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu b/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu index 8bf9587b94..196a94197b 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include // CHECK: #include "hip/hip_complex.h" diff --git a/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu b/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu index e58116e22a..6aae90e403 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu b/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu index 2356080e4d..a2052b3de6 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_softmax.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu b/tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu index 906fe93853..9c05a53fa8 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuFFT/simple_cufft.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args // CHECK: #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp index af30969c6f..9fb77dcf3e 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp +++ b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp @@ -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. // diff --git a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp index 222e30570a..4442fb634c 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp +++ b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp @@ -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. // diff --git a/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu b/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu index d4cfd90e1f..4e748973d5 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu @@ -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 /* diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu index df2499c041..6e163d47a6 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include // CHECK: #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu index de7629367f..57b2c61098 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu index ef52072576..349d6471db 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu index 32760f4fa7..18ba9b006c 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu index c6e7374d0b..1155af09e1 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu index d38dcd98e4..872750a5d5 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu index 371f836a5f..cc938d7da4 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu index cf788c8f33..2c826935ef 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu index 9b1aaf623f..3bcbd96bb0 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu index 326c231de4..6690a7c296 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu index 2d905bcf6d..868f3be69e 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu +++ b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include #include diff --git a/tests/hipify-clang/unit_tests/samples/allocators.cu b/tests/hipify-clang/unit_tests/samples/allocators.cu index 3f130be227..a1bce9d4cc 100644 --- a/tests/hipify-clang/unit_tests/samples/allocators.cu +++ b/tests/hipify-clang/unit_tests/samples/allocators.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #pragma once diff --git a/tests/hipify-clang/unit_tests/samples/axpy.cu b/tests/hipify-clang/unit_tests/samples/axpy.cu index 2e59fc021a..6bd1065c83 100644 --- a/tests/hipify-clang/unit_tests/samples/axpy.cu +++ b/tests/hipify-clang/unit_tests/samples/axpy.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include diff --git a/tests/hipify-clang/unit_tests/samples/coalescing.cu b/tests/hipify-clang/unit_tests/samples/coalescing.cu index 4c04289044..ec4645d673 100644 --- a/tests/hipify-clang/unit_tests/samples/coalescing.cu +++ b/tests/hipify-clang/unit_tests/samples/coalescing.cu @@ -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 diff --git a/tests/hipify-clang/unit_tests/samples/cudaRegister.cu b/tests/hipify-clang/unit_tests/samples/cudaRegister.cu index 11192c452e..2cc754300a 100644 --- a/tests/hipify-clang/unit_tests/samples/cudaRegister.cu +++ b/tests/hipify-clang/unit_tests/samples/cudaRegister.cu @@ -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. diff --git a/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu b/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu index d35718d714..9e00a9b8e3 100644 --- a/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu +++ b/tests/hipify-clang/unit_tests/samples/dynamic_shared_memory.cu @@ -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 diff --git a/tests/hipify-clang/unit_tests/samples/intro.cu b/tests/hipify-clang/unit_tests/samples/intro.cu index da797eb2ec..5ae5479aa9 100644 --- a/tests/hipify-clang/unit_tests/samples/intro.cu +++ b/tests/hipify-clang/unit_tests/samples/intro.cu @@ -1,4 +1,4 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args #include #include diff --git a/tests/hipify-clang/unit_tests/samples/square.cu b/tests/hipify-clang/unit_tests/samples/square.cu index 076f9191f8..4609fd37e4 100644 --- a/tests/hipify-clang/unit_tests/samples/square.cu +++ b/tests/hipify-clang/unit_tests/samples/square.cu @@ -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. diff --git a/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu b/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu index ea6b50488f..f53932bf74 100644 --- a/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu +++ b/tests/hipify-clang/unit_tests/samples/static_shared_memory.cu @@ -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 diff --git a/tests/hipify-clang/unit_tests/samples/vec_add.cu b/tests/hipify-clang/unit_tests/samples/vec_add.cu index ec813e8bad..bc8219bf8c 100644 --- a/tests/hipify-clang/unit_tests/samples/vec_add.cu +++ b/tests/hipify-clang/unit_tests/samples/vec_add.cu @@ -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)