From 1397bf79d85b141df0d2a966c07542ffc799a731 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 6 Mar 2019 19:26:05 +0300 Subject: [PATCH 1/3] [HIPIFY][tests] Update lit testing infrastructure + Set -D__LP64__ in case of 64-bit hipify-clang binary [partial workaround for clang's bug https://bugs.llvm.org/show_bug.cgi?id=38811] C:/GIT/LLVM/trunk/llvm-64-release-vs2017/dist/lib/clang/9.0.0\include\__clang_cuda_device_functions.h(1609,45): error GEF7559A7: no matching function for call to 'roundf' __DEVICE__ long lroundf(float __a) { return roundf(__a); } #if defined(__LP64__) __DEVICE__ long lround(double __a) { return llround(__a); } __DEVICE__ long lroundf(float __a) { return llroundf(__a); } // ok: llroundf should be used when 64-bit #else __DEVICE__ long lround(double __a) { return round(__a); } __DEVICE__ long lroundf(float __a) { return roundf(__a); } // error #endif + Print more system info while testing in the following form: ======================================== CUDA 9.0 - will be used for testing LLVM 9.0.0svn - will be used for testing AMD64 - Platform architecture Windows 10 - Platform OS 64 - hipify-clang binary bitness 32 - python 3.7.2 binary bitness ======================================== [ROCm/hip commit: e899ee0e060b113535d91e4a1d478300483a0afd] --- projects/hip/tests/hipify-clang/lit.cfg | 12 +++++++++++- projects/hip/tests/hipify-clang/lit.site.cfg.in | 3 ++- 2 files changed, 13 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 00a3eeb873..594d5fcf69 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/tests/hipify-clang/lit.cfg @@ -3,6 +3,7 @@ import os import platform import re import subprocess +import struct import lit.formats import lit.util @@ -11,7 +12,14 @@ import lit.util site_cfg = lit_config.params.get('site_config', None) lit_config.load_config(config, site_cfg) -print("CUDA " + config.cuda_version + " will be used for testing.") +print(str("========================================")) +print("CUDA " + config.cuda_version + " - will be used for testing") +print("LLVM " + config.llvm_version + " - will be used for testing") +print(platform.machine() + " - Platform architecture") +print(platform.system() + " " + platform.release() + " - Platform OS") +print(str(config.pointer_size * 8) + " - hipify-clang binary bitness") +print(str(struct.calcsize("P") * 8) + " - python " + str(platform.python_version()) + " binary bitness") +print(str("========================================")) config.excludes = ['cmdparser.hpp'] config.excludes.append('spatial_batch_norm_op.h') @@ -77,6 +85,8 @@ else: run_test_ext = ".sh" clang_arguments += " -isystem'%s'/samples/common/inc" clang_arguments += " -I'%s'/include" +if config.pointer_size == 8: + clang_arguments += " -D__LP64__" hipify_arguments = "--cuda-path='%s'" diff --git a/projects/hip/tests/hipify-clang/lit.site.cfg.in b/projects/hip/tests/hipify-clang/lit.site.cfg.in index f93c206be4..71cc3e08d3 100644 --- a/projects/hip/tests/hipify-clang/lit.site.cfg.in +++ b/projects/hip/tests/hipify-clang/lit.site.cfg.in @@ -1,6 +1,8 @@ import sys import os +config.pointer_size = @CMAKE_SIZEOF_VOID_P@ +config.llvm_version = "@LLVM_PACKAGE_VERSION@" config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@" config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" @@ -32,4 +34,3 @@ except KeyError: e = sys.exc_info()[1] key, = e.args lit_config.fatal("unable to find %r parameter, use '--param=%s=VALUE'" % (key,key)) - From 17e2bfe3f88fd06dd83946cdb2b6015f5489c9c8 Mon Sep 17 00:00:00 2001 From: Steven Noonan Date: Wed, 6 Mar 2019 11:21:10 -0800 Subject: [PATCH 2/3] nvcc_detail/hip_runtime_api.h: add missing hipDeviceSetCacheConfig API Signed-off-by: Steven Noonan [ROCm/hip commit: ee750d5ea4f7cd375302e3b086ccefd10dd64c0a] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index add4c3f238..4d84cfbc51 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -79,6 +79,12 @@ typedef enum hipChannelFormatKind { #define hipBoundaryModeTrap cudaBoundaryModeTrap #define hipBoundaryModeClamp cudaBoundaryModeClamp +// hipFuncCache +#define hipFuncCachePreferNone cudaFuncCachePreferNone +#define hipFuncCachePreferShared cudaFuncCachePreferShared +#define hipFuncCachePreferL1 cudaFuncCachePreferL1 +#define hipFuncCachePreferEqual cudaFuncCachePreferEqual + // hipResourceType #define hipResourceType cudaResourceType #define hipResourceTypeArray cudaResourceTypeArray @@ -614,6 +620,10 @@ inline static hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* pCacheConfig) { return hipCUDAErrorTohipError(cudaDeviceGetCacheConfig(pCacheConfig)); } +inline static hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig) { + return hipCUDAErrorTohipError(cudaDeviceSetCacheConfig(cacheConfig)); +} + inline static const char* hipGetErrorString(hipError_t error) { return cudaGetErrorString(hipErrorToCudaError(error)); } From a431326408f5ea714db073cec9f782d31823e240 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Thu, 7 Mar 2019 14:45:45 -0500 Subject: [PATCH 3/3] Fix HIP/VDI build failure due to dlopen change [ROCm/hip commit: 6e9e90addd618b04d67854a51fba5ebf77c8f723] --- .../hip/include/hip/hcc_detail/hip_runtime_api.h | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index cb4e073e4e..7b7ffb0638 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -35,14 +35,20 @@ THE SOFTWARE. #define GENERIC_GRID_LAUNCH 1 #endif -#include +#ifndef __HIP_VDI__ +#define __HIP_VDI__ 0 +#endif #include #include #include #include #include + +#if !__HIP_VDI__ +#include #include +#endif #if defined(_MSC_VER) #define DEPRECATED(msg) __declspec(deprecated(msg)) @@ -1371,6 +1377,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); +#if !__HIP_VDI__ __attribute__((visibility("hidden"))) hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*); @@ -1551,6 +1558,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, (const char*)symbolName); } +#endif // __HIP_VDI__ /** * @brief Copy data from src to dst asynchronously. * @@ -2502,6 +2510,8 @@ struct Agent_global { hipDeviceptr_t address; uint32_t byte_cnt; }; + +#if !__HIP_VDI__ #if defined(__cplusplus) } // extern "C" #endif @@ -2593,7 +2603,6 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, #if defined(__cplusplus) extern "C" { #endif - /** * @brief returns device memory pointer and size of the kernel present in the module with symbol @p * name @@ -2619,6 +2628,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, return r; } +#endif // __HIP_VDI__ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name); /**