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); /** 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)); } 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)) -