Merge 'master' into 'amd-master'

Change-Id: I8491766916813e6f344bdc8aa65cda30bf9f3f63


[ROCm/hip commit: b3c5577266]

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -35,14 +35,20 @@ THE SOFTWARE.
 #define GENERIC_GRID_LAUNCH 1
 #endif
 
-#include <hsa/hsa.h>
+#ifndef __HIP_VDI__
+#define __HIP_VDI__ 0
+#endif
 
 #include <hip/hcc_detail/host_defines.h>
 #include <hip/hip_runtime_api.h>
 #include <hip/hcc_detail/driver_types.h>
 #include <hip/hcc_detail/hip_texture_types.h>
 #include <hip/hcc_detail/hip_surface_types.h>
+
+#if !__HIP_VDI__
+#include <hsa/hsa.h>
 #include <hip/hcc_detail/program_state.hpp>
+#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);
 /**

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

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'"
 

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))
-