Merge branch 'amd-develop' into amd-master

Change-Id: Ie5aa8e70607936f63cf4c763298140e83a375a68
2016-11-28 09:18:11 +05:30
parent 4b3d0f0e7e a504df955e
commit 8a464ac850
57 changed files with 3836 additions and 407 deletions
@@ -173,32 +173,38 @@ if(HIP_PLATFORM STREQUAL "hcc")
    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
    set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}")

-    set(SOURCE_FILES src/device_util.cpp
+    set(SOURCE_FILES_SHARED
        src/hip_hcc.cpp
        src/hip_context.cpp
        src/hip_device.cpp
        src/hip_error.cpp
        src/hip_event.cpp
-        src/hip_ldg.cpp
        src/hip_memory.cpp
        src/hip_peer.cpp
        src/hip_stream.cpp
-        src/hip_fp16.cpp
        src/hip_module.cpp)

+    set(SOURCE_FILES_STATIC
+        src/device_util.cpp
+        src/hip_ldg.cpp
+        src/hip_fp16.cpp)
+
    if(${HIP_LIB_TYPE} EQUAL 0)
-        add_library(hip_hcc OBJECT ${SOURCE_FILES})
+        add_library(hip_hcc OBJECT ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
    elseif(${HIP_LIB_TYPE} EQUAL 1)
-        add_library(hip_hcc STATIC ${SOURCE_FILES})
+        add_library(hip_hcc STATIC ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
    else()
-        add_library(hip_hcc SHARED ${SOURCE_FILES})
+        set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
+        add_library(hip_hcc SHARED ${SOURCE_FILES_SHARED})
+        add_library(hip_device STATIC ${SOURCE_FILES_STATIC})
+        add_dependencies(hip_device hip_hcc)
    endif()

    # Generate hcc_version.txt
    add_custom_target(query_hcc_version COMMAND ${HCC_HOME}/bin/hcc --version > ${PROJECT_BINARY_DIR}/hcc_version.tmp)
    add_custom_target(check_hcc_version COMMAND ${CMAKE_COMMAND} -E copy_if_different ${PROJECT_BINARY_DIR}/hcc_version.tmp ${PROJECT_BINARY_DIR}/hcc_version.txt DEPENDS query_hcc_version)
    set_source_files_properties(${PROJECT_BINARY_DIR}/hcc_version.txt PROPERTIES GENERATED TRUE)
-    set_source_files_properties(${SOURCE_FILES} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
+    set_source_files_properties(${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
    add_dependencies(hip_hcc check_hcc_version update_build_and_version_info)

    # Generate .hipInfo
@@ -219,8 +225,10 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA
 if(HIP_PLATFORM STREQUAL "hcc")
    if(${HIP_LIB_TYPE} EQUAL 0)
        install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
-    else()
+    elseif(${HIP_LIB_TYPE} EQUAL 1)
        install(TARGETS hip_hcc DESTINATION lib)
+    else()
+        install(TARGETS hip_hcc hip_device DESTINATION lib)
    endif()
    install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_ir.ll DESTINATION lib)

@@ -274,6 +274,11 @@ foreach $arg (@ARGV)
        $buildDeps = 1;
    }

+    if($arg eq '-use_fast_math') {
+        print "In fast Math";
+        $HIPCXXFLAGS .= " -DHIP_FAST_MATH ";
+    }
+
    if ($arg =~ m/^-/) {
        # options start with -

@@ -328,7 +333,7 @@ if ($needHipHcc) {
    } elsif ($HIP_LIB_TYPE eq 1) {
        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ;
    } else {
-        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc ";
+        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc -lhip_device ";
    }
 }

@@ -20,19 +20,19 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 ##
-#usage hipify [OPTIONS] INPUT_FILE 
+#usage hipify [OPTIONS] INPUT_FILE
 use Getopt::Long;

 my $warn_whitelist ="";
 GetOptions(
      "print-stats" => \$print_stats                   # print the command-line, like a header.
    , "count-conversions" => \$count_conversions       # count conversions.
-    , "quiet-warnings" => \$quiet_warnings           # don't print warnings on unknown CUDA functions. 
-    , "warn-whitelist=s"=> \$warn_whitelist 
+    , "quiet-warnings" => \$quiet_warnings           # don't print warnings on unknown CUDA functions.
+    , "warn-whitelist=s"=> \$warn_whitelist
    , "no-translate-builtins" => \$no_translate_builtins      # don't translate math functions.
    , "no-translate-textures" => \$no_translate_textures
    , "no-output" => \$no_output                      # don't write any translated output to stdout.
-    , "inplace" => \$inplace                          # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file. 
+    , "inplace" => \$inplace                          # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file.
 														# If .prehip file exists, use that as input to hip.
    , "n" => \$n                                      # combination of print_stats + no-output.
 );
@@ -89,7 +89,7 @@ push (@warn_whitelist, split(',',$warn_whitelist));

 #---
 #Compute total of all individual counts:
-sub totalStats 
+sub totalStats
 {
    my %count = %{ shift() };

@@ -113,11 +113,11 @@ sub printStats
    my $total = totalStats(\%counts);

    printf STDERR "%s %d CUDA->HIP refs( ", $label, $total;
-    
+
    foreach $stat (@statNames) {
        printf STDERR "%s:%d ", $stat, $counts{$stat};
-    } 
-    
+    }
+
    printf STDERR ") warn:%d LOC:%d", $warnings, $loc;
 }

@@ -196,11 +196,11 @@ while (@ARGV) {

    # Note : \b is used in perl to indicate the start of a word - typically that is what we want in this case:
    #
-    
+
    # count of transforms in this file, init to 0 here:
    my %ft;
    clearStats(\%ft, \@statNames);
-    my $countIncludes = 0;  
+    my $countIncludes = 0;
    my $countKeywords = 0; # keywords like __global__, __shared__ - not converted by hipify but counted here.
    my $warnings = 0;
    my $warningsCublas = 0;
@@ -218,11 +218,11 @@ while (@ARGV) {
        # __CUDACC__ is set by NVCC to indicate it is treating the input file as CUDA code (as opposed to host)
        # Typically we want any code treated as CUDA code to be treated as accelerator code by Kalmar too
        # __HIPCC__ will set KALMARCC
-        $ft{'def'} += s/\b__CUDACC__\b/__HIPCC__/g;  
+        $ft{'def'} += s/\b__CUDACC__\b/__HIPCC__/g;

        # __CUDA_ARCH is often used to detect when a function or kernel is being compiled for the device.
        # Don't automaticall convert this - likely these will need special attention with HIP_ARCH_HAS_* macros
-        #$ft{'def'} += s/\b__CUDA_ARCH__\b/__HIP_ARCH__/g;  
+        #$ft{'def'} += s/\b__CUDA_ARCH__\b/__HIP_ARCH__/g;



@@ -230,6 +230,7 @@ while (@ARGV) {
        #Includes:
        $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime\.h[>"]/$1<hip\/hip_runtime.h>/;
        $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime_api\.h[>"]/$1<hip\/hip_runtime_api.h>/;
+        $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_fp16\.h[>"]/$1<hip\/hip_fp16.h>/;


        #--------
@@ -249,6 +250,7 @@ while (@ARGV) {
        $ft{'err'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g;
        $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g;
        $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g;
+        $ft{'err'} += s/\bcudaErrorPeerAccessAlreadyEnabled\b/hipErrorPeerAccessAlreadyEnabled/g;

 		# error APIs:
        $ft{'err'} += s/\bcudaGetLastError\b/hipGetLastError/g;
@@ -294,6 +296,10 @@ 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/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g;
+        $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g;
+        $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g;
+        $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;

        $ft{'mem'} += s/\bcudaMallocArray\b/hipMallocArray/g;
        $ft{'mem'} += s/\bcudaMallocPitch\b/hipMallocPitch/g;
@@ -308,7 +314,7 @@ while (@ARGV) {
        $ft{'coord_func'} += s/\bblockIdx\.x\b/hipBlockIdx_x/g;
        $ft{'coord_func'} += s/\bblockIdx\.y\b/hipBlockIdx_y/g;
        $ft{'coord_func'} += s/\bblockIdx\.z\b/hipBlockIdx_z/g;
-        
+
        $ft{'coord_func'} += s/\bblockDim\.x\b/hipBlockDim_x/g;
        $ft{'coord_func'} += s/\bblockDim\.y\b/hipBlockDim_y/g;
        $ft{'coord_func'} += s/\bblockDim\.z\b/hipBlockDim_z/g;
@@ -322,7 +328,7 @@ while (@ARGV) {


        #--------
-        # Events 
+        # Events
        $ft{'event'} += s/\bcudaEvent_t\b/hipEvent_t/g;
        $ft{'event'} += s/\bcudaEventCreate\b/hipEventCreate/g;
        $ft{'event'} += s/\bcudaEventCreateWithFlags\b/hipEventCreateWithFlags/g;
@@ -330,6 +336,7 @@ while (@ARGV) {
        $ft{'event'} += s/\bcudaEventRecord\b/hipEventRecord/g;
        $ft{'event'} += s/\bcudaEventElapsedTime\b/hipEventElapsedTime/g;
        $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g;
+        $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g;

        #--------
        # Streams
@@ -341,10 +348,10 @@ while (@ARGV) {
        $ft{'stream'} += s/\bcudaStreamSynchronize\b/hipStreamSynchronize/g;
        $ft{'stream'} += s/\bcudaStreamDefault\b/hipStreamDefault/g;
        $ft{'stream'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g;
-        
+

        #--------
-        # Other synchronization 
+        # Other synchronization
        $ft{'dev'} += s/\bcudaDeviceSynchronize\b/hipDeviceSynchronize/g;
        $ft{'dev'} += s/\bcudaThreadSynchronize\b/hipDeviceSynchronize/g;  # translate deprecated cudaThreadSynchronize
        $ft{'dev'} += s/\bcudaDeviceReset\b/hipDeviceReset/g;
@@ -353,9 +360,10 @@ while (@ARGV) {
        $ft{'dev'} += s/\bcudaGetDevice\b/hipGetDevice/g;

        #--------
-        # Device 
+        # Device
        $ft{'dev'} += s/\bcudaDeviceProp\b/hipDeviceProp_t/g;
        $ft{'dev'} += s/\bcudaGetDeviceProperties\b/hipGetDeviceProperties/g;
+        $ft{'dev'} += s/\bcudaDeviceGetPCIBusId\b/hipDeviceGetPCIBusId/g;

        # Attribute
        $ft{'err'} += s/\bcudaDevAttrMaxThreadsPerBlock\b/hipDeviceAttributeMaxThreadsPerBlock/g;
@@ -408,6 +416,11 @@ while (@ARGV) {
        $ft{'dev'} += s/\bcudaDeviceEnablePeerAccess\b/hipDeviceEnablePeerAccess/g;
        $ft{'mem'} += s/\bcudaMemcpyPeerAsync\b/hipMemcpyPeerAsync/g;
        $ft{'mem'} += s/\bcudaMemcpyPeer\b/hipMemcpyPeer/g;
+        $ft{'mem'} += s/\bcudaIpcOpenMemHandle\b/hipIpcOpenMemHandle/g;
+        $ft{'mem'} += s/\bcudaIpcCloseMemHandle\b/hipIpcCloseMemHandle/g;
+        $ft{'mem'} += s/\bcudaIpcGetMemHandle\b/hipIpcGetMemHandle/g;
+        $ft{'mem'} += s/\bcudaIpcMemHandle_t\b/hipIpcMemHandle_t/g;
+        $ft{'mem'} += s/\bcudaIpcMemLazyEnablePeerAccess\b/hipIpcMemLazyEnablePeerAccess/g;


        # Shared mem:
@@ -441,7 +454,7 @@ while (@ARGV) {
            no warnings qw/uninitialized/;

            my $k = 0;
- 
+
            # Match extern __shared__ type foo[]; syntax
            # Replace as HIP_DYNAMIC_SHARED() macro
            $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;
@@ -464,7 +477,7 @@ while (@ARGV) {
        #--------
        # CUDA Launch Syntax
        # Note these only work if launch is on a single line.
-        
+
        # Handle the <<numBlocks, blockDim>>> syntax:
        {
            # match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
@@ -524,7 +537,7 @@ while (@ARGV) {
        unless ($quiet_warnings) {
            #print STDERR "Check WARNINGs\n";
            # copy into array of lines, process line-by-line to show warnings:
-            if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/) or (/(\bcublas[A-Z]\w+)/) or (/(\bcurand[A-Z]\w+)/) ) { 
+            if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/) or (/(\bcublas[A-Z]\w+)/) or (/(\bcurand[A-Z]\w+)/) ) {
                my @lines = split /\n/, $_;
                my $tmp = $_;  # copies the whole file, could be a little smarter here...
                my $line_num = 0;
@@ -541,7 +554,7 @@ while (@ARGV) {
 					my $tag ;
                    if ((/(\bcuda[A-Z]\w+)/) or (/<<<.*>>>/)) {
                        # flag any remaining code that look like cuda API calls, may want to add these to hipify
-                        $tag = (defined $1) ? $1 : "Launch"; 
+                        $tag = (defined $1) ? $1 : "Launch";
 					} elsif (/(\bcublas[A-Z]\w+)/) {
                        $warningsCublas++;
 						$tag = $1;
@@ -552,7 +565,7 @@ while (@ARGV) {

 					if (defined $tag) {
                        $warnings++;
-                        $warningTags{$tag}++; 
+                        $warningTags{$tag}++;
                        print STDERR "  warning: $fileName:#$line_num : $_";
                        print STDERR "\n";
                    }
@@ -599,7 +612,7 @@ while (@ARGV) {
            print $OUTFILE  "$_";
        }

-        $lineCount = $_ =~ tr/\n//; 
+        $lineCount = $_ =~ tr/\n//;

    }

@@ -653,7 +666,7 @@ if ($count_conversions) {



-sub countSupportedSpecialFunctions 
+sub countSupportedSpecialFunctions
 {
    my $m = 0;

@@ -671,9 +684,9 @@ sub countSupportedSpecialFunctions
    return $m;
 }

-sub warnUnsupportedSpecialFunctions 
+sub warnUnsupportedSpecialFunctions
 {
-    my $line_num =  shift; 
+    my $line_num =  shift;

    my $m = 0;

@@ -0,0 +1,207 @@
+hsa_amd_coherency_get_type
+hsa_amd_coherency_set_type
+hsa_amd_profiling_get_dispatch_time
+hsa_amd_profiling_get_async_copy_time
+hsa_amd_image_get_info_max_dim
+hsa_amd_queue_cu_set_mask
+hsa_amd_agent_iterate_memory_pools
+hsa_amd_memory_pool_free
+hsa_amd_agent_memory_pool_get_info
+hsa_amd_memory_migrate
+hsa_amd_memory_lock
+hsa_amd_memory_unlock
+hsa_amd_interop_map_buffer
+hsa_amd_interop_unmap_buffer
+hsa_amd_image_create
+hsa_ext_program_create
+hsa_ext_program_destroy
+hsa_ext_program_add_module
+hsa_ext_program_iterate_modules
+hsa_ext_program_get_info
+hsa_ext_program_finalize
+hsa_ext_image_data_get_info
+hsa_ext_sampler_create
+hsa_status_string
+hsa_init
+hsa_shut_down
+hsa_system_get_info
+hsa_extension_get_name
+hsa_system_extension_supported
+hsa_system_major_extension_supported
+hsa_system_get_extension_table
+hsa_system_get_major_extension_table
+hsa_agent_get_info
+hsa_iterate_agents
+hsa_agent_get_info_size
+hsa_agent_set_info
+hsa_agent_get_exception_policies
+hsa_cache_get_info
+hsa_agent_iterate_caches
+hsa_agent_extension_supported
+hsa_agent_major_extension_supported
+hsa_signal_create
+hsa_signal_destroy
+hsa_signal_load_scacquire
+hsa_signal_load_relaxed
+hsa_signal_load_acquire
+hsa_signal_store_relaxed
+hsa_signal_store_screlease
+hsa_signal_store_release
+hsa_signal_silent_store_relaxed
+hsa_signal_silent_store_screlease
+hsa_signal_exchange_scacq_screl
+hsa_signal_exchange_acq_rel
+hsa_signal_exchange_scacquire
+hsa_signal_exchange_acquire
+hsa_signal_exchange_relaxed
+hsa_signal_exchange_screlease
+hsa_signal_exchange_release
+hsa_signal_cas_scacq_screl
+hsa_signal_cas_acq_rel
+hsa_signal_cas_scacquire
+hsa_signal_cas_acquire
+hsa_signal_cas_relaxed
+hsa_signal_cas_screlease
+hsa_signal_cas_release
+hsa_signal_add_scacq_screl
+hsa_signal_add_acq_rel
+hsa_signal_add_scacquire
+hsa_signal_add_acquire
+hsa_signal_add_relaxed
+hsa_signal_add_screlease
+hsa_signal_add_release
+hsa_signal_subtract_scacq_screl
+hsa_signal_subtract_acq_rel
+hsa_signal_subtract_scacquire
+hsa_signal_subtract_acquire
+hsa_signal_subtract_relaxed
+hsa_signal_subtract_screlease
+hsa_signal_subtract_release
+hsa_signal_and_scacq_screl
+hsa_signal_and_acq_rel
+hsa_signal_and_scacquire
+hsa_signal_and_acquire
+hsa_signal_and_relaxed
+hsa_signal_and_screlease
+hsa_signal_and_release
+hsa_signal_or_scacq_screl
+hsa_signal_or_acq_rel
+hsa_signal_or_scacquire
+hsa_signal_or_acquire
+hsa_signal_or_relaxed
+hsa_signal_or_screlease
+hsa_signal_or_release
+hsa_signal_xor_scacq_screl
+hsa_signal_xor_acq_rel
+hsa_signal_xor_scacquire
+hsa_signal_xor_acquire
+hsa_signal_xor_relaxed
+hsa_signal_xor_screlease
+hsa_signal_xor_release
+hsa_signal_wait_scacquire
+hsa_signal_wait_relaxed
+hsa_signal_wait_acquire
+hsa_signal_group_create
+hsa_signal_group_destroy
+hsa_signal_group_wait_any_scacquire
+hsa_signal_group_wait_any_relaxed
+hsa_queue_create
+hsa_soft_queue_create
+hsa_queue_destroy
+hsa_queue_inactivate
+hsa_queue_load_read_index_acquire
+hsa_queue_load_read_index_scacquire
+hsa_queue_load_read_index_relaxed
+hsa_queue_load_write_index_acquire
+hsa_queue_load_write_index_scacquire
+hsa_queue_load_write_index_relaxed
+hsa_queue_store_write_index_relaxed
+hsa_queue_store_write_index_release
+hsa_queue_store_write_index_screlease
+hsa_queue_cas_write_index_acq_rel
+hsa_queue_cas_write_index_scacq_screl
+hsa_queue_cas_write_index_acquire
+hsa_queue_cas_write_index_scacquire
+hsa_queue_cas_write_index_relaxed
+hsa_queue_cas_write_index_release
+hsa_queue_cas_write_index_screlease
+hsa_queue_add_write_index_acq_rel
+hsa_queue_add_write_index_scacq_screl
+hsa_queue_add_write_index_acquire
+hsa_queue_add_write_index_scacquire
+hsa_queue_add_write_index_relaxed
+hsa_queue_add_write_index_release
+hsa_queue_add_write_index_screlease
+hsa_queue_store_read_index_relaxed
+hsa_queue_store_read_index_release
+hsa_queue_store_read_index_screlease
+hsa_region_get_info
+hsa_agent_iterate_regions
+hsa_memory_allocate
+hsa_memory_free
+hsa_memory_copy
+hsa_memory_assign_agent
+hsa_memory_register
+hsa_memory_deregister
+hsa_isa_from_name
+hsa_agent_iterate_isas
+hsa_isa_get_info
+hsa_isa_get_info_alt
+hsa_isa_get_exception_policies
+hsa_isa_get_round_method
+hsa_wavefront_get_info
+hsa_isa_iterate_wavefronts
+hsa_isa_compatible
+hsa_code_object_reader_create_from_file
+hsa_code_object_reader_create_from_memory
+hsa_code_object_reader_destroy
+hsa_executable_create
+hsa_executable_create_alt
+hsa_executable_destroy
+hsa_executable_load_program_code_object
+hsa_executable_load_agent_code_object
+hsa_executable_freeze
+hsa_executable_get_info
+hsa_executable_global_variable_define
+hsa_executable_agent_global_variable_define
+hsa_executable_readonly_variable_define
+hsa_executable_validate
+hsa_executable_validate_alt
+hsa_executable_get_symbol
+hsa_executable_get_symbol_by_name
+hsa_executable_symbol_get_info
+hsa_executable_iterate_symbols
+hsa_executable_iterate_agent_symbols
+hsa_executable_iterate_program_symbols
+hsa_code_object_serialize
+hsa_code_object_deserialize
+hsa_code_object_destroy
+hsa_code_object_get_info
+hsa_executable_load_code_object
+hsa_code_object_get_symbol
+hsa_code_object_get_symbol_from_name
+hsa_code_symbol_get_info
+hsa_code_object_iterate_symbols
+hsa_ven_amd_loader_query_host_address
+hsa_ven_amd_loader_query_segment_descriptors
+hsa_amd_profiling_set_profiler_enabled
+hsa_amd_profiling_async_copy_enable
+hsa_amd_profiling_convert_tick_to_system_domain
+hsa_amd_signal_async_handler
+hsa_amd_async_function
+hsa_amd_signal_wait_any
+hsa_amd_memory_pool_get_info
+hsa_amd_memory_pool_allocate
+hsa_amd_memory_async_copy
+hsa_amd_agents_allow_access
+hsa_amd_memory_pool_can_migrate
+hsa_amd_memory_fill
+
+hsa_ext_image_get_capability
+hsa_ext_image_create
+hsa_ext_image_destroy
+hsa_ext_image_copy
+hsa_ext_image_import
+hsa_ext_image_export
+hsa_ext_image_clear
+hsa_ext_sampler_destroy
@@ -1,31 +1,211 @@
-# Profiling HIP Code
+# Profiling and Debugging HIP Code

-HIP provides several capabilities to support debugging and profiling.  Profiling information can be displayed to stderr or viewed in the CodeXl visualization tool.
+This section describes the profiling and debugging capabilities that HIP provides.  
+Profiling information can viewed in the CodeXL visualization tool or printed directly to stderr as the application runs.
+This document starts with some of the general capabilities of CodeXL and then describes some of the additional HIP marker and debug features.

-### Usign CodeXL to profile a HIP Application
-By defauly, CodeXL can trace all kernel commands, data transfer commands, and HSA Runtime (ROCr) API calls.  
-/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
+ * [CodeXL Profiling](#codexl-profiling)
+   * [Collecting and Viewing Traces](#collecting-and-viewing-traces)
+     * [Using rocm-profiler timestamp profiling](#using-rocm-profiler-timestamp-profiling)
+     * [Using rocm-profiler performance counter collection:](#using-rocm-profiler-performance-counter-collection)
+     * [Using CodeXL to view profiling results:](#using-codexl-to-view-profiling-results)
+     * [More information on CodeXL](#more-information-on-codexl)
+   * [HIP Markers](#hip-markers)
+     * [Profiling HIP APIs](#profiling-hip-apis)
+     * [Adding markers to applications](#adding-markers-to-applications)
+   * [Additional HIP Profiling Features](#additional-hip-profiling-features)
+     * [Demangling C   Kernel Names](#demangling-c-kernel-names)
+     * [Controlling when profiling starts and ends](#controlling-when-profiling-starts-and-ends)
+     * [Reducing timeline trace output file size](#reducing-timeline-trace-output-file-size)
+     * [How to enable profiling at HIP build time](#how-to-enable-profiling-at-hip-build-time)
+ * [Tracing and Debug](#tracing-and-debug)
+   * [Tracing HIP APIs](#tracing-hip-apis)
+     * [Color](#color)
+   * [Using HIP_DB](#using-hip_db)
+   * [Using ltrace](#using-ltrace)
+   * [Chicken bits](#chicken-bits)

-### Using CodeXL markers for HIP Functions
-HIP can generate markers at function being/end which are displayed on the CodeXL timeline view.
+## CodeXL Profiling
+
+### Collecting and Viewing Traces
+
+#### Using rocm-profiler timestamp profiling
+rocm-profiler is a command-line tool for tracing any application that uses ROCr API, including HCC and HIP.
+rocm-profiler's timeline trace will show the beginning and end for all kernel commands, data transfer commands, and HSA Runtime (ROCr) API calls.  The trace results are saved into a file, which by convention uses the "atp" extension.  Here is an example that shows how to run the command-line profiler:
+```shell
+$ /opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A -T  <applicationName> <applicationArguments>
+```
+
+#### Using rocm-profiler performance counter collection:
+rocm-profiler can record performance counter information to provide greater insight inside a kernel, such as the memory bandwidth, ALU busy percentage, and cache statistics.
+Collecting the common set of useful counters requires passing the counter configuration files for two passes:
+```
+$ /opt/rocm/bin/rocm-profiler -C -O --counterfile /opt/rocm/profiler/counterfiles/counters_HSA_Fiji_pass1 --counterfile /opt/rocm/profiler/counterfiles/counters_HSA_Fiji_pass2  <applicationName> <applicationArguments>
+```
+
+
+#### Using CodeXL to view profiling results:
+The trace can be loaded and viewed in the CodeXL visualization tool:
+
+- Open the CodeXL GUI, create an new project, and switch to "Profile Mode":
+  - $ CodeXL & 
+  - [File->New Project, leave fields as is, just click "OK"]
+  - [Profile->Switch to Profile Mode]
+- Load timestamp tracing results into a timeline view:
+  - Right click on the project in the CodeXL Explorer view
+  - Click "Import Session..."
+  - Select to $HOME/apitrace.atp (or appropriate .atp file if you used another file name)
+
+- Load the performance counter results
+  - Right click on the project in the CodeXL Explorer view
+  - Click "Import Session..."
+  - Select  $HOME/Session1.csv  (or appropriate .csv file if you used another file name)
+
+
+#### More information on CodeXL
+rocm-profiler --help will show additional options and usage guidelines.
+
+See this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) for more information on profiling ROCm apps (including HIP) with CodeXL.
+
+The 2.2 version of Windows CodeXL does not correctly handle Linux line-endings.  If you are collecting a trace on Linux and then viewing it with the 2.2 Windows CodeXL, first convert the line ending in the .atp file to Windows-style line endings.
+
+### HIP Markers
+#### Profiling HIP APIs  
+HIP can generate markers at function beginning and end which are displayed on the CodeXL timeline view.
 HIP 1.0 compiles marker support by default, and you can enable it by setting the HIP_PROFILE_API environment variable and then running the rocm-profiler:

 ```shell

 # Use profile to generate timeline view:
 export HIP_PROFILE_API=1
-/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
+$ /opt/rocm/bin/rocm-profiler -A -T <applicationName> <applicationArguments>

 Or
-/opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -o <outputATPFileName> -A <applicationName> <applicationArguments>
+$ /opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -A -T <applicationName> <applicationArguments>
 ```

-#### Developer Builds
+HIP_PROFILE_API supports two levels of information.
+- HIP_PROFILE_API=1 : Short format.  Print name of API but no arguments.  For example:  
+`hipMemcpy`
+- HIP_PROFILE_API=2 : Long format.  Print name of API + values of all function arguments. For example:  
+`hipMemcpy (0x7f32154db010, 0x50446e000, 4000000, hipMemcpyDeviceToHost)`
+
+#### Adding markers to applications
+
+Markers can be used to define application-specific events that will be recorded in the ATP file and displayed in the CodeXL gui.
+This can be particularly useful for visualizing how the higher-level phases of application behavior relate to the lower level HIP APIs, kernel launches, and data transfers.
+For example, an instrumented machine learning framework could show the beginning and ending of each layer in the network.
+
+Markers have a specific begin and end time, and can be nested.  Nested calls are displayed hierarchically in the CodeXL gui, with each level of the hierarchy occupying a different row.
+
+The HIP APis are defined in "hip_profile.h":
+```
+#include <hip/hip_profile.h>
+
+HIP_BEGIN_MARKER(const char *markerName, const char *groupName);
+HIP_END_MARKER();
+
+HIP_BEGIN_MARKER("Setup", "MyAppGroup");
+// ...
+// application code for setup
+// ...
+HIP_END_MARKER();
+```
+
+For C++ codes, HIP also provides a scoped marker which records the start time when constructed and the end time when the scoped marker is destructed at the end of the scope.  This provides a convenient, single-line mechanism to record an event that neatly corresponds to a region of code.  
+
+```cxx
+void FunctionFoo(...) 
+{
+  HIP_SCOPED_MARKER("FunctionFoo", "MyAppGroup"); // Marker starts recording here.
+
+  // ...
+  // Function implementation
+  // ...
+
+  // Marker destroyed here and records end time stamp.
+};
+``` 
+
+The HIP marker API is only supported on ROCm platform.  The marker macros are defined on CUDA platforms and will compile, but are silently ignored at runtime.
+
+This [HIP sample](samples/2_Cookbook/2_Profiler/) shows the profiler marker API used in a small application.
+
+More information on the marker API can be found in the profiler header file and PDF in a ROCM installation:
+- /opt/rocm/profiler/CXLActivityLogger/include/CXLActivityLogger.h
+- /opt/rocm/profiler/CXLActivityLogger/doc/CXLActivityLogger.pdf
+
+### Additional HIP Profiling Features
+#### Demangling C++ Kernel Names
+HIP includes the `hipdemangleatp` tool which can post-process an ATP file to "demangle" C++ names.
+Mangled kernel names encode the C++ arguments and other information, and are guaranteed to be unique even for cases such as operator overloading.  However, the mangled names can be quite verbose.  For example:
+
+`ZZ39gemm_NoTransA_MICRO_NBK_M_N_K_TS16XMTS4RN2hc16accelerator_viewEPKflS3_lPfliiiiiiffEN3_EC__719__cxxamp_trampolineElililiiiiiiS3_iS3_S4_ff`
+
+`hipdemangleatp` will convert this into the more readable:
+`gemm_NoTransA_MICRO_NBK_M_N_K_TS16XMTS4`
+
+The `hipdemangleatp` tool operates on the ATP file "in-place" and thus replaces the input file with the demangled version.  
+
+```
+$ hipdemangleatp myfile.atp
+```
+
+The kernel name is also shown in some of the summary htlm files (Top10 kernels).  These can be regenerated from the demangled ATP file by re-running rocm-profiler:
+```
+$ rocm-profiler -T --atpfile myfile.atp
+```
+
+A future version of CodeXL may directly integrate demangle functionality.
+
+
+#### Controlling when profiling starts and ends
+hipProfilerStart() and hipProfilerEnd() can be inserted into an application to control which phases of the applications are profiled.
+These APIs can be used to skip initialization code or to focus profiling on a desired region, and are particularly useful for large long-running applications.
+See the API documentation for more information.  These APIs work on both ROCm and CUDA paths.
+
+On ROCm, the following environment variables can be used to control when profiling occurs:
+
+```
+HIP_DB_START_API  : Comma-separated list of tid.api_seq_num for when to start debug and profiling.
+HIP_DB_STOP_API   : Comma-separated list of tid.api_seq_num for when to stop debug and profiling.
+```
+
+HIP/ROCm assigns a monotonically increasing sequence number to the APIs called from each thread.  The thread and API sequence number can be used in the above API to control when tracing starts and stops.  These flags also control the HIP_DB messages (described below).
+
+When using these options, start the profiler with profiling disabled:
+```
+# ROCm:
+$ rocm-profiler --startdisabled ...
+
+# CUDA:
+$ nvprof --profile-from-start-off ...
+```
+
+This feature is under development.
+
+#### Reducing timeline trace output file size
+If the application is already recording the HIP APIs, the HSA APIs are somewhat redundant and the ATP file size can be substantially reduced by not recording these APIs.  HIP includes a text file that lists all of the HSA APis and can assist in this filtering:
+
+```
+$ rocm-profiler -F hip/bin/hsa-api-filter-cxl.txt 
+```
+
+This file can be copied and edited to provide more selective HSA event recording.
+
+
+#### How to enable profiling at HIP build time
+Recent pre-built packages of HIP are always built with profiling support enabled.
 For developer builds, you must enable marker support manually when compiling HIP.

 1. Build HIP with ATP markers enabled
 HIP pre-built packages are enabled with ATP marker support by default.
-To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step.
+To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step.   Build and install HIP.
+```shell
+$ mkdir build && cd build
+$ cmake .. -DCOMPILE_HIP_ATP_MARKER
+$ make install
+```

 2. Install ROCm-Profiler
 Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well.
@@ -36,20 +216,57 @@ Alternatively, you can build ROCm-Profiler using the instructions [here](https:/
 Then follow the steps above to collect a marker-enabled trace.


-### Using HIP_TRACE_API
-You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided
-by the HIP_DB switch. For example:
-```shell
-# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function.
-HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp
+## Tracing and Debug
+
+### Tracing HIP APIs
+The HIP runtime can print the HIP function strings to stderr using HIP_TRACE_API environment variable.
+The trace prints two messages for each API - one at the beginning of the API call (line starts with "<<") and one at the end of the API call (line ends with ">>").
+Here's an example for one API followed by a description for the sections of the trace:
+
+```
+<<hip-api tid:1.6 hipMemcpy (0x7f32154db010, 0x50446e000, 4000000, hipMemcpyDeviceToHost)
+  hip-api tid:1.6 hipMemcpy                      ret= 0 (hipSuccess)>>
 ```

-#### Color
-Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.  
-You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable.  Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White.
-None will disable use of color control codes and may be useful when saving the trace file or when a pure text trace is desired.
+- `<<hip-api` is the header used for all HIP API debug messages.  The message is also shown in a specific color.  This can be used to distinguish this API from other HIP or application messages.
+-  `tid:1.6` indicates that this API call came from thread #1 and is the 6th API call in that thread.   When the first API in a new thread is called, HIP will associates a short sequential ID with that thread.  You can see the full thread ID (reported by C++) as 0x7f6183b097c0 in the example below.  
+- `hipMemcpy` is the name of the API.
+- The first line then prints a comma-separated list of the arguments to the function.  APIs which return values to the caller by writing to pointers will show the pointer addresses rather than the pointer contents.  This behavior may change in the future.
+- The second line shows the completion of the API, including the numeric return value (`ret= 0`) as well as an string representation for the error code (`hipSuccess`).  If the returned error code is non-zero, then the csecond line message is shown in red (unless HIP_TRACE_API_COLOR is "none" - see below).
+
+
+Heres a specific example showing the output of the [square](samples/0_Intro/square) program running on HIP:
+
+```
+$ HIP_TRACE_API=1  ./square.hip.out 
+  hip-api tid:1:HIP initialized short_tid#1 (maps to full_tid: 0x7f6183b097c0)
+<<hip-api tid:1.1 hipGetDeviceProperties (0x7ffddb673e08, 0)
+  hip-api tid:1.1 hipGetDeviceProperties         ret= 0 (hipSuccess)>>
+info: running on device gfx803
+info: allocate host mem (  7.63 MB)
+info: allocate device mem (  7.63 MB)
+<<hip-api tid:1.2 hipMalloc (0x7ffddb673fb8, 4000000)
+  hip-api tid:1.2 hipMalloc                      ret= 0 (hipSuccess)>>
+<<hip-api tid:1.3 hipMalloc (0x7ffddb673fb0, 4000000)
+  hip-api tid:1.3 hipMalloc                      ret= 0 (hipSuccess)>>
+info: copy Host2Device
+<<hip-api tid:1.4 hipMemcpy (0x50409d000, 0x7f32158ac010, 4000000, hipMemcpyHostToDevice)
+  hip-api tid:1.4 hipMemcpy                      ret= 0 (hipSuccess)>>
+info: launch 'vector_square' kernel
+1.5 hipLaunchKernel 'HIP_KERNEL_NAME(vector_square)' gridDim:{512,1,1} groupDim:{256,1,1} sharedMem:+0 stream#0.0
+info: copy Device2Host
+<<hip-api tid:1.6 hipMemcpy (0x7f32154db010, 0x50446e000, 4000000, hipMemcpyDeviceToHost)
+  hip-api tid:1.6 hipMemcpy                      ret= 0 (hipSuccess)>>
+info: check result
+PASSED!
+```
+
+
+#### Color
+Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
+You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable.  Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White.
+None will disable use of color control codes for both the opening and closing and may be useful when saving the trace file or when a pure text trace is desired.

-#### 


 ### Using HIP_DB
@@ -57,39 +274,60 @@ None will disable use of color control codes and may be useful when saving the t
 This flag is primarily targeted to assist HIP development team in the development of the HIP runtime, but in some situations may be useful to HIP application developers as well.
 The HIP debug information is designed to print important information during the execution of a HIP API.  HIP provides
 different color-coded levels of debug informaton:
-  - api  : Print the beginning and end of each HIP API, including the arguments and return codes.
+  - api  : Print the beginning and end of each HIP API, including the arguments and return codes.  This is equivalent to setting HIP_TRACE_API=1.
  - sync : Print multi-thread and other synchronization debug information.
  - copy : Print which engine is doing the copy, which copy flavor is selected, information on source and destination memory.
  - mem  : Print information about memory allocation - which pointers are allocated, where they are allocated, peer mappings, and more.

 DB_MEM format is flags separated by '+' sign, or a hex code for the bitmask.  Generally the + format is preferred.  
 For example:
-```shell
-HIP_DB=api+copy+mem  my-application
-HIP_DB=0xF  my-application
 ```
-HIP_DB=1 same as HIP_TRACE_API=1
+$ HIP_DB=api+copy+mem  my-application
+$ HIP_DB=0xF  my-application
+```
+
+### Using ltrace
+ltrace is a standard linux tool which provides a message to stderr on every dynamic library call.  Since ROCr and the ROCt (the ROC thunk, which is the thin user-space interface to the ROC kernel driver) are both dynamic libraries, this provides an easy way to trace the activity in these libraries.  Tracing can be a powerful way to quickly observe the flow of the application before diving into the details with a command-line debugger.
+The trace can also show performance issues related to accidental calls to expensive API calls on the critical path.
+
+ltrace can be easily combined with the HIP_DB switches to visualize the runtime behavior of the entire ROCm software stack.  Here's a sample command-line and output:
+
+```
+$ HIP_DB=api ltrace -C -e 'hsa*'   <applicationName> <applicationArguments>
+
+...
+
+<<hip-api tid:1.17 hipMemcpy (0x7f7776d3e010, 0x503d1d000, 4194304, hipMemcpyDeviceToHost)
+libmcwamp_hsa.so->hsa_signal_store_relaxed(0x1804000, 0, 0, 0x400000) = 0
+libmcwamp_hsa.so->hsa_signal_store_relaxed(0x1816000, 0, 0x7f777f85f2a0, 0x400000) = 0
+libmcwamp_hsa.so->hsa_amd_memory_lock(0x7f7776d3e010, 0x400000, 0x1213b70, 1 <unfinished ...>
+libhsa-runtime64.so.1->hsaKmtRegisterMemoryToNodes(0x7f7776d3e010, 0x400000, 1, 0x1220c10) = 0
+libhsa-runtime64.so.1->hsaKmtMapMemoryToGPUNodes(0x7f7776d3e010, 0x400000, 0x7ffc32865400, 64) = 0
+<... hsa_amd_memory_lock resumed> )              = 0
+libmcwamp_hsa.so->hsa_signal_store_relaxed(0x1804000, 1, 0x7f777e95a770, 0x12205b0) = 0
+libmcwamp_hsa.so->hsa_amd_memory_async_copy(0x50411d010, 0x11e70d0, 0x503d1d000, 0x11e70d0) = 0
+libmcwamp_hsa.so->hsa_signal_wait_acquire(0x1804000, 2, 1, -1) = 0
+libmcwamp_hsa.so->hsa_amd_memory_unlock(0x7f7776d3e010, 0x1213c6c, 0x12c3c600000000, 0x1804000 <unfinished ...>
+libhsa-runtime64.so.1->hsaKmtUnmapMemoryToGPU(0x7f7776d3e010, 0x7f7776d3e010, 0x12c3c600000000, 0x1804000) = 0
+libhsa-runtime64.so.1->hsaKmtDeregisterMemory(0x7f7776d3e010, 0x7f7776d3e010, 0x7f777f60f9e8, 0x1220580) = 0
+<... hsa_amd_memory_unlock resumed> )            = 0
+  hip-api tid:1.17 hipMemcpy                      ret= 0 (hipSuccess)>>
+```
+
+Some key information from the trace above.
+  - The trace snippet shows the execution of a hipMemcpy API, bracketed by the first and last message in the trace output.  The messages show the thread id and API sequence number (`1.17`).  ltrace output intermixes messages from all threads, so the HIP debug information can be useful to determine which threads are executing.
+  - The code flows through HIP APIs into ROCr (HSA) APIs (hsa*) and into the thunk (hsaKmt*) calls.
+  - The HCC runtime is "libmcwamp_hsa.so" and the HSA/ROCr runtime is "libhsa-runtime64.so".
+  - In this particular case, the memory copy is for unpinned memory, and the selected copy algorithm is to pin the host memory "in-place" before performing the copy.  The signaling APIs and calls to pin ("lock", "register") the memory are readily apparent in the trace output.


+### Chicken bits
+Chicken bits are environment variables which cause the HIP, HCC, or HSA driver to disable some feature or optimization.
+These are not intended for production but can be useful diagnose synchronization problems in the application (or driver).

+Some of the most useful chicken bits are described here:

-Trace provides quick look at API.
-Explain output of 
-Reference the cookbook example.
-Command-line profile.
-/// disable profiling at the start of the application you can start CodeXLGpuProfiler with the --startdisabled flag.
-
-Can use strace interleaved with HSA Debug calls .
-
-HIP_PROFILE_API=1
-HIP_PROFILE_API=2  : Will show the full API in the trace.  This can be useful for lower-level debugging when you want to see all the parameters that are passed to a specific API.
-
-demangle atp
-
-Write how to collect performance counters.
- include how to compute bandwidth for copy and kernel activity.
-
- How to disable HSA APIs.
- Do I need to use profiler with HSA enabled?  Do I need to enable HSA profiling on the command line?
-
-Offline compile, how to visualize.  
+- HIP_LAUNCH_BLOCKING=1 : On ROCm, this flag waits on the host after each kernel launches and after each memory copy command.  On CUDA, the waits are only enforced after each kernel launch.  This is useful to isolate synchronization problems.  Specifically, if the code works with this flag set, then it indicates the kernels and memory management code are correct, and any failures likely are causes by improper or missing synchronization.
+- HSA_ENABLE_SDMA=0     : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines.  Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine.  This flag is useful to isolate issues with the hardware copy engines.
+- HSA_ENABLE_INTERRUPT=0 : Causes completion signals to be detected with memory-based polling rather than interrupts.  Can be useful to diagnose interrupt storm issues in the driver.
+- HSA_DISABLE_CACHE=1  : Disables the GPU L2 data cache.
@@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */

-#ifndef HCC_H
-#define HCC_H
+#ifndef HIP_HCC_H
+#define HIP_HCC_H

 #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
 #include "hip/hcc_detail/hcc_acc.h"
@@ -58,17 +58,13 @@ THE SOFTWARE.
 #error (HCC must support GRID_LAUNCH_20)
 #endif

-#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01)
-#define HIP_LAUNCH_PARAM_BUFFER_SIZE    ((void*) 0x02)
-#define HIP_LAUNCH_PARAM_END            ((void*) 0x03)
-
 extern int HIP_TRACE_API;

 //TODO-HCC-GL - change this to typedef.
 //typedef grid_launch_parm hipLaunchParm ;
 #define hipLaunchParm grid_launch_parm
 #ifdef __cplusplus
-#include <hip/hcc_detail/hip_texture.h>
+//#include <hip/hcc_detail/hip_texture.h>
 #include <hip/hcc_detail/hip_ldg.h>
 #endif
 #include <hip/hcc_detail/host_defines.h>
@@ -134,7 +130,6 @@ __device__ float atanhf(float x);
 __device__ float cbrtf(float x);
 __device__ float ceilf(float x);
 __device__ float copysignf(float x, float y);
-__device__ float cosf(float x);
 __device__ float coshf(float x);
 __device__ float cyl_bessel_i0f(float x);
 __device__ float cyl_bessel_i1f(float x);
@@ -146,9 +141,7 @@ __host__ float erfcxf(float x);
 __device__ float erff(float x);
 __device__ float erfinvf(float y);
 __host__ float erfinvf(float y);
-__device__ float exp10f(float x);
 __device__ float exp2f(float x);
-__device__ float expf(float x);
 __device__ float expm1f(float x);
 __device__ float fabsf(float x);
 __device__ float fdimf(float x, float y);
@@ -171,11 +164,8 @@ __device__ float ldexpf(float x, int exp);
 __device__ float lgammaf(float x);
 __device__ long long int llrintf(float x);
 __device__ long long int llroundf(float x);
-__device__ float log10f(float x);
 __device__ float log1pf(float x);
-__device__ float log2f(float x);
 __device__ float logbf(float x);
-__device__ float logf(float x);
 __device__ long int lrintf(float x);
 __device__ long int lroundf(float x);
 __device__ float modff(float x, float *iptr);
@@ -191,7 +181,6 @@ __host__ float normcdff(float y);
 __device__ float normcdfinvf(float y);
 __host__ float normcdfinvf(float y);
 __device__ float normf(int dim, const float *a);
-__device__ float powf(float x, float y);
 __device__ float rcbrtf(float x);
 __host__ float rcbrtf(float x);
 __device__ float remainderf(float x, float y);
@@ -210,14 +199,11 @@ __device__ float rsqrtf(float x);
 __device__ float scalblnf(float x, long int n);
 __device__ float scalbnf(float x, int n);
 __host__ __device__ unsigned signbit(float a);
-__device__ void sincosf(float x, float *sptr, float *cptr);
 __device__ void sincospif(float x, float *sptr, float *cptr);
 __host__ void sincospif(float x, float *sptr, float *cptr);
-__device__ float sinf(float x);
 __device__ float sinhf(float x);
 __device__ float sinpif(float x);
 __device__ float sqrtf(float x);
-__device__ float tanf(float x);
 __device__ float tanhf(float x);
 __device__ float tgammaf(float x);
 __device__ float truncf(float x);
@@ -482,24 +468,251 @@ __host__ __device__ int max(int arg1, int arg2);
 __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();

 //TODO - add a couple fast math operations here, the set here will grow :
-__device__ float __cosf(float x);
-__device__ float __expf(float x);
-__device__ float __frsqrt_rn(float x);
-__device__ float __fsqrt_rd(float x);
-__device__ float __fsqrt_rn(float x);
-__device__ float __fsqrt_ru(float x);
-__device__ float __fsqrt_rz(float x);
-__device__ float __log10f(float x);
-__device__ float __log2f(float x);
-__device__ float __logf(float x);
-__device__ float __powf(float base, float exponent);
-__device__ void __sincosf(float x, float *s, float *c) ;
-__device__ float __sinf(float x);
-__device__ float __tanf(float x);
-__device__ float __dsqrt_rd(double x);
-__device__ float __dsqrt_rn(double x);
-__device__ float __dsqrt_ru(double x);
-__device__ float __dsqrt_rz(double x);
+
+// Single Precision Precise Math
+__device__ float __hip_precise_cosf(float);
+__device__ float __hip_precise_exp10f(float);
+__device__ float __hip_precise_expf(float);
+__device__ float __hip_precise_frsqrt_rn(float);
+__device__ float __hip_precise_fsqrt_rd(float);
+__device__ float __hip_precise_fsqrt_rn(float);
+__device__ float __hip_precise_fsqrt_ru(float);
+__device__ float __hip_precise_fsqrt_rz(float);
+__device__ float __hip_precise_log10f(float);
+__device__ float __hip_precise_log2f(float);
+__device__ float __hip_precise_logf(float);
+__device__ float __hip_precise_powf(float, float);
+__device__ void __hip_precise_sincosf(float,float*,float*);
+__device__ float __hip_precise_sinf(float);
+__device__ float __hip_precise_tanf(float);
+
+// Double Precision Precise Math
+__device__ double __hip_precise_dsqrt_rd(double);
+__device__ double __hip_precise_dsqrt_rn(double);
+__device__ double __hip_precise_dsqrt_ru(double);
+__device__ double __hip_precise_dsqrt_rz(double);
+
+// Single Precision Fast Math
+extern __attribute__((const)) float __hip_fast_cosf(float) __asm("llvm.cos.f32");
+extern __attribute__((const)) float __hip_fast_exp2f(float) __asm("llvm.exp2.f32");
+__device__ float __hip_fast_exp10f(float);
+__device__ float __hip_fast_expf(float);
+__device__ float __hip_fast_frsqrt_rn(float);
+extern __attribute__((const)) float __hip_fast_fsqrt_rd(float) __asm("llvm.sqrt.f32");
+__device__ float __hip_fast_fsqrt_rn(float);
+__device__ float __hip_fast_fsqrt_ru(float);
+__device__ float __hip_fast_fsqrt_rz(float);
+__device__ float __hip_fast_log10f(float);
+extern __attribute__((const)) float __hip_fast_log2f(float) __asm("llvm.log2.f32");
+__device__ float __hip_fast_logf(float);
+__device__ float __hip_fast_powf(float, float);
+__device__ void __hip_fast_sincosf(float,float*,float*);
+extern __attribute__((const)) float __hip_fast_sinf(float) __asm("llvm.sin.f32");
+__device__ float __hip_fast_tanf(float);
+extern __attribute__((const)) float __hip_fast_fmaf(float,float,float) __asm("llvm.fma.f32");
+extern __attribute__((const)) float __hip_fast_frcp(float) __asm("llvm.amdgcn.rcp.f32");
+
+extern __attribute__((const)) double __hip_fast_dsqrt(double) __asm("llvm.sqrt.f64");
+extern __attribute__((const)) double __hip_fast_fma(double,double,double) __asm("llvm.fma.f64");
+extern __attribute__((const)) double __hip_fast_drcp(double) __asm("llvm.amdgcn.rcp.f64");
+
+#ifdef HIP_FAST_MATH
+// Single Precision Precise Math when enabled
+
+__device__ inline float cosf(float x) {
+  return __hip_fast_cosf(x);
+}
+
+__device__ inline float exp10f(float x) {
+  return __hip_fast_exp10f(x);
+}
+
+__device__ inline float expf(float x) {
+  return __hip_fast_expf(x);
+}
+
+__device__ inline float log10f(float x) {
+  return __hip_fast_log10f(x);
+}
+
+__device__ inline float log2f(float x) {
+  return __hip_fast_log2f(x);
+}
+
+__device__ inline float logf(float x) {
+  return __hip_fast_logf(x);
+}
+
+__device__ inline float powf(float base, float exponent) {
+  return __hip_fast_powf(base, exponent);
+}
+
+__device__ inline void sincosf(float x, float *s, float *c) {
+  return __hip_fast_sincosf(x, s, c);
+}
+
+__device__ inline float sinf(float x) {
+  return __hip_fast_sinf(x);
+}
+
+__device__ inline float tanf(float x) {
+  return __hip_fast_tanf(x);
+}
+
+#else
+
+__device__ float sinf(float);
+__device__ float cosf(float);
+__device__ float tanf(float);
+__device__ void sincosf(float, float*, float*);
+__device__ float logf(float);
+__device__ float log2f(float);
+__device__ float log10f(float);
+__device__ float expf(float);
+__device__ float exp10f(float);
+__device__ float powf(float, float);
+
+#endif
+// Single Precision Fast Math
+__device__ inline float __cosf(float x) {
+  return __hip_fast_cosf(x);
+}
+
+__device__ inline float __exp10f(float x) {
+  return __hip_fast_exp10f(x);
+}
+
+__device__ inline float __expf(float x) {
+  return __hip_fast_expf(x);
+}
+
+__device__ inline float __frsqrt_rn(float x) {
+  return __hip_fast_frsqrt_rn(x);
+}
+
+__device__ inline float __fsqrt_rd(float x) {
+  return __hip_fast_fsqrt_rd(x);
+}
+
+__device__ inline float __fsqrt_rn(float x) {
+  return __hip_fast_fsqrt_rn(x);
+}
+
+__device__ inline float __fsqrt_ru(float x) {
+  return __hip_fast_fsqrt_ru(x);
+}
+
+__device__ inline float __fsqrt_rz(float x) {
+  return __hip_fast_fsqrt_rz(x);
+}
+
+__device__ inline float __log10f(float x) {
+  return __hip_fast_log10f(x);
+}
+
+__device__ inline float __log2f(float x) {
+  return __hip_fast_log2f(x);
+}
+
+__device__ inline float __logf(float x) {
+  return __hip_fast_logf(x);
+}
+
+__device__ inline float __powf(float base, float exponent) {
+  return __hip_fast_powf(base, exponent);
+}
+
+__device__ inline void __sincosf(float x, float *s, float *c) {
+  return __hip_fast_sincosf(x, s, c);
+}
+
+__device__ inline float __sinf(float x) {
+  return __hip_fast_sinf(x);
+}
+
+__device__ inline float __tanf(float x) {
+  return __hip_fast_tanf(x);
+}
+
+__device__ inline float __fmaf_rd(float x, float y, float z) {
+  return __hip_fast_fmaf(x, y, z);
+}
+
+__device__ inline float __fmaf_rn(float x, float y, float z) {
+  return __hip_fast_fmaf(x, y, z);
+}
+
+__device__ inline float __fmaf_ru(float x, float y, float z) {
+  return __hip_fast_fmaf(x, y, z);
+}
+
+__device__ inline float __fmaf_rz(float x, float y, float z) {
+  return __hip_fast_fmaf(x, y, z);
+}
+
+__device__ inline float __frcp_rd(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline float __frcp_rn(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline float __frcp_ru(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline float __frcp_rz(float x) {
+  return __hip_fast_frcp(x);
+}
+
+__device__ inline double __dsqrt_rd(double x) {
+  return __hip_fast_dsqrt(x);
+}
+
+__device__ inline double __dsqrt_rn(double x) {
+  return __hip_fast_dsqrt(x);
+}
+
+__device__ inline double __dsqrt_ru(double x) {
+  return __hip_fast_dsqrt(x);
+}
+
+__device__ inline double __dsqrt_rz(double x) {
+  return __hip_fast_dsqrt(x);
+}
+
+__device__ inline double __fma_rd(double x, double y, double z) {
+  return __hip_fast_fma(x, y, z);
+}
+
+__device__ inline double __fma_rn(double x, double y, double z) {
+  return __hip_fast_fma(x, y, z);
+}
+
+__device__ inline double __fma_ru(double x, double y, double z) {
+  return __hip_fast_fma(x, y, z);
+}
+
+__device__ inline double __fma_rz(double x, double y, double z) {
+  return __hip_fast_fma(x, y, z);
+}
+
+__device__ inline double __drcp_rd(double x) {
+  return __hip_fast_drcp(x);
+}
+
+__device__ inline double __drcp_rn(double x) {
+  return __hip_fast_drcp(x);
+}
+
+__device__ inline double __drcp_ru(double x) {
+  return __hip_fast_drcp(x);
+}
+
+__device__ inline double __drcp_rz(double x) {
+  return __hip_fast_drcp(x);
+}

 /**
 * CUDA 8 device function features
@@ -592,40 +805,18 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
 #define hipGridDim_y   (hc_get_num_groups(1))
 #define hipGridDim_z   (hc_get_num_groups(2))

-// loop unrolling
-__device__ static inline void* memcpy(void* dst, void* src, size_t size)
-{
-    uint8_t *dstPtr, *srcPtr;
-    dstPtr = (uint8_t*)dst;
-    srcPtr = (uint8_t*)src;
-    for(uint32_t i=0;i<size;i++) {
-        dstPtr[i] = srcPtr[i];
-    }
-    return nullptr;
-}
-
-__device__ static inline void* memset(void* ptr, uint8_t val, size_t size)
-{
-    uint8_t *dstPtr;
-    dstPtr = (uint8_t*)ptr;
-    for(uint32_t i=0;i<size;i++) {
-        dstPtr[i] = val;
-    }
-    return nullptr;
-}
+//extern "C" __device__ void* memcpy(void* dst, void* src, size_t size);
+//extern "C" __device__ void* memset(void* ptr, uint8_t val, size_t size);

 extern "C" __device__ void* __hip_hc_malloc(size_t);
 extern "C" __device__ void* __hip_hc_free(void *ptr);

-__device__ static inline void* malloc(size_t size)
-{
-    return __hip_hc_malloc(size);
-}
+//extern "C" __device__ void* malloc(size_t size);
+//extern "C" __device__ void* free(void *ptr);

-__device__ static inline void* free(void *ptr)
-{
-    return __hip_hc_free(ptr);
-}
+extern "C" __device__ char4 __hip_hc_add8pk(char4, char4);
+extern "C" __device__ char4 __hip_hc_sub8pk(char4, char4);
+extern "C" __device__ char4 __hip_hc_mul8pk(char4, char4);

 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)

@@ -39,6 +39,10 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif

+#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01)
+#define HIP_LAUNCH_PARAM_BUFFER_SIZE    ((void*) 0x02)
+#define HIP_LAUNCH_PARAM_END            ((void*) 0x03)
+
 // Structure definitions:
 #ifdef __cplusplus
 extern "C" {
@@ -108,12 +112,12 @@ enum hipLimit_t
 /**
 * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored.
 */
-typedef enum hipFuncCache {
+typedef enum hipFuncCache_t {
    hipFuncCachePreferNone, ///< no preference for shared memory or L1 (default)
    hipFuncCachePreferShared, ///< prefer larger shared memory and smaller L1 cache
    hipFuncCachePreferL1, ///< prefer larger L1 cache and smaller shared memory
    hipFuncCachePreferEqual, ///< prefer equal size L1 cache and shared memory
-} hipFuncCache;
+} hipFuncCache_t;


 /**
@@ -317,7 +321,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId);
 * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache.  This hint is ignored on those architectures.
 *
 */
-hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig );
+hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig );


 /**
@@ -329,7 +333,7 @@ hipError_t hipDeviceSetCacheConfig ( hipFuncCache cacheConfig );
 * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache.  This hint is ignored on those architectures.
 *
 */
-hipError_t hipDeviceGetCacheConfig ( hipFuncCache *cacheConfig );
+hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t *cacheConfig );

 /**
 * @brief Get Resource limits of current device
@@ -353,7 +357,7 @@ hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit_t limit);
 * Note: AMD devices and recent Nvidia GPUS do not support reconfigurable cache.  This hint is ignored on those architectures.
 *
 */
-hipError_t hipFuncSetCacheConfig ( hipFuncCache config );
+hipError_t hipFuncSetCacheConfig ( hipFuncCache_t config );

 /**
 * @brief Returns bank width of shared memory for current device
@@ -409,7 +413,7 @@ hipError_t hipSetDeviceFlags ( unsigned flags);
 *
 * @returns #hipSuccess, #hipErrorInvalidValue
 */
-hipError_t hipChooseDevice(int *device, hipDeviceProp_t* prop);
+hipError_t hipChooseDevice(int *device, const hipDeviceProp_t* prop);

 // end doxygen Device
 /**
@@ -1456,7 +1460,7 @@ hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion);
 *
 * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
 */
-hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig );
+hipError_t hipCtxGetCacheConfig ( hipFuncCache_t *cacheConfig );

 /**
 * @brief Set L1/Shared cache partition.
@@ -1469,7 +1473,7 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig );
 *
 * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
 */
-hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig );
+hipError_t hipCtxSetCacheConfig ( hipFuncCache_t cacheConfig );

 /**
 * @brief Set Shared memory bank configuration.
@@ -1757,6 +1761,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 // TODO - expand descriptions:
 /**
 * @brief Start recording of profiling information
+ * When using this API, start the profiler with profiling disabled.  (--startdisabled)
 * @warning : hipProfilerStart API is under development.
 */
 hipError_t hipProfilerStart();
@@ -1764,6 +1769,7 @@ hipError_t hipProfilerStart();

 /**
 * @brief Stop recording of profiling information.
+ * When using this API, start the profiler with profiling disabled.  (--startdisabled)
 * @warning : hipProfilerStop API is under development.
 */
 hipError_t hipProfilerStop();
@@ -22,8 +22,8 @@ THE SOFTWARE.

 //#pragma once

-#ifndef HIP_TEXTURE_H
-#define HIP_TEXTURE_H
+#ifndef HIP_HCC_DETAIL_TEXTURE_H
+#define HIP_HCC_DETAIL_TEXTURE_H

 /**
 *  @file  hcc_detail/hip_texture.h
@@ -32,7 +32,7 @@ THE SOFTWARE.

 #include <limits.h>

-#include <hip/hip_runtime.h>
+//#include <hip/hcc_detail/hip_runtime.h>

 //----
 //Texture - TODO - likely need to move this to a separate file only included with kernel compilation.
@@ -84,7 +84,7 @@ struct texture : public textureReference {
 };
 #endif

-typedef struct hipArray {
+typedef struct {
  unsigned int width;
  unsigned int height;
  hipChannelFormatKind f;
@@ -32,7 +32,7 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif

-#if __HCC__
+#if 0
 #include <hc_short_vector.hpp>

 using namespace hc::short_vector;
@@ -137,8 +137,24 @@ struct uchar3
    unsigned char x, y, z;
 };

-struct __hip_align(char4, 4, signed char x; signed char y; signed char z; signed char w;);
-struct __hip_align(uchar4, 4, unsigned char x; unsigned char y; unsigned char z; unsigned char w;);
+struct char4
+{
+    union {
+        signed char x, y, z, w;
+        unsigned int val;
+    };
+};
+
+struct uchar4
+{
+    union {
+        unsigned char x, y, z, w;
+        unsigned int val;
+    };
+};
+
+//struct __hip_align(char4, 4, signed char x; signed char y; signed char z; signed char w;);
+//struct __hip_align(uchar4, 4, unsigned char x; unsigned char y; unsigned char z; unsigned char w;);

 struct __hip_align(short1, 2, signed short x;);
 struct __hip_align(ushort1, 2, unsigned short x;);
@@ -0,0 +1,37 @@
+/*
+Copyright (c) 2015-2016 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.
+*/
+
+
+
+#ifndef HIP_HIP_TEXTURE_H
+#define HIP_HIP_TEXTURE_H
+
+#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
+#include <hip/hcc_detail/hip_texture.h>
+#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
+#include <hip/nvcc_detail/hip_texture.h>
+#else 
+#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
+#endif 
+
+
+#endif
@@ -70,12 +70,14 @@ hipMemcpyHostToHost
 #define HIP_LAUNCH_PARAM_BUFFER_SIZE     CU_LAUNCH_PARAM_BUFFER_SIZE
 #define HIP_LAUNCH_PARAM_END            CU_LAUNCH_PARAM_END
 #define hipLimitMallocHeapSize          cudaLimitMallocHeapSize
+#define hipIpcMemLazyEnablePeerAccess          cudaIpcMemLazyEnablePeerAccess

 typedef cudaEvent_t hipEvent_t;
 typedef cudaStream_t hipStream_t;
 typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
 typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
 typedef cudaLimit hipLimit_t;
+typedef cudaFuncCache hipFuncCache_t;
 typedef CUcontext hipCtx_t;
 typedef CUsharedconfig hipSharedMemConfig;
 typedef CUfunc_cache hipFuncCache;
@@ -323,6 +325,10 @@ inline static hipError_t hipDeviceSynchronize() {
    return hipCUDAErrorTohipError(cudaDeviceSynchronize());
 }

+inline static hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *pCacheConfig) {
+    return hipCUDAErrorTohipError(cudaDeviceGetCacheConfig(pCacheConfig));
+}
+
 inline static const char* hipGetErrorString(hipError_t error){
    return cudaGetErrorString(hipErrorToCudaError(error));
 }
@@ -0,0 +1,6 @@
+#ifndef HIP_TEXTURE_H
+#define HIP_TEXTURE_H
+
+#include <texture_types.h>
+
+#endif
@@ -7,6 +7,7 @@ elseif(@HIP_LIB_TYPE@ EQUAL 1)
    install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.a DESTINATION lib)
 else()
    install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
+    install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
 endif()
 install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
 install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib)
@@ -0,0 +1,27 @@
+
+Copyright (c) 2011, UT-Battelle, LLC
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are met:
+
+* Redistributions of source code must retain the above copyright
+  notice, this list of conditions and the following disclaimer.
+* Redistributions in binary form must reproduce the above copyright
+  notice, this list of conditions and the following disclaimer in the
+  documentation and/or other materials provided with the distribution.
+* Neither the name of Oak Ridge National Laboratory, nor UT-Battelle, LLC, nor
+  the names of its contributors may be used to endorse or promote products
+  derived from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
@@ -0,0 +1,33 @@
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+	HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=hipCommander
+OPT=-O3
+#CXXFLAGS = -O3 -g
+CXXFLAGS =  $(OPT) --std=c++11
+
+HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
+ifeq (${HIP_PLATFORM}, hcc)
+	CXXFLAGS += " -stdlib=libc++"
+endif
+
+CODE_OBJECTS=nullkernel.hsaco
+
+all: ${EXE} ${CODE_OBJECTS}
+
+$(EXE): hipCommander.cpp 
+	$(HIPCC) $(CXXFLAGS) $^ -o $@
+
+nullkernel.hsaco : nullkernel.hip.cpp
+	$(HIPCC) --genco nullkernel.hip -o nullkernel.hsaco
+
+
+install: $(EXE)
+	cp $(EXE) $(HIP_PATH)/bin
+
+
+clean:
+	rm -f *.o *.co $(EXE)
@@ -0,0 +1,527 @@
+#include "ResultDatabase.h"
+
+#include <cfloat>
+#include <algorithm>
+#include <cmath>
+#include <iomanip>
+
+using namespace std;
+
+bool ResultDatabase::Result::operator<(const Result &rhs) const
+{
+    if (test < rhs.test)
+        return true;
+    if (test > rhs.test)
+        return false;
+    if (atts < rhs.atts)
+        return true;
+    if (atts > rhs.atts)
+        return false;
+    return false; // less-operator returns false on equal
+}
+
+double ResultDatabase::Result::GetMin() const
+{
+    double r = FLT_MAX;
+    for (int i=0; i<value.size(); i++)
+    {
+        r = min(r, value[i]);
+    }
+    return r;
+}
+
+double ResultDatabase::Result::GetMax() const
+{
+    double r = -FLT_MAX;
+    for (int i=0; i<value.size(); i++)
+    {
+        r = max(r, value[i]);
+    }
+    return r;
+}
+
+double ResultDatabase::Result::GetMedian() const
+{
+    return GetPercentile(50);
+}
+
+double ResultDatabase::Result::GetPercentile(double q) const
+{
+    int n = value.size();
+    if (n == 0)
+        return FLT_MAX;
+    if (n == 1)
+        return value[0];
+
+    if (q <= 0)
+        return value[0];
+    if (q >= 100)
+        return value[n-1];
+
+    double index = ((n + 1.) * q / 100.) - 1;
+
+    vector<double> sorted = value;
+    sort(sorted.begin(), sorted.end());
+
+    if (n == 2)
+        return (sorted[0] * (1 - q/100.)  +  sorted[1] * (q/100.));
+
+    int index_lo = int(index);
+    double frac = index - index_lo;
+    if (frac == 0)
+        return sorted[index_lo];
+
+    double lo = sorted[index_lo];
+    double hi = sorted[index_lo + 1];
+    return lo + (hi-lo)*frac;
+}
+
+double ResultDatabase::Result::GetMean() const
+{
+    double r = 0;
+    for (int i=0; i<value.size(); i++)
+    {
+        r += value[i];
+    }
+    return r / double(value.size());
+}
+
+double ResultDatabase::Result::GetStdDev() const
+{
+    double r = 0;
+    double u = GetMean();
+    if (u == FLT_MAX)
+        return FLT_MAX;
+    for (int i=0; i<value.size(); i++)
+    {
+        r += (value[i] - u) * (value[i] - u);
+    }
+    r = sqrt(r / value.size());
+    return r;
+}
+
+
+void ResultDatabase::AddResults(const string &test,
+                                const string &atts,
+                                const string &unit,
+                                const vector<double> &values)
+{
+    for (int i=0; i<values.size(); i++)
+    {
+        AddResult(test, atts, unit, values[i]);
+    }
+}
+
+static string RemoveAllButLeadingSpaces(const string &a)
+{
+    string b;
+    int n = a.length();
+    int i = 0;
+    while (i<n && a[i] == ' ')
+    {
+        b += a[i];
+        ++i;
+    }
+    for (; i<n; i++)
+    {
+        if (a[i] != ' ' && a[i] != '\t')
+            b += a[i];
+    }
+    return b;
+}
+
+void ResultDatabase::AddResult(const string &test_orig,
+                               const string &atts_orig,
+                               const string &unit_orig,
+                               double value)
+{
+    string test = RemoveAllButLeadingSpaces(test_orig);
+    string atts = RemoveAllButLeadingSpaces(atts_orig);
+    string unit = RemoveAllButLeadingSpaces(unit_orig);
+    int index;
+    for (index = 0; index < results.size(); index++)
+    {
+        if (results[index].test == test &&
+            results[index].atts == atts)
+        {
+            if (results[index].unit != unit)
+                throw "Internal error: mixed units";
+
+            break;
+        }
+    }
+
+    if (index >= results.size())
+    {
+        Result r;
+        r.test = test;
+        r.atts = atts;
+        r.unit = unit;
+        results.push_back(r);
+    }
+
+    results[index].value.push_back(value);
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::DumpDetailed
+//
+//  Purpose:
+//    Writes the full results, including all trials.
+//
+//  Arguments:
+//    out        where to print
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    August 14, 2009
+//
+//  Modifications:
+//    Jeremy Meredith, Wed Nov 10 14:25:17 EST 2010
+//    Renamed to DumpDetailed to make room for a DumpSummary.
+//
+//    Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
+//    Added note about (*) missing value tag.
+//
+//    Jeremy Meredith, Tue Nov 23 13:57:02 EST 2010
+//    Changed note about missing values to be worded a little better.
+//
+// ****************************************************************************
+void ResultDatabase::DumpDetailed(ostream &out)
+{
+    vector<Result> sorted(results);
+    sort(sorted.begin(), sorted.end());
+
+    const int testNameW = 24 ;
+    const int attW = 12;
+    const int fieldW = 11;
+    out << std::fixed << right << std::setprecision(4);
+
+    int maxtrials = 1;
+    for (int i=0; i<sorted.size(); i++)
+    {
+        if (sorted[i].value.size() > maxtrials)
+            maxtrials = sorted[i].value.size();
+    }
+
+    // TODO: in big parallel runs, the "trials" are the procs
+    // and we really don't want to print them all out....
+    out << setw(testNameW) << "test\t"  
+        << setw(attW) << "atts\t"
+        << setw(fieldW) 
+        << "median\t"
+        << "mean\t"
+        << "stddev\t"
+        << "min\t"
+        << "max\t";
+    for (int i=0; i<maxtrials; i++)
+        out << "trial"<<i<<"\t";
+    out << endl;
+
+    for (int i=0; i<sorted.size(); i++)
+    {
+        Result &r = sorted[i];
+        out << setw(testNameW) << r.test + "\t";
+        out << setw(attW) << r.atts + "\t";
+        out << setw(fieldW) << r.unit + "\t";
+        if (r.GetMedian() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMedian() << "\t";
+        if (r.GetMean() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMean()   << "\t";
+        if (r.GetStdDev() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetStdDev() << "\t";
+        if (r.GetMin() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMin()    << "\t";
+        if (r.GetMax() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMax()    << "\t";
+        for (int j=0; j<r.value.size(); j++)
+        {
+            if (r.value[j] == FLT_MAX)
+                out << "N/A\t";
+            else
+                out << r.value[j] << "\t";
+        }
+
+        out << endl;
+    }
+    out << endl
+        << "Note: Any results marked with (*) had missing values." << endl
+        << "      This can occur on systems with a mixture of" << endl
+        << "      device types or architectural capabilities." << endl;
+}
+
+
+// ****************************************************************************
+//  Method:  ResultDatabase::DumpDetailed
+//
+//  Purpose:
+//    Writes the summary results (min/max/stddev/med/mean), but not
+//    every individual trial.
+//
+//  Arguments:
+//    out        where to print
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    November 10, 2010
+//
+//  Modifications:
+//    Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
+//    Added note about (*) missing value tag.
+//
+// ****************************************************************************
+void ResultDatabase::DumpSummary(ostream &out)
+{
+    vector<Result> sorted(results);
+    sort(sorted.begin(), sorted.end());
+
+    const int testNameW = 24 ;
+    const int attW = 12;
+    const int fieldW = 9;
+    out << std::fixed << right << std::setprecision(4);
+
+    // TODO: in big parallel runs, the "trials" are the procs
+    // and we really don't want to print them all out....
+    out << setw(testNameW) << "test\t"  
+        << setw(attW) << "atts\t"
+        << setw(fieldW) 
+        << "units\t"
+        << "median\t"
+        << "mean\t"
+        << "stddev\t"
+        << "min\t"
+        << "max\t";
+    out << endl;
+
+    for (int i=0; i<sorted.size(); i++)
+    {
+        Result &r = sorted[i];
+        out << setw(testNameW) << r.test + "\t";
+        out << setw(attW) << r.atts + "\t";
+        out << setw(fieldW) << r.unit + "\t";
+        if (r.GetMedian() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMedian() << "\t";
+        if (r.GetMean() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMean()   << "\t";
+        if (r.GetStdDev() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetStdDev() << "\t";
+        if (r.GetMin() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMin()    << "\t";
+        if (r.GetMax() == FLT_MAX)
+            out << "N/A\t";
+        else
+            out << r.GetMax()    << "\t";
+
+        out << endl;
+    }
+    out << endl
+        << "Note: results marked with (*) had missing values such as" << endl
+        << "might occur with a mixture of architectural capabilities." << endl;
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::ClearAllResults
+//
+//  Purpose:
+//    Clears all existing results from the ResultDatabase; used for multiple passes
+//    of the same test or multiple tests.
+//
+//  Arguments:
+//
+//  Programmer:  Jeffrey Young
+//  Creation:    September 10th, 2014
+//
+//  Modifications:
+//
+//
+// ****************************************************************************
+void ResultDatabase::ClearAllResults()
+{
+	results.clear();	
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::DumpCsv
+//
+//  Purpose:
+//    Writes either detailed or summary results (min/max/stddev/med/mean), but not
+//    every individual trial.
+//
+//  Arguments:
+//    out        file to print CSV results
+//
+//  Programmer:  Jeffrey Young
+//  Creation:    August 28th, 2014
+//
+//  Modifications:
+//
+// ****************************************************************************
+void ResultDatabase::DumpCsv(string fileName)
+{
+    bool emptyFile;
+    vector<Result> sorted(results);
+
+    sort(sorted.begin(), sorted.end());
+
+    //Check to see if the file is empty - if so, add the headers
+    emptyFile = this->IsFileEmpty(fileName);
+
+    //Open file and append by default
+    ofstream out;
+    out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app); 
+
+    //Add headers only for empty files
+    if(emptyFile)
+    {
+    // TODO: in big parallel runs, the "trials" are the procs
+    // and we really don't want to print them all out....
+    out << "test, "
+        << "atts, "
+        << "units, "
+        << "median, "
+        << "mean, "
+        << "stddev, "
+        << "min, "
+        << "max, ";
+    out << endl;
+    }
+
+    for (int i=0; i<sorted.size(); i++)
+    {
+        Result &r = sorted[i];
+        out << r.test << ", ";
+        out << r.atts << ", ";
+        out << r.unit << ", ";
+        if (r.GetMedian() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMedian() << ", ";
+        if (r.GetMean() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMean()   << ", ";
+        if (r.GetStdDev() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetStdDev() << ", ";
+        if (r.GetMin() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMin()    << ", ";
+        if (r.GetMax() == FLT_MAX)
+            out << "N/A, ";
+        else
+            out << r.GetMax()    << ", ";
+
+        out << endl;
+    }
+    out << endl;
+
+    out.close();
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::IsFileEmpty
+//
+//  Purpose:
+//    Returns whether a file is empty - used as a helper for CSV printing
+//
+//  Arguments:
+//    file  The input file to check for emptiness
+//
+//  Programmer:  Jeffrey Young
+//  Creation:    August 28th, 2014
+//
+//  Modifications:
+//
+// ****************************************************************************
+
+bool ResultDatabase::IsFileEmpty(string fileName)
+{
+      bool fileEmpty;
+
+      ifstream file(fileName.c_str());
+
+      //If the file doesn't exist it is by definition empty
+      if(!file.good())
+      {
+        return true;
+      }
+      else
+      {
+        fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
+        file.close();
+        
+	return fileEmpty;
+      }
+  
+      //Otherwise, return false  
+        return false;
+}
+
+
+
+// ****************************************************************************
+//  Method:  ResultDatabase::GetResultsForTest
+//
+//  Purpose:
+//    Returns a vector of results for just one test name.
+//
+//  Arguments:
+//    test       the name of the test results to search for
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    December  3, 2010
+//
+//  Modifications:
+//
+// ****************************************************************************
+vector<ResultDatabase::Result>
+ResultDatabase::GetResultsForTest(const string &test)
+{
+    // get only the given test results
+    vector<Result> retval;
+    for (int i=0; i<results.size(); i++)
+    {
+        Result &r = results[i];
+        if (r.test == test)
+            retval.push_back(r);
+    }
+    return retval;
+}
+
+// ****************************************************************************
+//  Method:  ResultDatabase::GetResults
+//
+//  Purpose:
+//    Returns all the results.
+//
+//  Arguments:
+//
+//  Programmer:  Jeremy Meredith
+//  Creation:    December  3, 2010
+//
+//  Modifications:
+//
+// ****************************************************************************
+const vector<ResultDatabase::Result> &
+ResultDatabase::GetResults() const
+{
+    return results;
+}
@@ -0,0 +1,100 @@
+#ifndef RESULT_DATABASE_H
+#define RESULT_DATABASE_H
+
+#include <string>
+#include <vector>
+#include <iostream>
+#include <fstream>
+#include <cfloat>
+using std::string;
+using std::vector;
+using std::ostream;
+using std::ofstream;
+using std::ifstream;
+
+
+// ****************************************************************************
+// Class:  ResultDatabase
+//
+// Purpose:
+//   Track numerical results as they are generated.
+//   Print statistics of raw results.
+//
+// Programmer:  Jeremy Meredith
+// Creation:    June 12, 2009
+//
+// Modifications:
+//    Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010
+//    Split timing reports into detailed and summary.  E.g. for serial code,
+//    we might report all trial values, but skip them in parallel.
+//
+//    Jeremy Meredith, Thu Nov 11 11:40:18 EST 2010
+//    Added check for missing value tag.
+//
+//    Jeremy Meredith, Mon Nov 22 13:37:10 EST 2010
+//    Added percentile statistic.
+//
+//    Jeremy Meredith, Fri Dec  3 16:30:31 EST 2010
+//    Added a method to extract a subset of results based on test name.  Also,
+//    the Result class is now public, so that clients can use them directly.
+//    Added a GetResults method as well, and made several functions const.
+//
+// ****************************************************************************
+class ResultDatabase
+{
+  public:
+    //
+    // A performance result for a single SHOC benchmark run.
+    //
+    struct Result
+    {
+        string test;  // e.g. "readback"
+        string atts;  // e.g. "pagelocked 4k^2"
+        string unit;  // e.g. "MB/sec"
+        vector<double> value; // e.g. "837.14"
+        double GetMin() const;
+        double GetMax() const;
+        double GetMedian() const;
+        double GetPercentile(double q) const;
+        double GetMean() const;
+        double GetStdDev() const;
+
+        bool operator<(const Result &rhs) const;
+
+        bool HadAnyFLTMAXValues() const
+        {
+            for (int i=0; i<value.size(); ++i)
+            {
+                if (value[i] >= FLT_MAX)
+                    return true;
+            }
+            return false;
+        }
+    };
+
+  protected:
+    vector<Result> results;
+
+  public:
+    void AddResult(const string &test,
+                   const string &atts,
+                   const string &unit,
+                   double value);
+    void AddResults(const string &test,
+                    const string &atts,
+                    const string &unit,
+                    const vector<double> &values);
+    vector<Result>        GetResultsForTest(const string &test);
+    const vector<Result> &GetResults() const;
+    void ClearAllResults();
+    void DumpDetailed(ostream&);
+    void DumpSummary(ostream&);
+    void DumpCsv(string fileName);
+
+  private:
+    bool IsFileEmpty(string fileName);
+
+};
+
+
+#endif
@@ -0,0 +1,50 @@
+_ Add AQL kernel.
+_ Fix &*kernel command so the kernel name/type is an argument not a new command.
+
+_ Add command to parse only.
+_ Add regression to parse all the hcm files.
+
+_ Partition HCC, HIP, HSA, OpenCL commands into separate files.
+
+
+_ Show time for back-to-back copies.
+_ Add variables.
+  %loopcnt
+
+    ./hipCommander %loopcnt=4  
+
+_ Add datasize command.
+
+
+_ Add ( ) to parsing.
+_ Add argument parsing and checking.
+
+_ Add verbose option to print each step of setup.
+  - print deliniater between setup and run.  Add run start message.
+
+  - print sizes of all buffers.
+  - print each command before running.
+  - show start/stop of timer routine.
+
+_ 
+_ Clear documentation on what each oepration does.
+_ Add time instrumentation for each command.
+_ Add pcie atomic.
+
+
+_ Add tests for negative cases, ie endloop w/o opening loop.
+
+
+README tips
+---
+- HIP_API_TRACE combined with -v is useful to track the exact commands generates by hipCommander.
+
+
+Other ideas:
+---
+[ ] Perf guide : stream creation very slow on HCC and should be avoided.
+
+
+Scratch:
+
+
@@ -0,0 +1,3 @@
+loop,1000; H2D; NullKernel; D2H; endloop;
+streamsync;
+printTiming, 1000
@@ -0,0 +1 @@
+H2D; NullKernel, D2H, streamsync
@@ -0,0 +1,3 @@
+setstream,1;
+NullKernel; streamsync;
+loop,10000; H2D;  NullKernel; streamsync; endloop,1;
@@ -0,0 +1,3 @@
+loop,1000; H2D; NullKernel; D2H; endloop;
+streamsync;
+printTiming, 1000
@@ -0,0 +1,2 @@
+setstream,1;
+loop,1000; NullKernel; syncstream; endloop,1,
@@ -0,0 +1,7 @@
+#include "hip/hip_runtime.h"
+
+extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad){
+    if (Ad) {
+        Ad[0] = 42;
+    }
+}
@@ -0,0 +1,10 @@
+setstream(1);
+NullKernel; streamsync;
+loop(30000); NullKernel; streamsync; endloop(1);
+loop(30000); H2D; H2D;  NullKernel; streamsync; endloop(1);
+loop(30000); H2D; H2D; H2D;  NullKernel; streamsync; endloop(1);
+
+loop(30000); H2D; NullKernel; D2H; streamsync; endloop(1);
+loop(30000); NullKernel; D2H; streamsync; endloop(1);
+loop(30000); NullKernel; D2H; D2H; streamsync; endloop(1);
+loop(30000); NullKernel; D2H; D2H; D2H; streamsync; endloop(1);
@@ -0,0 +1,8 @@
+setstream,1;
+NullKernel; streamsync;
+loop,100000; NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  streamsync; NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  streamsync; NullKernel; streamsync; D2H;  streamsync; endloop,1;
+
@@ -0,0 +1,5 @@
+setstream,1;
+NullKernel; streamsync;
+loop,100000; NullKernel; streamsync; endloop,1;
+loop,100000; H2D;  NullKernel; streamsync; endloop,1;
+loop,100000; H2D;  NullKernel; D2H;  streamsync; endloop,1;
@@ -0,0 +1,7 @@
+setstream,0;
+NullKernel; streamsync;
+loop,100000; NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  NullKernel; streamsync; endloop,1;
+
+loop,100000; H2D;  NullKernel; D2H;  streamsync; endloop,1;
@@ -0,0 +1,5 @@
+setstream(1);
+NullKernel; streamsync;
+loop(100);  ModuleKernel; streamsync; endloop(1);
+loop(100);  AqlKernel; streamsync; endloop(1);
+loop(3000); NullKernel; streamsync; endloop(1);
@@ -0,0 +1,3 @@
+setstream,1; 
+setstream,2;  H2D; NullKernel; D2H; 
+streamsync
@@ -0,0 +1,21 @@
+#include <hip/hip_runtime.h>
+
+static const int BLOCKSIZEX=32;
+static const int BLOCKSIZEY=16;
+
+__global__ void fails(hipLaunchParm lp, float* pErrorI)
+{
+    if(pErrorI!=0)
+    {
+        pErrorI[0]=1;
+    }
+}
+
+int main()
+{
+    dim3 blocks(1,1);
+    dim3 threads(BLOCKSIZEX,BLOCKSIZEY);
+    float error;
+
+    hipLaunchKernel(HIP_KERNEL_NAME(fails), blocks, threads, 0, 0, &error);
+}
@@ -1,36 +1,35 @@
-HIP_PATH?= $(wildcard /opt/rocm/hip)
-ifeq (,$(HIP_PATH))
-	HIP_PATH=../../..
-endif
-
-HIPCC=$(HIP_PATH)/bin/hipcc
-
-TARGET=hcc
-
-SOURCES = peer2peer.cpp
-OBJECTS = $(SOURCES:.cpp=.o)
-
-EXECUTABLE=./peer2peer
-
-.PHONY: test
-
-
-all: $(EXECUTABLE) test
-
-CXXFLAGS =-g
-CXX=$(HIPCC)
-
-
-$(EXECUTABLE): $(OBJECTS)
-	$(HIPCC) $(OBJECTS) -o $@
-
-
-test: $(EXECUTABLE)
-	$(EXECUTABLE)
-
-
-clean:
-	rm -f $(EXECUTABLE)
-	rm -f $(OBJECTS)
-	rm -f $(HIP_PATH)/src/*.o
-
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+    HIP_PATH=../../..
+endif
+
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+TARGET=hcc
+
+SOURCES = peer2peer.cpp
+OBJECTS = $(SOURCES:.cpp=.o)
+
+EXECUTABLE=./peer2peer
+
+.PHONY: test
+
+
+all: $(EXECUTABLE) test
+
+CXXFLAGS =-g
+CXX=$(HIPCC)
+
+
+$(EXECUTABLE): $(OBJECTS)
+    $(HIPCC) $(OBJECTS) -o $@
+
+test: $(EXECUTABLE)
+    $(EXECUTABLE)
+
+
+clean:
+rm -f $(EXECUTABLE)
+rm -f $(OBJECTS)
+rm -f $(HIP_PATH)/src/*.o
+
@@ -108,7 +108,7 @@ void disablePeer2Peer(int currentGpu, int peerGpu)

    HIPCHECK(hipSetDevice(currentGpu));
    hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
-	
+
    if(canAccessPeer==1){
        HIPCHECK(hipDeviceDisablePeerAccess(peerGpu));
    }
@@ -155,7 +155,7 @@ __global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
 int main(){

    checkPeer2PeerSupport();
-    
+
    int gpuCount;
    int currentGpu, peerGpu;

@@ -191,10 +191,10 @@ int main(){
                    0, 0,
                    gpuTransposeMatrix[0], data[0], width);

-	HIPCHECK(hipSetDevice(peerGpu));
-	TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
-	hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
-	hipMalloc((void**)&data[1], NUM * sizeof(float));
+    HIPCHECK(hipSetDevice(peerGpu));
+    TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
+    hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
+    hipMalloc((void**)&data[1], NUM * sizeof(float));
    hipMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), hipMemcpyDeviceToDevice);

    hipLaunchKernel(matrixTranspose_dynamic_shared,
@@ -26,6 +26,184 @@ THE SOFTWARE.

 #include "hip/hip_runtime.h"

+//=================================================================================================
+/*
+ Implementation of malloc and free device functions.
+
+ This is the best place to put them because the device
+ global variables need to be initialized at the start.
+*/
+
+#define NUM_PAGES_PER_THREAD  16
+#define SIZE_OF_PAGE          64
+#define NUM_THREADS_PER_CU    64
+#define NUM_CUS_PER_GPU       64
+#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
+#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
+#define SIZE_OF_HEAP SIZE_MALLOC
+
+size_t g_malloc_heap_size = SIZE_OF_HEAP;
+
+__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
+__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
+
+__device__ void *__hip_hc_malloc(size_t size)
+{
+    char *heap = (char*)gpuHeap;
+    if(size > SIZE_OF_HEAP)
+    {
+        return (void*)nullptr;
+    }
+    uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
+    uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+
+    uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
+    uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
+
+    uint32_t stride = size / SIZE_OF_PAGE;
+    uint32_t start = numHeapsPerWorkItem * currentWorkItem;
+
+    uint32_t k=0;
+
+    while(gpuFlags[k] > 0)
+    {
+        k++;
+    }
+
+    for(uint32_t i=0;i<stride-1;i++)
+    {
+        gpuFlags[i+start+k] = 1;
+    }
+
+    gpuFlags[start+stride-1+k] = 2;
+
+    void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
+
+    return ptr;
+}
+
+__device__ void* __hip_hc_free(void *ptr)
+{
+    if(ptr == nullptr)
+    {
+       return nullptr;
+    }
+
+    uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap;
+    uint32_t offsetPage = offsetByte / SIZE_OF_PAGE;
+
+    while(gpuFlags[offsetPage] != 0) {
+        if(gpuFlags[offsetPage] == 2) {
+            gpuFlags[offsetPage] = 0;
+            offsetPage++;
+            break;
+        } else {
+            gpuFlags[offsetPage] = 0;
+            offsetPage++;
+        }
+    }
+
+    return nullptr;
+}
+
+__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
+    return hc::__amdgcn_ds_bpermute(index, src);
+}
+
+__device__ float __hip_ds_bpermutef(int index, float src) {
+    return hc::__amdgcn_ds_bpermute(index, src);
+}
+
+__device__ unsigned __hip_ds_permute(int index, unsigned src) {
+    return hc::__amdgcn_ds_permute(index, src);
+}
+
+__device__ float __hip_ds_permutef(int index, float src) {
+    return hc::__amdgcn_ds_permute(index, src);
+}
+
+__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
+    return hc::__amdgcn_ds_swizzle(src, pattern);
+}
+
+__device__ float __hip_ds_swizzlef(float src, int pattern) {
+    return hc::__amdgcn_ds_swizzle(src, pattern);
+}
+
+__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
+    return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
+}
+
+#define MASK1 0x00ff00ff
+#define MASK2 0xff00ff00
+
+__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) {
+    char4 out;
+    unsigned one1 = in1.val & MASK1;
+    unsigned one2 = in2.val & MASK1;
+    out.val = (one1 + one2) & MASK1;
+    one1 = in1.val & MASK2;
+    one2 = in2.val & MASK2;
+    out.val = out.val | ((one1 + one2) & MASK2);
+    return out;
+}
+
+__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
+    char4 out;
+    unsigned one1 = in1.val & MASK1;
+    unsigned one2 = in2.val & MASK1;
+    out.val = (one1 - one2) & MASK1;
+    one1 = in1.val & MASK2;
+    one2 = in2.val & MASK2;
+    out.val = out.val | ((one1 - one2) & MASK2);
+    return out;
+}
+
+__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
+    char4 out;
+    unsigned one1 = in1.val & MASK1;
+    unsigned one2 = in2.val & MASK1;
+    out.val = (one1 * one2) & MASK1;
+    one1 = in1.val & MASK2;
+    one2 = in2.val & MASK2;
+    out.val = out.val | ((one1 * one2) & MASK2);
+    return out;
+}
+
+// loop unrolling
+__device__ void* memcpy(void* dst, void* src, size_t size)
+{
+    uint8_t *dstPtr, *srcPtr;
+    dstPtr = (uint8_t*)dst;
+    srcPtr = (uint8_t*)src;
+    for(uint32_t i=0;i<size;i++) {
+        dstPtr[i] = srcPtr[i];
+    }
+    return nullptr;
+}
+
+__device__ void* memset(void* ptr, uint8_t val, size_t size)
+{
+    uint8_t *dstPtr;
+    dstPtr = (uint8_t*)ptr;
+    for(uint32_t i=0;i<size;i++) {
+        dstPtr[i] = val;
+    }
+    return nullptr;
+}
+
+__device__ void* malloc(size_t size)
+{
+    return __hip_hc_malloc(size);
+}
+
+__device__ void* free(void *ptr)
+{
+    return __hip_hc_free(ptr);
+}
+
+//=================================================================================================
+
 // TODO: Choose whether default is precise math or fast math based on compilation flag.
 #ifdef __HCC_ACCELERATOR__
 using namespace hc::precise_math;
@@ -1865,26 +2043,151 @@ __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
 }


+// Precise Math Functions
+__device__ float __hip_precise_cosf(float x) {
+  return hc::precise_math::cosf(x);
+}

-//TODO - add a couple fast math operations here, the set here will grow :
-__device__  float __cosf(float x) {return hc::fast_math::cosf(x); };
-__device__  float __expf(float x) {return hc::fast_math::expf(x); };
-__device__  float __frsqrt_rn(float x) {return hc::fast_math::rsqrt(x); };
-__device__  float __fsqrt_rd(float x) {return hc::fast_math::sqrt(x); };
-__device__  float __fsqrt_rn(float x) {return hc::fast_math::sqrt(x); };
-__device__  float __fsqrt_ru(float x) {return hc::fast_math::sqrt(x); };
-__device__  float __fsqrt_rz(float x) {return hc::fast_math::sqrt(x); };
-__device__  float __log10f(float x) {return hc::fast_math::log10f(x); };
-__device__  float __log2f(float x) {return hc::fast_math::log2f(x); };
-__device__  float __logf(float x) {return hc::fast_math::logf(x); };
-__device__  float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); };
-__device__  void __sincosf(float x, float *s, float *c) {return hc::fast_math::sincosf(x, s, c); };
-__device__  float __sinf(float x) {return hc::fast_math::sinf(x); };
-__device__  float __tanf(float x) {return hc::fast_math::tanf(x); };
-__device__  float __dsqrt_rd(double x) {return hc::fast_math::sqrt(x); };
-__device__  float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); };
-__device__  float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); };
-__device__  float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
+__device__ float __hip_precise_exp10f(float x) {
+  return hc::precise_math::exp10f(x);
+}
+
+__device__ float __hip_precise_expf(float x) {
+  return hc::precise_math::expf(x);
+}
+
+__device__ float __hip_precise_frsqrt_rn(float x) {
+  return hc::precise_math::rsqrt(x);
+}
+
+__device__ float __hip_precise_fsqrt_rd(float x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ float __hip_precise_fsqrt_rn(float x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ float __hip_precise_fsqrt_ru(float x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ float __hip_precise_fsqrt_rz(float x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ float __hip_precise_log10f(float x) {
+  return hc::precise_math::log10(x);
+}
+
+__device__ float __hip_precise_log2f(float x) {
+  return hc::precise_math::log2(x);
+}
+
+__device__ float __hip_precise_logf(float x) {
+  return hc::precise_math::logf(x);
+}
+
+__device__ float __hip_precise_powf(float base, float exponent) {
+  return hc::precise_math::powf(base, exponent);
+}
+
+__device__ void __hip_precise_sincosf(float x, float *s, float *c) {
+  hc::precise_math::sincosf(x, s, c);
+}
+
+__device__ float __hip_precise_sinf(float x) {
+  return hc::precise_math::sinf(x);
+}
+
+__device__ float __hip_precise_tanf(float x) {
+  return hc::precise_math::tanf(x);
+}
+
+// Double Precision Math
+__device__ double __hip_precise_dsqrt_rd(double x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ double __hip_precise_dsqrt_rn(double x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ double __hip_precise_dsqrt_ru(double x) {
+  return hc::precise_math::sqrt(x);
+}
+
+__device__ double __hip_precise_dsqrt_rz(double x) {
+  return hc::precise_math::sqrt(x);
+}
+
+#define LOG_BASE2_E_DIV_2 0.4426950408894701
+#define LOG_BASE2_5 2.321928094887362
+#define ONE_DIV_LOG_BASE2_E 0.69314718056
+#define ONE_DIV_LOG_BASE2_10 0.30102999566
+
+// Fast Math Intrinsics
+__device__ float __hip_fast_exp10f(float x) {
+  return __hip_fast_exp2f(x*LOG_BASE2_E_DIV_2);
+}
+
+__device__ float __hip_fast_expf(float x) {
+  return __hip_fast_expf(x*LOG_BASE2_5);
+}
+
+__device__ float __hip_fast_frsqrt_rn(float x) {
+  return 1 / __hip_fast_fsqrt_rd(x);;
+}
+
+__device__ float __hip_fast_fsqrt_rn(float x) {
+  return __hip_fast_fsqrt_rd(x);
+}
+
+__device__ float __hip_fast_fsqrt_ru(float x) {
+  return __hip_fast_fsqrt_rd(x);
+}
+
+__device__ float __hip_fast_fsqrt_rz(float x) {
+  return __hip_fast_fsqrt_rd(x);
+}
+
+__device__ float __hip_fast_log10f(float x) {
+  return ONE_DIV_LOG_BASE2_E * __hip_fast_log2f(x);
+}
+
+__device__ float __hip_fast_logf(float x) {
+  return ONE_DIV_LOG_BASE2_10 * __hip_fast_log2f(x);
+}
+
+__device__ float __hip_fast_powf(float base, float exponent) {
+  return hc::fast_math::powf(base, exponent);
+}
+
+__device__ void __hip_fast_sincosf(float x, float *s, float *c) {
+  *s = __hip_fast_sinf(x);
+  *c = __hip_fast_cosf(x);
+}
+
+__device__ float __hip_fast_tanf(float x) {
+  return hc::fast_math::tanf(x);
+}
+
+// Double Precision Math
+__device__ double __hip_fast_dsqrt_rd(double x) {
+  return hc::fast_math::sqrt(x);
+}
+
+__device__ double __hip_fast_dsqrt_rn(double x) {
+  return hc::fast_math::sqrt(x);
+}
+
+__device__ double __hip_fast_dsqrt_ru(double x) {
+  return hc::fast_math::sqrt(x);
+}
+
+__device__ double __hip_fast_dsqrt_rz(double x) {
+  return hc::fast_math::sqrt(x);
+}

 __HIP_DEVICE__ char1 make_char1(signed char x)
 {
@@ -2294,6 +2597,10 @@ __HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
    return d4;
 }

+__device__ void  __threadfence_system(void){
+    // no-op
+}
+
 float __hip_host_erfinvf(float x)
 {
    float ret;
@@ -217,7 +217,7 @@ hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion)
    return ihipLogStatus(hipSuccess);
 }

-hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig )
+hipError_t hipCtxGetCacheConfig ( hipFuncCache_t *cacheConfig )
 {
    HIP_INIT_API(cacheConfig);

@@ -226,7 +226,7 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig )
    return ihipLogStatus(hipSuccess);
 }

-hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig )
+hipError_t hipCtxSetCacheConfig ( hipFuncCache_t cacheConfig )
 {
    HIP_INIT_API(cacheConfig);

@@ -44,17 +44,15 @@ hipError_t hipGetDevice(int *deviceId)
            *deviceId = ctx->getDevice()->_deviceId;
        }
    }else{
-        e = hipErrorInvalidDevice;
+        e = hipErrorInvalidValue;
    }

    return ihipLogStatus(e);
 }

 // TODO - does this initialize HIP runtime?
-hipError_t hipGetDeviceCount(int *count)
+hipError_t ihipGetDeviceCount(int *count)
 {
-    HIP_INIT_API(count);
-
    hipError_t e = hipSuccess;

    if(count != nullptr) {
@@ -66,12 +64,18 @@ hipError_t hipGetDeviceCount(int *count)
            e = ihipLogStatus(hipErrorNoDevice);
        }
    } else {
-        e = ihipLogStatus(hipErrorNoDevice);
+        e = ihipLogStatus(hipErrorInvalidValue);
    }
    return e;
 }

-hipError_t hipDeviceSetCacheConfig(hipFuncCache cacheConfig)
+hipError_t hipGetDeviceCount(int *count)
+{
+    HIP_INIT_API(count);
+    return ihipGetDeviceCount(count);
+}
+
+hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig)
 {
    HIP_INIT_API(cacheConfig);

@@ -80,10 +84,14 @@ hipError_t hipDeviceSetCacheConfig(hipFuncCache cacheConfig)
    return ihipLogStatus(hipSuccess);
 }

-hipError_t hipDeviceGetCacheConfig(hipFuncCache *cacheConfig)
+hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig)
 {
    HIP_INIT_API(cacheConfig);

+    if(cacheConfig == nullptr) {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+
    *cacheConfig = hipFuncCachePreferNone;

    return ihipLogStatus(hipSuccess);
@@ -105,7 +113,7 @@ hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit_t limit)
    }
 }

-hipError_t hipFuncSetCacheConfig (hipFuncCache cacheConfig)
+hipError_t hipFuncSetCacheConfig (hipFuncCache_t cacheConfig)
 {
    HIP_INIT_API(cacheConfig);

@@ -168,13 +176,13 @@ hipError_t hipDeviceReset(void)
    return ihipLogStatus(hipSuccess);
 }

-hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
+hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
 {
-    HIP_INIT_API(pi, attr, device);
-
    hipError_t e = hipSuccess;

-    if(pi != nullptr) {
+    if(pi == nullptr) {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }

    auto * hipDevice = ihipGetDevice(device);
    hipDeviceProp_t *prop = &hipDevice->_props;
@@ -236,16 +244,17 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
    } else {
        e = hipErrorInvalidDevice;
    }
-    }else{
-        e = hipErrorInvalidDevice;
-    }
    return ihipLogStatus(e);
 }

-hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
+hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
 {
-    HIP_INIT_API(props, device);
+    HIP_INIT_API(pi, attr, device);
+    return ihipDeviceGetAttribute(pi,attr,device);
+}

+hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device)
+{
    hipError_t e;

    if(props != nullptr){
@@ -264,6 +273,12 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
    return ihipLogStatus(e);
 }

+hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
+{
+    HIP_INIT_API(props, device);
+    return ihipGetDeviceProperties(props, device);
+}
+
 hipError_t hipSetDeviceFlags( unsigned int flags)
 {
    HIP_INIT_API(flags);
@@ -308,8 +323,8 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device
    HIP_INIT_API(major,minor, device);
    hipError_t e = hipSuccess;
    int deviceId= device->_deviceId;
-    e = hipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId);
-    e = hipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId);
+    e = ihipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId);
+    e = ihipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId);
    return ihipLogStatus(e);
 }

@@ -328,7 +343,7 @@ hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
    HIP_INIT_API(pciBusId,len, device);
    hipError_t e = hipSuccess;
    int deviceId= device->_deviceId;
-    e = hipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId);
+    e = ihipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId);
    return ihipLogStatus(e);
 }

@@ -346,10 +361,10 @@ hipError_t hipDeviceGetByPCIBusId (int*  device, const int* pciBusId )
    hipDeviceProp_t  tempProp;
    int deviceCount;
    hipError_t e = hipErrorInvalidValue;
-    hipGetDeviceCount( &deviceCount );
+    ihipGetDeviceCount( &deviceCount );
    *device = 0;
-    for (int i=0; i< deviceCount; i++) {
-        hipGetDeviceProperties( &tempProp, i );
+    for (int i = 0; i< deviceCount; i++) {
+        ihipGetDeviceProperties( &tempProp, i );
        if(tempProp.pciBusID == *pciBusId) {
            *device =i;
            e = hipSuccess;
@@ -364,19 +379,19 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop )
    HIP_INIT_API(device,prop);
    hipDeviceProp_t  tempProp;
    int deviceCount;
-    int inPropCount=0;
-    int matchedPropCount=0;
+    int inPropCount = 0;
+    int matchedPropCount = 0;
    hipError_t e = hipSuccess;
-    hipGetDeviceCount( &deviceCount );
+    ihipGetDeviceCount( &deviceCount );
    *device = 0;
-    for (int i=0; i< deviceCount; i++) {
-        hipGetDeviceProperties( &tempProp, i );
-        if(prop->major !=0) {
+    for (int i = 0; i < deviceCount; i++) {
+        ihipGetDeviceProperties( &tempProp, i );
+        if(prop->major != 0) {
            inPropCount++;
            if(tempProp.major >= prop->major) {
                matchedPropCount++;
            }
-            if(prop->minor !=0) {
+            if(prop->minor != 0) {
                inPropCount++;
                 if(tempProp.minor >= prop->minor) {
                     matchedPropCount++;
@@ -395,31 +410,31 @@ hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop )
                matchedPropCount++;
            }
        }
-        if(prop->maxThreadsPerBlock  != 0) {
+        if(prop->maxThreadsPerBlock != 0) {
            inPropCount++;
            if(tempProp.maxThreadsPerBlock >= prop->maxThreadsPerBlock ) {
                matchedPropCount++;
            }
        }
-        if(prop->totalConstMem  != 0) {
+        if(prop->totalConstMem != 0) {
            inPropCount++;
            if(tempProp.totalConstMem >= prop->totalConstMem ) {
                matchedPropCount++;
            }
        }
-        if(prop->multiProcessorCount  != 0) {
+        if(prop->multiProcessorCount != 0) {
            inPropCount++;
            if(tempProp.multiProcessorCount >= prop->multiProcessorCount ) {
                matchedPropCount++;
            }
        }
-        if(prop->maxThreadsPerMultiProcessor  != 0) {
+        if(prop->maxThreadsPerMultiProcessor != 0) {
            inPropCount++;
            if(tempProp.maxThreadsPerMultiProcessor >= prop->maxThreadsPerMultiProcessor ) {
                matchedPropCount++;
            }
        }
-        if(prop->memoryClockRate  != 0) {
+        if(prop->memoryClockRate != 0) {
            inPropCount++;
            if(tempProp.memoryClockRate >= prop->memoryClockRate ) {
                matchedPropCount++;
@@ -109,114 +109,6 @@ std::atomic<int> g_lastShortTid(1);
 std::vector<ProfTrigger> g_dbStartTriggers;
 std::vector<ProfTrigger> g_dbStopTriggers;

-
-
-/*
- Implementation of malloc and free device functions.
-
- This is the best place to put them because the device
- global variables need to be initialized at the start.
-*/
-
-#define NUM_PAGES_PER_THREAD  16
-#define SIZE_OF_PAGE          64
-#define NUM_THREADS_PER_CU    64
-#define NUM_CUS_PER_GPU       64
-#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
-#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
-#define SIZE_OF_HEAP SIZE_MALLOC
-
-size_t g_malloc_heap_size = SIZE_OF_HEAP;
-
-__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
-__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
-
-__device__ void *__hip_hc_malloc(size_t size)
-{
-    char *heap = (char*)gpuHeap;
-    if(size > SIZE_OF_HEAP)
-    {
-        return (void*)nullptr;
-    }
-    uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
-    uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
-
-    uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
-    uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
-
-    uint32_t stride = size / SIZE_OF_PAGE;
-    uint32_t start = numHeapsPerWorkItem * currentWorkItem;
-
-    uint32_t k=0;
-
-    while(gpuFlags[k] > 0)
-    {
-        k++;
-    }
-
-    for(uint32_t i=0;i<stride-1;i++)
-    {
-        gpuFlags[i+start+k] = 1;
-    }
-
-    gpuFlags[start+stride-1+k] = 2;
-
-    void* ptr = (void*)(heap + heapSizePerWorkItem * currentWorkItem + k*SIZE_OF_PAGE);
-
-    return ptr;
-}
-
-__device__ void* __hip_hc_free(void *ptr)
-{
-    if(ptr == nullptr)
-    {
-       return nullptr;
-    }
-
-    uint32_t offsetByte = (uint64_t)ptr - (uint64_t)gpuHeap;
-    uint32_t offsetPage = offsetByte / SIZE_OF_PAGE;
-
-    while(gpuFlags[offsetPage] != 0) {
-        if(gpuFlags[offsetPage] == 2) {
-            gpuFlags[offsetPage] = 0;
-            offsetPage++;
-            break;
-        } else {
-            gpuFlags[offsetPage] = 0;
-            offsetPage++;
-        }
-    }
-
-    return nullptr;
-}
-
-__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
-    return hc::__amdgcn_ds_bpermute(index, src);
-}
-
-__device__ float __hip_ds_bpermutef(int index, float src) {
-    return hc::__amdgcn_ds_bpermute(index, src);
-}
-
-__device__ unsigned __hip_ds_permute(int index, unsigned src) {
-    return hc::__amdgcn_ds_permute(index, src);
-}
-
-__device__ float __hip_ds_permutef(int index, float src) {
-    return hc::__amdgcn_ds_permute(index, src);
-}
-
-__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
-    return hc::__amdgcn_ds_swizzle(src, pattern);
-}
-
-__device__ float __hip_ds_swizzlef(float src, int pattern) {
-    return hc::__amdgcn_ds_swizzle(src, pattern);
-}
-
-__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
-    return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
-}
 //=================================================================================================
 // Thread-local storage:
 //=================================================================================================
@@ -322,7 +214,7 @@ ShortTid::ShortTid()  :
 { 
    _shortTid = g_lastShortTid.fetch_add(1); 

-    if (HIP_DB & (1<<DB_API)) {
+    if (COMPILE_HIP_DB && HIP_TRACE_API) {
        std::stringstream tid_ss;
        std::stringstream tid_ss_num;
        tid_ss_num << std::this_thread::get_id();
@@ -882,18 +774,13 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)

    // Get Max Threads Per Multiprocessor

-    HsaSystemProperties props;
-    hsaKmtReleaseSystemProperties();
-    if(HSAKMT_STATUS_SUCCESS == hsaKmtAcquireSystemProperties(&props)) {
-        HsaNodeProperties node_prop = {0};
-        if(HSAKMT_STATUS_SUCCESS == hsaKmtGetNodeProperties(node, &node_prop)) {
-            uint32_t waves_per_cu = node_prop.MaxWavesPerSIMD;
-            uint32_t simd_per_cu = node_prop.NumSIMDPerCU;
-            prop-> maxThreadsPerMultiProcessor = prop->warpSize*waves_per_cu*simd_per_cu;
-        }
+    HsaNodeProperties node_prop = {0};
+    if(HSAKMT_STATUS_SUCCESS == hsaKmtGetNodeProperties(node, &node_prop)) {
+        uint32_t waves_per_cu = node_prop.MaxWavesPerSIMD;
+        uint32_t simd_per_cu = node_prop.NumSIMDPerCU;
+        prop-> maxThreadsPerMultiProcessor = prop->warpSize*waves_per_cu*simd_per_cu;
    }

-
    // Get memory properties
    err = hsa_agent_iterate_regions(_hsaAgent, get_region_info, prop);
    DeviceErrorCheck(err);
@@ -1368,7 +1255,7 @@ void ihipInit()
    READ_ENV_I(release, HIP_TRACE_API, 0,  "Trace each HIP API call.  Print function name and return code to stderr as program executes.");
    READ_ENV_S(release, HIP_TRACE_API_COLOR, 0,  "Color to use for HIP_API.  None/Red/Green/Yellow/Blue/Magenta/Cyan/White");
    READ_ENV_I(release, HIP_PROFILE_API, 0,  "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, 0x2=full API name including args.");
-    READ_ENV_S(release, HIP_DB_START_API, 0,  "Comma-separted list of tid.api_seq_num for when to start debug and profiling.");
+    READ_ENV_S(release, HIP_DB_START_API, 0,  "Comma-separated list of tid.api_seq_num for when to start debug and profiling.");
    READ_ENV_S(release, HIP_DB_STOP_API, 0,  "Comma-separated list of tid.api_seq_num for when to stop debug and profiling.");
    
    READ_ENV_C(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the sequence are visible to HIP applications and they are enumerated in the order of sequence.", HIP_VISIBLE_DEVICES_callback );
@@ -1387,6 +1274,9 @@ void ihipInit()
    if (HIP_TRACE_API && !COMPILE_HIP_TRACE_API) {
        fprintf (stderr, "warning: env var HIP_TRACE_API=0x%x but COMPILE_HIP_TRACE_API=0.  (perhaps enable COMPILE_HIP_TRACE_API in src code before compiling?)\n", HIP_DB);
    }
+    if (HIP_TRACE_API) {
+        HIP_DB|=0x1;
+    }

    if (HIP_PROFILE_API && !COMPILE_HIP_ATP_MARKER) {
        fprintf (stderr, "warning: env var HIP_PROFILE_API=0x%x but COMPILE_HIP_ATP_MARKER=0.  (perhaps enable COMPILE_HIP_ATP_MARKER in src code before compiling?)\n", HIP_PROFILE_API);
@@ -1859,13 +1749,17 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,

        if (HIP_FORCE_P2P_HOST & 0x1) {
            *forceUnpinnedCopy = true;
-            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", (*copyDevice)->getDeviceNum());
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", 
+                    (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
+
        } else {
-            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d) can see src and dst.\n",  (*copyDevice)->getDeviceNum());
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst.\n",  
+                    (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
        }
    } else {
        *forceUnpinnedCopy = true;
-        tprintf (DB_COPY, "P2P: copy engine(dev:%d) cannot see both host and device pointers - forcing copy with unpinned engine.\n", (*copyDevice)->getDeviceNum());
+        tprintf (DB_COPY, "P2P: copy engine(dev:%d agent=0x%lx) cannot see both host and device pointers - forcing copy with unpinned engine.\n",
+                    (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
    }
 }

@@ -1894,11 +1788,18 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,

    {
        LockedAccessor_StreamCrit_t crit (_criticalData);
-        tprintf (DB_COPY, "copySync copyDev:%d  dst=%p(home_dev:%d, tracked:%d, isDevMem:%d)  src=%p(home_dev:%d, tracked:%d, isDevMem:%d)   sz=%zu dir=%s forceUnpinnedCopy=%d\n",
+        tprintf (DB_COPY, "copySync copyDev:%d  dst=%p (phys_dev:%d, isDevMem:%d)  src=%p(phys_dev:%d, isDevMem:%d)   sz=%zu dir=%s forceUnpinnedCopy=%d\n",
                 copyDevice ? copyDevice->getDeviceNum():-1, 
-                 dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, 
-                 src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem,  
+                 dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
+                 src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
                 sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
+        tprintf (DB_COPY, "  dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
+                 dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes,
+                 dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem);
+        tprintf (DB_COPY, "  src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
+                 src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes,
+                 srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem);
+

 #if USE_COPY_EXT_V2
        crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? &copyDevice->getDevice()->_acc : nullptr, forceUnpinnedCopy);
@@ -1944,13 +1845,17 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
        ihipCtx_t *copyDevice;
        bool forceUnpinnedCopy;
        resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &copyDevice, &forceUnpinnedCopy);
-
-
-        tprintf (DB_COPY, "copyASync copyEngine_dev:%d  dst=%p(home_dev:%d, tracked:%d, isDevMem:%d)  src=%p(home_dev:%d, tracked:%d, isDevMem:%d)   sz=%zu dir=%s.  forceUnpinnedCopy=%d \n",
-                 copyDevice->getDeviceNum(), 
-                 dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, 
-                 src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem,  
+        tprintf (DB_COPY, "copyASync copyDev:%d  dst=%p (phys_dev:%d, isDevMem:%d)  src=%p(phys_dev:%d, isDevMem:%d)   sz=%zu dir=%s forceUnpinnedCopy=%d\n",
+                 copyDevice ? copyDevice->getDeviceNum():-1, 
+                 dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
+                 src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
                 sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
+        tprintf (DB_COPY, "  dst=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
+                 dst, dstPtrInfo._hostPointer, dstPtrInfo._devicePointer, dstPtrInfo._sizeBytes,
+                 dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem);
+        tprintf (DB_COPY, "  src=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d\n",
+                 src, srcPtrInfo._hostPointer, srcPtrInfo._devicePointer, srcPtrInfo._sizeBytes,
+                 srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem);

        // "tracked" really indicates if the pointer's virtual address is available in the GPU address space.
        // If both pointers are not tracked, we need to fall back to a sync copy.
@@ -27,6 +27,7 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 #include "hip_hcc.h"
 #include "trace_helper.h"
+#include "hip/hcc_detail/hip_texture.h"

 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
@@ -1036,7 +1037,3 @@ hipError_t hipFreeArray(hipArray* array)
    return ihipLogStatus(hipStatus);
 }

-
-__device__ void  __threadfence_system(void){
-    // no-op
-}
@@ -7,15 +7,15 @@ int main()
    int pi;
    int attr = 0;
 //    hipDeviceAttribute_t attr = hipDeviceAttributeMaxThreadsPerBlock;
-    HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
+    HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
    HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
    attr = -1;
-    HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
+    HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
    HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), 0), hipDeviceGetAttribute);
    attr = 0;
-    HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
+    HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
    HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
    attr = -1;
-    HIP_CHECK(hipDeviceGetAttribute(nullptr, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
+    HIP_CHECK(hipDeviceGetAttribute(NULL, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
    HIP_CHECK(hipDeviceGetAttribute(&pi, hipDeviceAttribute_t(attr), -1), hipDeviceGetAttribute);
 }
@@ -39,10 +39,10 @@ __device__ void double_precision_intrinsics()
    //__dmul_rn(1.0, 2.0);
    //__dmul_ru(1.0, 2.0);
    //__dmul_rz(1.0, 2.0);
-    //__drcp_rd(2.0);
-    //__drcp_rn(2.0);
-    //__drcp_ru(2.0);
-    //__drcp_rz(2.0);
+    __drcp_rd(2.0);
+    __drcp_rn(2.0);
+    __drcp_ru(2.0);
+    __drcp_rz(2.0);
    __dsqrt_rd(4.0);
    __dsqrt_rn(4.0);
    __dsqrt_ru(4.0);
@@ -51,10 +51,10 @@ __device__ void double_precision_intrinsics()
    //__dsub_rn(2.0, 1.0);
    //__dsub_ru(2.0, 1.0);
    //__dsub_rz(2.0, 1.0);
-    //__fma_rd(1.0, 2.0, 3.0);
-    //__fma_rn(1.0, 2.0, 3.0);
-    //__fma_ru(1.0, 2.0, 3.0);
-    //__fma_rz(1.0, 2.0, 3.0);
+    __fma_rd(1.0, 2.0, 3.0);
+    __fma_rn(1.0, 2.0, 3.0);
+    __fma_ru(1.0, 2.0, 3.0);
+    __fma_rz(1.0, 2.0, 3.0);
 }

 __global__ void compileDoublePrecisionIntrinsics(hipLaunchParm lp, int ignored)
@@ -0,0 +1,62 @@
+/*
+Copyright (c) 2015-2016 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+#define LEN 512
+#define SIZE LEN<<2
+
+
+
+__global__ void floatMath(hipLaunchParm lp, float *In, float *Out) {
+  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  Out[tid] = __cosf(In[tid]);
+  Out[tid] = __exp10f(Out[tid]);
+  Out[tid] = __expf(Out[tid]);
+  Out[tid] = __frsqrt_rn(Out[tid]);
+  Out[tid] = __fsqrt_rd(Out[tid]);
+  Out[tid] = __fsqrt_rn(Out[tid]);
+  Out[tid] = __fsqrt_ru(Out[tid]);
+  Out[tid] = __fsqrt_rz(Out[tid]);
+  Out[tid] = __log10f(Out[tid]);
+  Out[tid] = __log2f(Out[tid]);
+  Out[tid] = __logf(Out[tid]);
+  Out[tid] = __powf(2.0f, Out[tid]);
+  __sincosf(Out[tid], &In[tid], &Out[tid]);
+  Out[tid] = __sinf(Out[tid]);
+  Out[tid] = __cosf(Out[tid]);
+  Out[tid] = __tanf(Out[tid]);
+}
+
+int main(){
+  float *Inh, *Outh, *Ind, *Outd;
+  hipMalloc((void**)&Ind, SIZE);
+  hipMalloc((void**)&Outd, SIZE);
+  hipLaunchKernel(floatMath, dim3(LEN,1,1), dim3(1,1,1), 0, 0, Ind, Outd);
+  passed();
+}
@@ -0,0 +1,122 @@
+/*
+Copyright (c) 2015-2016 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 "test_common.h"
+
+__global__ void FloatMathPrecise(hipLaunchParm lp)
+{
+    int iX;
+    float fX, fY;
+
+    acosf(1.0f);
+    acoshf(1.0f);
+    asinf(0.0f);
+    asinhf(0.0f);
+    atan2f(0.0f, 1.0f);
+    atanf(0.0f);
+    atanhf(0.0f);
+    cbrtf(0.0f);
+    fX = ceilf(0.0f);
+    fX = copysignf(1.0f, -2.0f);
+    cosf(0.0f);
+    coshf(0.0f);
+    cospif(0.0f);
+    //cyl_bessel_i0f(0.0f);
+    //cyl_bessel_i1f(0.0f);
+    erfcf(0.0f);
+    erfcinvf(2.0f);
+    erfcxf(0.0f);
+    erff(0.0f);
+    erfinvf(1.0f);
+    exp10f(0.0f);
+    exp2f(0.0f);
+    expf(0.0f);
+    expm1f(0.0f);
+    fX = fabsf(1.0f);
+    fdimf(1.0f, 0.0f);
+    fdividef(0.0f, 1.0f);
+    fX = floorf(0.0f);
+    fmaf(1.0f, 2.0f, 3.0f);
+    fX = fmaxf(0.0f, 0.0f);
+    fX = fminf(0.0f, 0.0f);
+    fmodf(0.0f, 1.0f);
+    //frexpf(0.0f, &iX);
+    hypotf(1.0f, 0.0f);
+    ilogbf(1.0f);
+    isfinite(0.0f);
+    fX = isinf(0.0f);
+    fX = isnan(0.0f);
+    j0f(0.0f);
+    j1f(0.0f);
+    jnf(-1.0f, 1.0f);
+    ldexpf(0.0f, 0);
+    //lgammaf(1.0f);
+    llrintf(0.0f);
+    llroundf(0.0f);
+    log10f(1.0f);
+    log1pf(-1.0f);
+    log2f(1.0f);
+    logbf(1.0f);
+    logf(1.0f);
+    lrintf(0.0f);
+    lroundf(0.0f);
+    //modff(0.0f, &fX);
+    fX = nanf("1");
+    fX = nearbyintf(0.0f);
+    //nextafterf(0.0f);
+    norm3df(1.0f, 0.0f, 0.0f);
+    norm4df(1.0f, 0.0f, 0.0f, 0.0f);
+    normcdff(0.0f);
+    normcdfinvf(1.0f);
+    fX = 1.0f; normf(1, &fX);
+    powf(1.0f, 0.0f);
+    rcbrtf(1.0f);
+    remainderf(2.0f, 1.0f);
+    //remquof(1.0f, 2.0f, &iX);
+    rhypotf(0.0f, 1.0f);
+    fY = rintf(1.0f);
+    rnorm3df(0.0f, 0.0f, 1.0f);
+    rnorm4df(0.0f, 0.0f, 0.0f, 1.0f);
+    fX = 1.0f; rnormf(1, &fX);
+    fY = roundf(0.0f);
+    rsqrtf(1.0f);
+    scalblnf(0.0f, 1);
+    scalbnf(0.0f, 1);
+    signbit(1.0f);
+    sincosf(0.0f, &fX, &fY);
+    sincospif(0.0f, &fX, &fY);
+    sinf(0.0f);
+    sinhf(0.0f);
+    sinpif(0.0f);
+    sqrtf(0.0f);
+    tanf(0.0f);
+    tanhf(0.0f);
+    tgammaf(2.0f);
+    fY = truncf(0.0f);
+    y0f(1.0f);
+    y1f(1.0f);
+    ynf(1, 1.0f);
+}
+
+int main() {
+    hipLaunchKernel(FloatMathPrecise, dim3(1,1,1), dim3(1,1,1), 0, 0);
+}
@@ -0,0 +1,81 @@
+/*
+Copyright (c) 2015-2016 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s 
+ * RUN: %t
+ * HIT_END
+ */
+
+#include<hip/hip_runtime_api.h>
+#include<hip/hip_runtime.h>
+#include<iostream>
+#include"test_common.h"
+
+#define LEN 512
+#define SIZE LEN<<2
+
+__global__ void kernel_trig(hipLaunchParm lp, float *In, float *sin_d, float *cos_d, float *tan_d, float *sin_pd, float *cos_pd){
+  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  sin_d[tid] = __sinf(In[tid]);
+  cos_d[tid] = __cosf(In[tid]);
+  tan_d[tid] = __tanf(In[tid]);
+  __sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]);
+}
+
+int main(){
+  float *In, *sin_h, *cos_h, *tan_h, *sin_ph, *cos_ph;
+  float *In_d, *sin_d, *cos_d, *tan_d, *sin_pd, *cos_pd;
+  In = new float[LEN];
+  sin_h = new float[LEN];
+  cos_h = new float[LEN];
+  tan_h = new float[LEN];
+  sin_ph = new float[LEN];
+  cos_ph = new float[LEN];
+  for(int i=0;i<LEN;i++){
+    In[i] = 1.0f;
+    sin_h[i] = 0.0f;
+    cos_h[i] = 0.0f;
+    tan_h[i] = 0.0f;
+    sin_ph[i] = 0.0f;
+    cos_ph[i] = 0.0f;
+  }
+  hipMalloc((void**)&In_d, SIZE);
+  hipMalloc((void**)&sin_d, SIZE);
+  hipMalloc((void**)&cos_d, SIZE);
+  hipMalloc((void**)&tan_d, SIZE);
+  hipMalloc((void**)&sin_pd, SIZE);
+  hipMalloc((void**)&cos_pd, SIZE);
+  hipMemcpy(In_d, In, SIZE, hipMemcpyHostToDevice);
+  hipLaunchKernel(kernel_trig, dim3(LEN,1,1), dim3(1,1,1), 0, 0, In_d, sin_d, cos_d, tan_d, sin_pd, cos_pd);
+  hipMemcpy(sin_h, sin_d, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(cos_h, cos_d, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(tan_h, tan_d, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(sin_ph, sin_pd, SIZE, hipMemcpyDeviceToHost);
+  hipMemcpy(cos_ph, cos_pd, SIZE, hipMemcpyDeviceToHost);
+  for(int i=0;i<LEN;i++) {
+    if(sin_h[i] != sin_ph[i] || cos_h[i] != cos_ph[i] || tan_h[i]*cos_h[i] != sin_h[i]){
+      std::cout<<"Failed!"<<std::endl;
+    }
+  }
+  passed();
+}
@@ -0,0 +1,30 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include"test_common.h"
+
+int main()
+{
+    int dev;
+    hipDeviceProp_t prop;
+    HIP_PRINT_STATUS(hipChooseDevice(&dev, &prop));
+    HIP_PRINT_STATUS(hipChooseDevice(0, &prop));
+    HIP_PRINT_STATUS(hipChooseDevice(0, 0));
+}
@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include"test_common.h"
+
+int main()
+{
+    int val;
+    hipDeviceAttribute_t attr = hipDeviceAttributeMaxThreadsPerBlock;                   ///< Maximum number of threads per block.
+
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0));
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0));
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, -1));
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, -1));
+    attr = hipDeviceAttribute_t(91);
+
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, 0));
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, 0));
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(NULL, attr, -1));
+    HIP_PRINT_STATUS(hipDeviceGetAttribute(&val, attr, -1));
+
+}
@@ -0,0 +1,28 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include"test_common.h"
+
+int main()
+{
+    hipFuncCache_t pCacheConfig;
+    HIP_PRINT_STATUS(hipDeviceGetCacheConfig(&pCacheConfig));
+    HIP_PRINT_STATUS(hipDeviceGetCacheConfig(NULL));
+}
@@ -0,0 +1,27 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include "test_common.h"
+
+int main()
+{
+    hipLimit_t lim = hipLimitMallocHeapSize;
+    HIP_PRINT_STATUS(hipDeviceGetLimit(NULL, lim));
+}
@@ -0,0 +1,28 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include"test_common.h"
+
+int main()
+{
+    hipSharedMemConfig_t config;
+    HIP_PRINT_STATUS(hipDeviceGetSharedMemConfig(NULL));
+    HIP_PRINT_STATUS(hipDeviceGetSharedMemConfig(&config));
+}
@@ -0,0 +1,27 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include "test_common.h"
+
+int main()
+{
+    HIP_PRINT_STATUS(hipGetDevice(NULL));
+    HIP_PRINT_STATUS(hipGetDevice(0));
+}
@@ -0,0 +1,26 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include "test_common.h"
+
+int main()
+{
+    HIP_PRINT_STATUS(hipGetDeviceCount(NULL));
+}
@@ -0,0 +1,31 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include "test_common.h"
+
+int main()
+{
+    hipDeviceProp_t prop;
+    HIP_PRINT_STATUS(hipGetDeviceProperties(&prop, -1));
+    int cnt;
+    hipGetDeviceCount(&cnt);
+    HIP_PRINT_STATUS(hipGetDeviceProperties(&prop, cnt+1));
+    HIP_PRINT_STATUS(hipGetDeviceProperties(NULL, 0));
+}
@@ -0,0 +1,29 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include "test_common.h"
+
+int main()
+{
+    HIP_PRINT_STATUS(hipSetDevice(-1));
+    int count;
+    hipGetDeviceCount(&count);
+    HIP_PRINT_STATUS(hipSetDevice(count+1));
+}
@@ -0,0 +1,27 @@
+/*
+Copyright (c) 2015-2016 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_api.h>
+#include "test_common.h"
+
+int main()
+{
+    HIP_PRINT_STATUS(hipSetDeviceFlags(-1));
+    HIP_PRINT_STATUS(hipSetDeviceFlags(11));
+}
@@ -23,6 +23,8 @@ THE SOFTWARE.
 #include <stddef.h>

 #include "hip/hip_runtime.h"
+#include "hip/hip_texture.h"
+#include "hip/hip_runtime_api.h"

 #define HC __attribute__((hc))

@@ -63,6 +65,8 @@ THE SOFTWARE.
    printf ("\n");\
    printf ("warn: TEST WARNING\n%s", KNRM );\

+#define HIP_PRINT_STATUS(status) \
+    std::cout<<hipGetErrorName(status)<<" at line: "<<__LINE__<<std::endl;

 #define HIPCHECK(error) \
 {\