diff --git a/CMakeLists.txt b/CMakeLists.txt
index 1ba58496f4..a1887d0cb5 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -63,6 +63,14 @@ if(HIP_PLATFORM STREQUAL "hcc")
             set(HCC_HOME $ENV{HCC_HOME} CACHE PATH "Path to which HCC has been installed")
         endif()
     endif()
+    # Determine LLVM_HOME
+    if(NOT DEFINED LLVM_HOME)
+        if(NOT DEFINED ENV{LLVM_HOME})
+            set(LLVM_HOME "/opt/rocm/llvm" CACHE PATH "Path to which LLVM has been installed")
+        else()
+            set(LLVM_HOME $ENV{LLVM_HOME} CACHE PATH "Path to which LLVM has been installed")
+        endif()
+    endif()
     if(DEFINED ENV{HIP_DEVELOPER})
         add_to_config(_buildInfo HCC_HOME)
     endif()
@@ -189,9 +197,11 @@ if(HIP_PLATFORM STREQUAL "hcc")
 
     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} -Wl,-Bsymbolic")
+    find_package(LLVM HINTS ${LLVM_HOME})
     set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803 --amdgpu-target=gfx900")
     add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
     target_link_libraries(hip_hcc PRIVATE hc_am)
+    target_link_libraries(hip_hcc PUBLIC LLVMAMDGPUUtils)
     add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
     target_link_libraries(hip_hcc_static PRIVATE hc_am)
     add_dependencies(hip_hcc_static hip_hcc)
