diff --git a/CMakeLists.txt b/CMakeLists.txt
index 187b32045b..8b9a603ad8 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -221,7 +221,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}")
 
     # Add remaining flags
-    execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --cxxflags OUTPUT_VARIABLE HCC_CXX_FLAGS)
+    set(HCC_CXX_FLAGS  "-hc -fno-gpu-rdc --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 ")
     set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC ${HCC_CXX_FLAGS} -I${HSA_PATH}/include")
 
     # Set compiler and compiler flags
@@ -252,9 +252,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/env.cpp
         src/h2f.cpp)
 
-    execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
-    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS}")
-    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908")
     if(COMPILE_HIP_ATP_MARKER)
         set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger")
     endif()
@@ -270,7 +267,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
                 PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include)
         target_link_libraries(hiprtc PUBLIC stdc++fs)
     endif()
-
+    set_target_properties(hip_hcc PROPERTIES CXX_VISIBILITY_PRESET hidden)
+    set_target_properties(hip_hcc PROPERTIES VISIBILITY_INLINES_HIDDEN 1)
 
     if(HIP_PLATFORM STREQUAL "hcc")
         find_package(amd_comgr REQUIRED CONFIG
diff --git a/README.md b/README.md
index 200b2f7474..145f950fef 100644
--- a/README.md
+++ b/README.md
@@ -37,6 +37,15 @@ HIP releases are typically of two types. The tag naming convention is different
 - [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/OpenCL)
 - [hipify-clang](hipify-clang/README.md)
+- Supported CUDA APIs:
+  * [Runtime API](docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md)
+  * [Driver API](docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md)
+  * [cuComplex API](docs/markdown/cuComplex_API_supported_by_HIP.md)
+  * [cuBLAS](docs/markdown/CUBLAS_API_supported_by_HIP.md)
+  * [cuRAND](docs/markdown/CURAND_API_supported_by_HIP.md)
+  * [cuDNN](docs/markdown/CUDNN_API_supported_by_HIP.md)
+  * [cuFFT](docs/markdown/CUFFT_API_supported_by_HIP.md)
+  * [cuSPARSE](docs/markdown/CUSPARSE_API_supported_by_HIP.md)
 - [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
 - [Release Notes](RELEASE.md)
 
diff --git a/bin/hipcc b/bin/hipcc
index 1eda99b936..61f4f1a34e 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -118,10 +118,10 @@ if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
 if (defined $HIP_VDI_HOME) {
     if (!defined $HIP_CLANG_PATH and (-e "$HIP_VDI_HOME/bin/clang" or -e "$HIP_VDI_HOME/bin/clang.exe")) {
         $HIP_CLANG_PATH = "$HIP_VDI_HOME/bin";
-    }
-    # With HIP_VDI_HOME defined, assume the installation of clang components, including headers.
-    if (!defined $HIP_CLANG_INCLUDE_PATH) {
-        $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
+        # With HIP_VDI_HOME defined, assume the installation of clang components, including headers.
+        if (!defined $HIP_CLANG_INCLUDE_PATH) {
+            $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
+        }
     }
     if (!defined $DEVICE_LIB_PATH and -e "$HIP_VDI_HOME/lib/bitcode") {
         $DEVICE_LIB_PATH = "$HIP_VDI_HOME/lib/bitcode";
@@ -624,7 +624,7 @@ foreach $arg (@ARGV)
         #print "O: <$arg>\n";
     } else {
         # input files and libraries
-        if (($arg =~ /\.cpp$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) {
+        if (($arg =~ /\.cpp$/) or ($arg =~ /\.cxx$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) {
             $hasC = 1;
             $needCXXFLAGS = 1;
             if ($HIP_PLATFORM eq 'clang') {
diff --git a/bin/hipify-perl b/bin/hipify-perl
index e375d22928..625d2b6fd2 100755
--- a/bin/hipify-perl
+++ b/bin/hipify-perl
@@ -252,6 +252,8 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g;
         $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g;
         $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g;
+        $ft{'mem'} += s/\bcudaHostRegisterDefault\b/hipHostRegisterDefault/g;
+        $ft{'mem'} += s/\bcudaHostRegisterPortable\b/hipHostRegisterPortable/g;
         $ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g;
         $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g;
         $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g;
diff --git a/include/hip/hcc_detail/channel_descriptor.h b/include/hip/hcc_detail/channel_descriptor.h
index de290fafc1..38acff9951 100644
--- a/include/hip/hcc_detail/channel_descriptor.h
+++ b/include/hip/hcc_detail/channel_descriptor.h
@@ -23,11 +23,13 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
 
+#include <hip/hip_common.h>
 #include <hip/hcc_detail/driver_types.h>
 #include <hip/hcc_detail/hip_vector_types.h>
 
 #ifdef __cplusplus
 
+HIP_PUBLIC_API
 hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
 
 static inline hipChannelFormatDesc hipCreateChannelDescHalf() {
diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h
index 18f344d642..510d3d058e 100644
--- a/include/hip/hcc_detail/driver_types.h
+++ b/include/hip/hcc_detail/driver_types.h
@@ -319,4 +319,18 @@ static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) {
     return e;
 }
 
+typedef enum hipFunction_attribute {
+    HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK,
+    HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_NUM_REGS,
+    HIP_FUNC_ATTRIBUTE_PTX_VERSION,
+    HIP_FUNC_ATTRIBUTE_BINARY_VERSION,
+    HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA,
+    HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES,
+    HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT,
+    HIP_FUNC_ATTRIBUTE_MAX
+}hipFunction_attribute;
+
 #endif
diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp
index f502fddf42..883470b1fb 100644
--- a/include/hip/hcc_detail/functional_grid_launch.hpp
+++ b/include/hip/hcc_detail/functional_grid_launch.hpp
@@ -25,25 +25,16 @@ THE SOFTWARE.
 #include "concepts.hpp"
 #include "helpers.hpp"
 #include "program_state.hpp"
+#include "hip_runtime_api.h"
 
-#include "hc.hpp"
-#include "hip/hip_hcc.h"
-#include "hip_runtime.h"
-
-#include <algorithm>
-#include <cstddef>
 #include <cstdint>
 #include <cstring>
-#include <functional>
-#include <iostream>
-#include <mutex>
 #include <stdexcept>
-#include <string>
 #include <tuple>
 #include <type_traits>
-#include <unordered_map>
 #include <utility>
-#include <vector>
+
+#pragma GCC visibility push(hidden)
 
 namespace hip_impl {
 template <typename T, typename std::enable_if<std::is_integral<T>{}>::type* = nullptr>
@@ -114,7 +105,7 @@ inline hip_impl::kernarg make_kernarg(
 }
 
 
-hsa_agent_t target_agent(hipStream_t stream);
+HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream);
 
 inline
 __attribute__((visibility("hidden")))
@@ -194,3 +185,5 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc
     hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
                        std::move(args)...);
 }
+
+#pragma GCC visibility pop
diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index e8871912fa..cec9eb9c94 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -67,19 +67,18 @@ THE SOFTWARE.
 #define HIP_LAUNCH_PARAM_END ((void*)0x03)
 
 #ifdef __cplusplus
-  #include <algorithm>
-  #include <mutex>
-  #include <string>
-  #include <unordered_map>
-  #include <vector>
-
   #define __dparm(x) \
           = x
 #else
   #define __dparm(x)
 #endif
 
+#ifdef __GNUC__
+#pragma GCC visibility push (default)
+#endif
+
 #ifdef __cplusplus
+
 namespace hip_impl {
 hipError_t hip_init();
 }  // namespace hip_impl
@@ -2603,16 +2602,27 @@ hipError_t hipModuleUnload(hipModule_t module);
 hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname);
 
 /**
- * @bried Find out attributes for a given function.
+ * @brief Find out attributes for a given function.
  *
  * @param [out] attr
  * @param [in] func
  *
- * @returns hipSuccess, hipErrorInvalidDeviceFunction
+ * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction
  */
 
 hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
 
+/**
+ * @brief Find out a specific attribute for a given function.
+ *
+ * @param [out] value
+ * @param [in]  attrib
+ * @param [in]  hfunc
+ *
+ * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction
+ */
+hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc);
+
 #if !__HIP_VDI__
 #if defined(__cplusplus)
 } // extern "C"
@@ -3222,6 +3232,9 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
 
 #endif
 
+#ifdef __GNUC__
+#pragma GCC visibility pop
+#endif
 
 /**
  *-------------------------------------------------------------------------------------------------
diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp
index 4e0fbb76a4..fca88f8094 100644
--- a/include/hip/hcc_detail/program_state.hpp
+++ b/include/hip/hcc_detail/program_state.hpp
@@ -31,11 +31,19 @@ THE SOFTWARE.
 #include <cstdint>
 #include <cstdlib>
 
+#include <hip/hip_common.h>
+
 struct ihipModuleSymbol_t;
 using hipFunction_t = ihipModuleSymbol_t*;
 
 namespace hip_impl {
 
+// This section contains internal APIs that
+// needs to be exported
+#ifdef __GNUC__
+#pragma GCC visibility push (default)
+#endif
+
 struct kernarg_impl;
 class kernarg {
 public:
@@ -83,6 +91,10 @@ private:
     friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);
 };
 
+#ifdef __GNUC__
+#pragma GCC visibility pop
+#endif
+
 inline
 __attribute__((visibility("hidden")))
 program_state& get_program_state() {
diff --git a/include/hip/hip_common.h b/include/hip/hip_common.h
index 71285fc8dc..79c787be4e 100644
--- a/include/hip/hip_common.h
+++ b/include/hip/hip_common.h
@@ -45,6 +45,14 @@ THE SOFTWARE.
 #define __HIP_DEVICE_COMPILE__ 1
 #endif
 
+#ifdef __GNUC__
+#define HIP_PUBLIC_API              __attribute__ ((visibility ("default")))
+#define HIP_INTERNAL_EXPORTED_API   __attribute__ ((visibility ("default")))
+#else
+#define HIP_PUBLIC_API
+#define HIP_INTERNAL_EXPORTED_API 
+#endif
+
 #if __HIP_DEVICE_COMPILE__ == 0
 // 32-bit Atomics
 #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)
diff --git a/include/hip/hip_hcc.h b/include/hip/hip_hcc.h
index 3935f4903b..a82a8c622b 100644
--- a/include/hip/hip_hcc.h
+++ b/include/hip/hip_hcc.h
@@ -48,6 +48,7 @@ class accelerator_view;
  * @brief Return hc::accelerator associated with the specified deviceId
  * @return #hipSuccess, #hipErrorInvalidDevice
  */
+HIP_PUBLIC_API
 hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
 
 /**
@@ -56,6 +57,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
  * If stream is 0, the accelerator_view for the default stream is returned.
  * @return #hipSuccess
  */
+HIP_PUBLIC_API
 hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av);
 
 
@@ -89,6 +91,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a
 
  * HIP/ROCm actually updates the start event when the associated kernel completes.
  */
+HIP_PUBLIC_API
 hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                     uint32_t localWorkSizeX, uint32_t localWorkSizeY,
@@ -98,6 +101,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     hipEvent_t stopEvent = nullptr,
                                     uint32_t flags = 0);
 
+HIP_PUBLIC_API
 hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                     uint32_t localWorkSizeX, uint32_t localWorkSizeY,
diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index 1c8e971545..1f7183acf0 100644
--- a/include/hip/hip_runtime_api.h
+++ b/include/hip/hip_runtime_api.h
@@ -117,13 +117,11 @@ typedef struct hipDeviceProp_t {
     int integrated;            ///< APU vs dGPU
     int cooperativeLaunch;            ///< HIP device supports cooperative launch
     int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
-#if !__HIP_VDI__ // Temporarily disable the following five new fields for HIP/VDI runtime
     int maxTexture1D;          ///< Maximum number of elements in 1D images
     int maxTexture2D[2];       ///< Maximum dimensions (width, height) of 2D images, in image elements
     int maxTexture3D[3];       ///< Maximum dimensions (width, height, depth) of 3D images, in image elements
     unsigned int* hdpMemFlushCntl;      ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
     unsigned int* hdpRegFlushCntl;      ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
-#endif
 } hipDeviceProp_t;
 
 
diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h
index fb4987feee..a77d5560d8 100644
--- a/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/include/hip/nvcc_detail/hip_runtime_api.h
@@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_const_t;
 typedef cudaFuncAttributes hipFuncAttributes;
+typedef CUfunction_attribute hipFunction_attribute;
 #define hip_Memcpy2D CUDA_MEMCPY2D
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
@@ -197,6 +198,19 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t;
 #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
 #define hipSharedMemBankSizeEightByte cudaSharedMemBankSizeEightByte
 
+//Function Attributes
+#define HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK
+#define HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_NUM_REGS CU_FUNC_ATTRIBUTE_NUM_REGS
+#define HIP_FUNC_ATTRIBUTE_PTX_VERSION CU_FUNC_ATTRIBUTE_PTX_VERSION
+#define HIP_FUNC_ATTRIBUTE_BINARY_VERSION CU_FUNC_ATTRIBUTE_BINARY_VERSION
+#define HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA CU_FUNC_ATTRIBUTE_CACHE_MODE_CA
+#define HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES
+#define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT
+#define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX
+
 inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
     switch (cuError) {
         case cudaSuccess:
@@ -1197,6 +1211,10 @@ inline static hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const voi
     return hipCUDAErrorTohipError(cudaFuncGetAttributes(attr, func));
 }
 
+inline static hipError_t hipFuncGetAttribute (int* value, hipFunction_attribute attrib, hipFunction_t hfunc) {
+    return hipCUResultTohipError(cuFuncGetAttribute(value, attrib, hfunc));
+}
+
 inline static hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
                                             const char* name) {
     return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name));
diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp
index 7031b9f898..7aa7f7ba2d 100644
--- a/samples/0_Intro/module_api_global/runKernel.cpp
+++ b/samples/0_Intro/module_api_global/runKernel.cpp
@@ -124,6 +124,11 @@ int main() {
     {
         hipFunction_t Function;
         HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals"));
+        int val =-1;
+        HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function));
+        printf("Shared Size Bytes = %d\n",val);
+        HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function));
+        printf("Num Regs = %d\n",val);
         HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
 
         hipMemcpyDtoH(B, Bd, SIZE);
diff --git a/src/functional_grid_launch.inl b/src/functional_grid_launch.inl
index 3c611ed468..a99a15db28 100644
--- a/src/functional_grid_launch.inl
+++ b/src/functional_grid_launch.inl
@@ -41,7 +41,7 @@ using namespace std;
 
 namespace hip_impl
 {
-    hsa_agent_t target_agent(hipStream_t stream)
+    HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream)
     {
         if (stream) {
             return *static_cast<hsa_agent_t*>(
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index dcfdfb6137..da0b38ed4f 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -46,6 +46,7 @@ THE SOFTWARE.
 #include "hsa/hsa_ext_image.h"
 #include "hip/hip_runtime.h"
 #include "hip_hcc_internal.h"
+#include "hip/hip_hcc.h"
 #include "trace_helper.h"
 #include "env.h"
 
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 04cdedbb75..93751167ee 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -803,6 +803,52 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
     return ihipLogStatus(hipSuccess);
 }
 
+hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc)
+{
+    HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc);
+    using namespace hip_impl;
+    
+    hipError_t retVal = hipSuccess;
+    if (!value) return ihipLogStatus(hipErrorInvalidValue);
+    hipFuncAttributes attr{};
+    attr = make_function_attributes(tls, *hfunc);
+    switch(attrib) {
+        case HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES:
+            *value = (int) attr.sharedSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK:
+            *value = attr.maxThreadsPerBlock;
+            break;
+        case HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES:
+            *value = (int) attr.constSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES:
+            *value = (int) attr.localSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_NUM_REGS:
+            *value = attr.numRegs;
+            break;
+        case HIP_FUNC_ATTRIBUTE_PTX_VERSION:
+            *value = attr.ptxVersion;
+            break;
+        case HIP_FUNC_ATTRIBUTE_BINARY_VERSION:
+            *value = attr.binaryVersion;
+            break;
+        case HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA:
+            *value = attr.cacheModeCA;
+            break;
+        case HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES:
+            *value = attr.maxDynamicSharedSizeBytes;
+            break;
+        case HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT:
+            *value = attr.preferredShmemCarveout;
+            break;
+        default:
+             retVal = hipErrorInvalidValue;
+    }
+    return ihipLogStatus(retVal);
+}
+
 hipError_t ihipModuleLoadData(TlsData *tls, hipModule_t* module, const void* image) {
     using namespace hip_impl;
 
diff --git a/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp b/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
index a3283622ff..295d2f5c95 100644
--- a/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
+++ b/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
@@ -59,7 +59,9 @@ __host__ void double_precision_math_functions() {
     // erfcx(0.0);
     // erfinv(1.0);
     exp(0.0);
+	#ifdef __unix__
     exp10(0.0);
+	#endif
     exp2(0.0);
     expm1(0.0);
     fabs(1.0);
@@ -75,9 +77,15 @@ __host__ void double_precision_math_functions() {
     std::isfinite(0.0);
     std::isinf(0.0);
     std::isnan(0.0);
+    #ifdef __unix__
     j0(0.0);
     j1(0.0);
     jn(-1.0, 1.0);
+    #elif _WIN64
+    _j0(0.0);
+    _j1(0.0);
+    _jn(-1.0, 1.0);
+    #endif
     ldexp(0.0, 0);
     //    lgamma(1.0);
     llrint(0.0);
@@ -120,7 +128,9 @@ __host__ void double_precision_math_functions() {
     scalbn(0.0, 1);
     std::signbit(1.0);
     sin(0.0);
+    #ifdef _unix__
     sincos(0.0, &fX, &fY);
+    #endif
     // sincospi(0.0, &fX, &fY);
     sinh(0.0);
     // sinpi(0.0);
@@ -129,9 +139,15 @@ __host__ void double_precision_math_functions() {
     tanh(0.0);
     tgamma(2.0);
     trunc(0.0);
+    #ifdef __unix__
     y0(1.0);
     y1(1.0);
     yn(1, 1.0);
+    #elif _WIN64
+    _y0(1.0);
+    _y1(1.0);
+    _yn(1, 1.0);
+    #endif
 }
 
 static void compileOnHost() { double_precision_math_functions(); }
diff --git a/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp b/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
index 1dd5737f86..78c2099e8f 100644
--- a/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
+++ b/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
@@ -58,7 +58,9 @@ __host__ void single_precision_math_functions() {
     // erfcxf(0.0f);
     erff(0.0f);
     // erfinvf(1.0f);
+    #ifdef __unix__
     exp10f(0.0f);
+    #endif
     exp2f(0.0f);
     expf(0.0f);
     expm1f(0.0f);
@@ -78,9 +80,11 @@ __host__ void single_precision_math_functions() {
     std::isfinite(0.0f);
     std::isinf(0.0f);
     std::isnan(0.0f);
+    #ifdef __unix__
     j0f(0.0f);
     j1f(0.0f);
     jnf(-1.0f, 1.0f);
+    #endif
     ldexpf(0.0f, 0);
     lgammaf(1.0f);
     llrintf(0.0f);
@@ -121,7 +125,9 @@ __host__ void single_precision_math_functions() {
     scalblnf(0.0f, 1);
     scalbnf(0.0f, 1);
     std::signbit(1.0f);
+    #ifdef __unix__
     sincosf(0.0f, &fX, &fY);
+    #endif
     // sincospif(0.0f, &fX, &fY);
     sinf(0.0f);
     sinhf(0.0f);
@@ -131,9 +137,11 @@ __host__ void single_precision_math_functions() {
     tanhf(0.0f);
     tgammaf(2.0f);
     truncf(0.0f);
+    #ifdef __unix__
     y0f(1.0f);
     y1f(1.0f);
     ynf(1, 1.0f);
+    #endif
 }
 
 static void compileOnHost() { single_precision_math_functions(); }