diff --git a/bin/hipcc b/bin/hipcc
index 361517d551..4c71347cf7 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -103,6 +103,8 @@ if ($HIP_PLATFORM eq "hcc") {
 
     $HIPLDFLAGS = `${HCC_HOME}/bin/hcc-config --ldflags`;
 
+    $LLVM_HOME = $ENV{'LLVM_HOME'};
+
     #### GCC system includes workaround ####
     $HCC_WA_FLAGS = " ";
     if ($HCC_VERSION_MAJOR eq 1) {
@@ -127,7 +129,7 @@ if ($HIP_PLATFORM eq "hcc") {
     $HIPCXXFLAGS .= " -Wno-deprecated-register";
 
     $HIPLDFLAGS .= " -lsupc++";
-    $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";
+    $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt `${LLVM_HOME}/bin/llvm-config --ldflags --libs`";
 
     # Add trace marker library:
     # TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library.
@@ -402,9 +404,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
 
 if ($needHipHcc) {
     if ($linkType eq 0) {
-        substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ;
+        substr($HIPLDFLAGS,0,0) = "  $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ;
     } else {
-        substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a ";
+        substr($HIPLDFLAGS,0,0) = "  -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a ";
     }
 }
 
diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile
index f2c0ce555a..632a8d3e70 100644
--- a/samples/0_Intro/module_api/Makefile
+++ b/samples/0_Intro/module_api/Makefile
@@ -5,14 +5,16 @@ endif
 HIPCC=$(HIP_PATH)/bin/hipcc
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
 
-all: vcpy_kernel.code runKernel.hip.out
+all: vcpy_kernel.code runKernel.hip.out defaultDriver.hip.out
 
 runKernel.hip.out: runKernel.cpp
 	$(HIPCC) $(HIPCC_FLAGS) $< -o $@
 
+defaultDriver.hip.out: defaultDriver.cpp
+	$(HIPCC)  $(HIPCC_FLAGS) $< -o $@
+
 vcpy_kernel.code: vcpy_kernel.cpp
-	$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
+	$(HIPCC)  --genco  $(GENCO_FLAGS) $^ -o $@
 
 clean:
 	rm -f *.code *.out
-
diff --git a/samples/0_Intro/module_api/defaultDriver.cpp b/samples/0_Intro/module_api/defaultDriver.cpp
new file mode 100644
index 0000000000..a29271a0dc
--- /dev/null
+++ b/samples/0_Intro/module_api/defaultDriver.cpp
@@ -0,0 +1,89 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include<iostream>
+#include<fstream>
+#include<vector>
+
+#define LEN 64
+#define SIZE LEN<<2
+
+#define fileName "test.co"
+#define kernel_name "vadd"
+
+int main(){
+    float *A, *B, *C;
+    hipDeviceptr_t Ad, Bd, Cd;
+    A = new float[LEN];
+    B = new float[LEN];
+    C = new float[LEN];
+
+    for(uint32_t i=0;i<LEN;i++){
+        A[i] = i*1.0f;
+        B[i] = 1.0f;
+        C[i] = 0.0f;
+    }
+
+    hipInit(0);
+    hipDevice_t device;
+    hipCtx_t context;
+    hipDeviceGet(&device, 0);
+    hipCtxCreate(&context, 0, device);
+
+    hipMalloc((void**)&Ad, SIZE);
+    hipMalloc((void**)&Bd, SIZE);
+    hipMalloc((void**)&Cd, SIZE);
+
+    hipMemcpyHtoD(Ad, A, SIZE);
+    hipMemcpyHtoD(Bd, B, SIZE);
+    hipMemcpyHtoD(Cd, C, SIZE);
+
+    hipModule_t Module;
+    hipFunction_t Function;
+    hipModuleLoad(&Module, fileName);
+    hipModuleGetFunction(&Function, Module, kernel_name);
+
+    int n = LEN;
+    void * args[4] = {&Ad, &Bd, &Cd, &n};
+
+    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
+
+    hipMemcpyDtoH(C, Cd, SIZE);
+    int mismatchCount = 0;
+    for(uint32_t i=0;i<LEN;i++){
+        if (A[i] + B[i] != C[i]) {
+            mismatchCount++;
+            std::cout<<"error: mismatch " << A[i]<<" + "<<B[i]<<" != "<<C[i]<<std::endl;
+        }
+    }
+
+    if (mismatchCount == 0) {
+        std::cout << "PASSED!\n";
+    } else {
+        std::cout << "FAILED!\n";
+    };
+
+    hipCtxDestroy(context);
+    return 0;
+}
diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp
index b91507aaa4..201892a4a1 100644
--- a/samples/0_Intro/module_api/runKernel.cpp
+++ b/samples/0_Intro/module_api/runKernel.cpp
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -98,7 +98,7 @@ int main(){
       HIP_LAUNCH_PARAM_END
     };
 
-    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); 
+    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
 
     hipMemcpyDtoH(B, Bd, SIZE);
     int mismatchCount = 0;
diff --git a/samples/0_Intro/module_api/test.cl b/samples/0_Intro/module_api/test.cl
new file mode 100644
index 0000000000..81b20cab0e
--- /dev/null
+++ b/samples/0_Intro/module_api/test.cl
@@ -0,0 +1,12 @@
+__kernel void memset(char in, __global int* out) {
+int tx = get_global_id(0);
+out[tx] = in;
+}
+
+
+__kernel void vadd(__global float *Ad, __global float *Bd, __global float *Cd, int N){
+int tx = get_global_id(0);
+if(tx < N){
+Cd[tx] = Ad[tx] + Bd[tx];
+}
+}
diff --git a/samples/0_Intro/module_api/test.co b/samples/0_Intro/module_api/test.co
new file mode 100755
index 0000000000..a3e6b991e3
Binary files /dev/null and b/samples/0_Intro/module_api/test.co differ
diff --git a/src/AMDGPUPTNote.h b/src/AMDGPUPTNote.h
new file mode 100644
index 0000000000..8f8c855a52
--- /dev/null
+++ b/src/AMDGPUPTNote.h
@@ -0,0 +1,45 @@
+//===-- AMDGPUNoteType.h - AMDGPU ELF PT_NOTE section info-------*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+/// \file
+///
+/// Enums and constants for AMDGPU PT_NOTE sections.
+///
+//
+//===----------------------------------------------------------------------===//
+//
+#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUPTNOTE_H
+#define LLVM_LIB_TARGET_AMDGPU_AMDGPUPTNOTE_H
+
+namespace AMDGPU {
+
+namespace ElfNote {
+
+const char SectionName[] = ".note";
+
+const char NoteName[] = "AMD";
+
+// TODO: Move this enum to include/llvm/Support so it can be used in tools?
+enum NoteType{
+    NT_AMDGPU_HSA_CODE_OBJECT_VERSION = 1,
+    NT_AMDGPU_HSA_HSAIL = 2,
+    NT_AMDGPU_HSA_ISA = 3,
+    NT_AMDGPU_HSA_PRODUCER = 4,
+    NT_AMDGPU_HSA_PRODUCER_OPTIONS = 5,
+    NT_AMDGPU_HSA_EXTENSION = 6,
+    NT_AMDGPU_HSA_RUNTIME_METADATA_V_1 = 7, // deprecated since 12/14/16.
+    NT_AMDGPU_HSA_RUNTIME_METADATA_V_2 = 8,
+    NT_AMDGPU_HSA_RUNTIME_METADATA = NT_AMDGPU_HSA_RUNTIME_METADATA_V_2,
+    NT_AMDGPU_HSA_HLDEBUG_DEBUG = 101,
+    NT_AMDGPU_HSA_HLDEBUG_TARGET = 102
+};
+}
+}
+
+#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUNOTETYPE_H
diff --git a/src/AMDGPURuntimeMetadata.h b/src/AMDGPURuntimeMetadata.h
new file mode 100644
index 0000000000..ed147ff4c4
--- /dev/null
+++ b/src/AMDGPURuntimeMetadata.h
@@ -0,0 +1,290 @@
+//===-- AMDGPURuntimeMetadata.h - AMDGPU Runtime Metadata -------*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+/// \file
+///
+/// Enums and structure types used by runtime metadata.
+///
+/// Runtime requests certain information (metadata) about kernels to be able
+/// to execute the kernels and answer the queries about the kernels.
+/// The metadata is represented as a note element in the .note ELF section of a
+/// binary (code object). The desc field of the note element is a YAML string
+/// consisting of key-value pairs. Each key is a string. Each value can be
+/// an integer, a string, or an YAML sequence. There are 3 levels of YAML maps.
+/// At the beginning of the YAML string is the module level YAML map. A
+/// kernel-level YAML map is in the amd.Kernels sequence. A
+/// kernel-argument-level map is in the amd.Args sequence.
+///
+/// The format should be kept backward compatible. New enum values and bit
+/// fields should be appended at the end. It is suggested to bump up the
+/// revision number whenever the format changes and document the change
+/// in the revision in this header.
+///
+//
+//===----------------------------------------------------------------------===//
+//
+#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H
+#define LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H
+
+#include <cstdint>
+#include <vector>
+#include <string>
+
+namespace AMDGPU {
+namespace RuntimeMD {
+
+  // Version and revision of runtime metadata
+  const unsigned char MDVersion   = 2;
+  const unsigned char MDRevision  = 1;
+
+  // Name of keys for runtime metadata.
+  namespace KeyName {
+
+    // Runtime metadata version
+    const char MDVersion[] = "amd.MDVersion";
+
+    // Instruction set architecture information
+    const char IsaInfo[] = "amd.IsaInfo";
+    // Wavefront size
+    const char IsaInfoWavefrontSize[] = "amd.IsaInfoWavefrontSize";
+    // Local memory size in bytes
+    const char IsaInfoLocalMemorySize[] = "amd.IsaInfoLocalMemorySize";
+    // Number of execution units per compute unit
+    const char IsaInfoEUsPerCU[] = "amd.IsaInfoEUsPerCU";
+    // Maximum number of waves per execution unit
+    const char IsaInfoMaxWavesPerEU[] = "amd.IsaInfoMaxWavesPerEU";
+    // Maximum flat work group size
+    const char IsaInfoMaxFlatWorkGroupSize[] = "amd.IsaInfoMaxFlatWorkGroupSize";
+    // SGPR allocation granularity
+    const char IsaInfoSGPRAllocGranule[] = "amd.IsaInfoSGPRAllocGranule";
+    // Total number of SGPRs
+    const char IsaInfoTotalNumSGPRs[] = "amd.IsaInfoTotalNumSGPRs";
+    // Addressable number of SGPRs
+    const char IsaInfoAddressableNumSGPRs[] = "amd.IsaInfoAddressableNumSGPRs";
+    // VGPR allocation granularity
+    const char IsaInfoVGPRAllocGranule[] = "amd.IsaInfoVGPRAllocGranule";
+    // Total number of VGPRs
+    const char IsaInfoTotalNumVGPRs[] = "amd.IsaInfoTotalNumVGPRs";
+    // Addressable number of VGPRs
+    const char IsaInfoAddressableNumVGPRs[] = "amd.IsaInfoAddressableNumVGPRs";
+
+    // Language
+    const char Language[] = "amd.Language";
+    // Language version
+    const char LanguageVersion[] = "amd.LanguageVersion";
+
+    // Kernels
+    const char Kernels[] = "amd.Kernels";
+    // Kernel name
+    const char KernelName[] = "amd.KernelName";
+    // Kernel arguments
+    const char Args[] = "amd.Args";
+    // Kernel argument size in bytes
+    const char ArgSize[] = "amd.ArgSize";
+    // Kernel argument alignment
+    const char ArgAlign[] = "amd.ArgAlign";
+    // Kernel argument type name
+    const char ArgTypeName[] = "amd.ArgTypeName";
+    // Kernel argument name
+    const char ArgName[] = "amd.ArgName";
+    // Kernel argument kind
+    const char ArgKind[] = "amd.ArgKind";
+    // Kernel argument value type
+    const char ArgValueType[] = "amd.ArgValueType";
+    // Kernel argument address qualifier
+    const char ArgAddrQual[] = "amd.ArgAddrQual";
+    // Kernel argument access qualifier
+    const char ArgAccQual[] = "amd.ArgAccQual";
+    // Kernel argument is const qualified
+    const char ArgIsConst[] = "amd.ArgIsConst";
+    // Kernel argument is restrict qualified
+    const char ArgIsRestrict[] = "amd.ArgIsRestrict";
+    // Kernel argument is volatile qualified
+    const char ArgIsVolatile[] = "amd.ArgIsVolatile";
+    // Kernel argument is pipe qualified
+    const char ArgIsPipe[] = "amd.ArgIsPipe";
+    // Required work group size
+    const char ReqdWorkGroupSize[] = "amd.ReqdWorkGroupSize";
+    // Work group size hint
+    const char WorkGroupSizeHint[] = "amd.WorkGroupSizeHint";
+    // Vector type hint
+    const char VecTypeHint[] = "amd.VecTypeHint";
+    // Kernel index for device enqueue
+    const char KernelIndex[] = "amd.KernelIndex";
+    // No partial work groups
+    const char NoPartialWorkGroups[] = "amd.NoPartialWorkGroups";
+    // Prinf function call information
+    const char PrintfInfo[] = "amd.PrintfInfo";
+    // The actual kernel argument access qualifier
+    const char ArgActualAcc[] = "amd.ArgActualAcc";
+    // Alignment of pointee type
+    const char ArgPointeeAlign[] = "amd.ArgPointeeAlign";
+
+  } // end namespace KeyName
+
+  namespace KernelArg {
+
+    enum Kind : uint8_t {
+      ByValue                 = 0,
+      GlobalBuffer            = 1,
+      DynamicSharedPointer    = 2,
+      Sampler                 = 3,
+      Image                   = 4,
+      Pipe                    = 5,
+      Queue                   = 6,
+      HiddenGlobalOffsetX     = 7,
+      HiddenGlobalOffsetY     = 8,
+      HiddenGlobalOffsetZ     = 9,
+      HiddenNone              = 10,
+      HiddenPrintfBuffer      = 11,
+      HiddenDefaultQueue      = 12,
+      HiddenCompletionAction  = 13,
+    };
+
+    enum ValueType : uint16_t {
+      Struct  = 0,
+      I8      = 1,
+      U8      = 2,
+      I16     = 3,
+      U16     = 4,
+      F16     = 5,
+      I32     = 6,
+      U32     = 7,
+      F32     = 8,
+      I64     = 9,
+      U64     = 10,
+      F64     = 11,
+    };
+
+    // Avoid using 'None' since it conflicts with a macro in X11 header file.
+    enum AccessQualifer : uint8_t {
+      AccNone    = 0,
+      ReadOnly   = 1,
+      WriteOnly  = 2,
+      ReadWrite  = 3,
+    };
+
+    enum AddressSpaceQualifer : uint8_t {
+      Private    = 0,
+      Global     = 1,
+      Constant   = 2,
+      Local      = 3,
+      Generic    = 4,
+      Region     = 5,
+    };
+
+  } // end namespace KernelArg
+
+  // Invalid values are used to indicate an optional key should not be emitted.
+  const uint8_t INVALID_ADDR_QUAL     = 0xff;
+  const uint8_t INVALID_ACC_QUAL      = 0xff;
+  const uint32_t INVALID_KERNEL_INDEX = ~0U;
+
+  namespace KernelArg {
+
+    // In-memory representation of kernel argument information.
+    struct Metadata {
+      uint32_t Size = 0;
+      uint32_t Align = 0;
+      uint32_t PointeeAlign = 0;
+      uint8_t Kind = 0;
+      uint16_t ValueType = 0;
+      std::string TypeName;
+      std::string Name;
+      uint8_t AddrQual = INVALID_ADDR_QUAL;
+      uint8_t AccQual = INVALID_ACC_QUAL;
+      uint8_t IsVolatile = 0;
+      uint8_t IsConst = 0;
+      uint8_t IsRestrict = 0;
+      uint8_t IsPipe = 0;
+
+      Metadata() = default;
+    };
+
+  } // end namespace KernelArg
+
+  namespace Kernel {
+
+    // In-memory representation of kernel information.
+    struct Metadata {
+      std::string Name;
+      std::string Language;
+      std::vector<uint8_t> LanguageVersion;
+      std::vector<uint32_t> ReqdWorkGroupSize;
+      std::vector<uint32_t> WorkGroupSizeHint;
+      std::string VecTypeHint;
+      uint32_t KernelIndex = INVALID_KERNEL_INDEX;
+      uint8_t NoPartialWorkGroups = 0;
+      std::vector<KernelArg::Metadata> Args;
+
+      Metadata() = default;
+    };
+
+  } // end namespace Kernel
+
+  namespace IsaInfo {
+
+    /// \brief In-memory representation of instruction set architecture
+    /// information.
+    struct Metadata {
+      /// \brief Wavefront size.
+      unsigned WavefrontSize = 0;
+      /// \brief Local memory size in bytes.
+      unsigned LocalMemorySize = 0;
+      /// \brief Number of execution units per compute unit.
+      unsigned EUsPerCU = 0;
+      /// \brief Maximum number of waves per execution unit.
+      unsigned MaxWavesPerEU = 0;
+      /// \brief Maximum flat work group size.
+      unsigned MaxFlatWorkGroupSize = 0;
+      /// \brief SGPR allocation granularity.
+      unsigned SGPRAllocGranule = 0;
+      /// \brief Total number of SGPRs.
+      unsigned TotalNumSGPRs = 0;
+      /// \brief Addressable number of SGPRs.
+      unsigned AddressableNumSGPRs = 0;
+      /// \brief VGPR allocation granularity.
+      unsigned VGPRAllocGranule = 0;
+      /// \brief Total number of VGPRs.
+      unsigned TotalNumVGPRs = 0;
+      /// \brief Addressable number of VGPRs.
+      unsigned AddressableNumVGPRs = 0;
+
+      Metadata() = default;
+    };
+
+  } // end namespace IsaInfo
+
+  namespace Program {
+
+    // In-memory representation of program information.
+    struct Metadata {
+      std::vector<uint8_t> MDVersionSeq;
+      IsaInfo::Metadata IsaInfo;
+      std::vector<std::string> PrintfInfo;
+      std::vector<Kernel::Metadata> Kernels;
+
+      explicit Metadata() = default;
+
+      // Construct from an YAML string.
+      explicit Metadata(const std::string &YAML);
+
+      // Convert to YAML string.
+      std::string toYAML();
+
+      // Convert from YAML string.
+      static Metadata fromYAML(const std::string &S);
+    };
+
+  } //end namespace Program
+
+} // end namespace RuntimeMD
+} // end namespace AMDGPU
+
+#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 1f20a47c13..dc0a681c6d 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -23,7 +23,12 @@ THE SOFTWARE.
 #include <fstream>
 #include <stdio.h>
 #include <stdlib.h>
+#include <libelf.h>
 #include <elf.h>
+#include <gelf.h>
+#include <map>
+#include "AMDGPUPTNote.h"
+#include "AMDGPURuntimeMetadata.h"
 
 #include "hsa/hsa.h"
 #include "hsa/hsa_ext_amd.h"
@@ -35,6 +40,30 @@ THE SOFTWARE.
 
 //TODO Use Pool APIs from HCC to get memory regions.
 
+#include <cassert>
+inline uint64_t alignTo(uint64_t Value, uint64_t Align, uint64_t Skew = 0) {
+  assert(Align != 0u && "Align can't be 0.");
+  Skew %= Align;
+  return (Value + Align - 1 - Skew) / Align * Align + Skew;
+}
+
+struct ihipKernArgInfo{
+  std::vector<uint32_t> Size;
+  std::vector<uint32_t> Align;
+  std::vector<std::string> ArgType;
+  std::vector<std::string> ArgName;
+  uint32_t totalSize;
+};
+
+std::map<std::string,struct ihipKernArgInfo> kernelArguments;
+
+struct MyElfNote {
+  uint32_t n_namesz = 0;
+  uint32_t n_descsz = 0;
+  uint32_t n_type = 0;
+
+  MyElfNote() = default;
+};
 
 struct ihipModuleSymbol_t{
     uint64_t    _object;             // The kernel object.
@@ -172,8 +201,56 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
             (*module)->ptr = p;
             (*module)->size = size;
             in.seekg(0, std::ios::beg);
-            std::copy(std::istreambuf_iterator<char>(in),
-                      std::istreambuf_iterator<char>(), ptr);
+            std::copy(std::istreambuf_iterator<char>(in), std::istreambuf_iterator<char>(), ptr);
+
+            Elf *e = elf_memory((char*)p, size);
+            if(elf_kind(e) != ELF_K_ELF){
+              return ihipLogStatus(hipErrorInvalidValue);
+            }
+            size_t numpHdrs;
+            if(elf_getphdrnum(e, &numpHdrs) != 0){
+              return ihipLogStatus(hipErrorInvalidValue);
+            }
+            for(size_t i=0;i<numpHdrs;i++){
+              GElf_Phdr pHdr;
+              if(gelf_getphdr(e, i, &pHdr) != &pHdr){
+                continue;
+              }
+              if(pHdr.p_type == PT_NOTE && pHdr.p_align >= sizeof(int)){
+                char* ptr = (char*) p + pHdr.p_offset;
+                char* segmentEnd = ptr + pHdr.p_filesz;
+                while(ptr < segmentEnd){
+                  MyElfNote *note = (MyElfNote*) ptr;
+                  char *name = (char*) &note[1];
+                  char *desc = name + alignTo(note->n_namesz, sizeof(int));
+                  if (note->n_type == 8) {
+                    std::string metadatastr((const char*)desc, (size_t)note->n_descsz);
+                    AMDGPU::RuntimeMD::Program::Metadata meta;
+                    meta = meta.fromYAML(metadatastr);
+                    for(int i=0;i<meta.Kernels.size();i++){
+                      struct ihipKernArgInfo pl;
+                      size_t totalSizeOfArgs = 0;
+                      for(int j=0;j<meta.Kernels[i].Args.size();j++){
+                        if(meta.Kernels[i].Args[j].Kind < 7) {
+                          pl.Size.push_back(meta.Kernels[i].Args[j].Size);
+                          pl.Align.push_back(meta.Kernels[i].Args[j].Align);
+                          pl.ArgType.push_back(meta.Kernels[i].Args[j].TypeName);
+                          pl.ArgName.push_back(meta.Kernels[i].Args[j].Name);
+                          totalSizeOfArgs += meta.Kernels[i].Args[j].Align;
+                        }
+                      }
+                      pl.totalSize = totalSizeOfArgs;
+                      kernelArguments[meta.Kernels[i].Name] = pl;
+                    }
+                    break;
+                  }
+
+                  ptr += sizeof(*note)
+                      + alignTo(note->n_namesz, sizeof(int))
+                      + alignTo(note->n_descsz, sizeof(int));
+                }
+              }
+            }
 
             status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object);
 
@@ -313,7 +390,19 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         void *config[5] = {0};
         size_t kernArgSize;
 
-        if(extra != NULL){
+        if(kernelParams != NULL){
+          std::string name = f->_name;
+          struct ihipKernArgInfo pl = kernelArguments[name];
+          char* argBuf = (char*)malloc(pl.totalSize);
+          memset(argBuf, 0, pl.totalSize);
+          int index = 0;
+          for(int i=0;i<pl.Size.size();i++){
+            memcpy(argBuf + index, kernelParams[i], pl.Size[i]);
+            index += pl.Align[i];
+          }
+          config[1] = (void*)argBuf;
+          kernArgSize = pl.totalSize;
+        } else if(extra != NULL){
             memcpy(config, extra, sizeof(size_t)*5);
             if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){
                 kernArgSize = *(size_t*)(config[3]);
@@ -364,7 +453,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 
         lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, nullptr/*completion_future*/);
 
-
+        if(kernelParams != NULL){
+          free(config[1]);
+        }
         ihipPostLaunchKernel(f->_name.c_str(), hStream, lp);
     }