diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index ccd390fbe5..1b95145b96 100644 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -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) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index dee0894869..09c4d813d0 100755 --- a/hipamd/bin/hipcc +++ b/hipamd/bin/hipcc @@ -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 "; } } diff --git a/hipamd/bin/hipify b/hipamd/bin/hipify index 5c56465040..4d77fad3ed 100755 --- a/hipamd/bin/hipify +++ b/hipamd/bin/hipify @@ -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/; $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime_api\.h[>"]/$1/; + $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_fp16\.h[>"]/$1/; #-------- @@ -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 <>> 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; diff --git a/hipamd/bin/hsa-api-filter-cxl.txt b/hipamd/bin/hsa-api-filter-cxl.txt new file mode 100644 index 0000000000..ea3847a791 --- /dev/null +++ b/hipamd/bin/hsa-api-filter-cxl.txt @@ -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 diff --git a/hipamd/docs/markdown/hip_profiling.md b/hipamd/docs/markdown/hip_profiling.md index e4b88945e5..db5d0fc425 100644 --- a/hipamd/docs/markdown/hip_profiling.md +++ b/hipamd/docs/markdown/hip_profiling.md @@ -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 -A + * [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 -A -T +``` + +#### 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 +``` + + +#### 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 -A +$ /opt/rocm/bin/rocm-profiler -A -T Or -/opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -o -A +$ /opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -A -T ``` -#### 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_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: + +``` +<> ``` -#### 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. +- `<> +info: running on device gfx803 +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +<> +<> +info: copy Host2Device +<> +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 +<> +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*' + +... + +<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 +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 +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. diff --git a/hipamd/include/hip/hcc.h b/hipamd/include/hip/hcc.h index 1542d5b4f2..dba26aeab3 100644 --- a/hipamd/include/hip/hcc.h +++ b/hipamd/include/hip/hcc.h @@ -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" diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index b1edef18d7..b1877ed0b3 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.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 +//#include #include #endif #include @@ -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 -#include +//#include //---- //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; diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 5c2e48026e..932e271527 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -32,7 +32,7 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif -#if __HCC__ +#if 0 #include 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;); diff --git a/hipamd/include/hip/hip_texture.h b/hipamd/include/hip/hip_texture.h new file mode 100644 index 0000000000..3e7802b457 --- /dev/null +++ b/hipamd/include/hip/hip_texture.h @@ -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 +#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) +#include +#else +#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif + + +#endif diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 0d15dfcb01..625448094b 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -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)); } diff --git a/hipamd/include/hip/nvcc_detail/hip_texture.h b/hipamd/include/hip/nvcc_detail/hip_texture.h new file mode 100644 index 0000000000..388733e492 --- /dev/null +++ b/hipamd/include/hip/nvcc_detail/hip_texture.h @@ -0,0 +1,6 @@ +#ifndef HIP_TEXTURE_H +#define HIP_TEXTURE_H + +#include + +#endif diff --git a/hipamd/packaging/hip_hcc.txt b/hipamd/packaging/hip_hcc.txt index d7b0877b62..78717247f0 100644 --- a/hipamd/packaging/hip_hcc.txt +++ b/hipamd/packaging/hip_hcc.txt @@ -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) diff --git a/hipamd/samples/1_Utils/hipCommander/LICENSE.txt b/hipamd/samples/1_Utils/hipCommander/LICENSE.txt new file mode 100644 index 0000000000..5d0d603232 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/LICENSE.txt @@ -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. + diff --git a/hipamd/samples/1_Utils/hipCommander/Makefile b/hipamd/samples/1_Utils/hipCommander/Makefile new file mode 100644 index 0000000000..e770c636a4 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/Makefile @@ -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) diff --git a/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp new file mode 100644 index 0000000000..2ec686f260 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.cpp @@ -0,0 +1,527 @@ +#include "ResultDatabase.h" + +#include +#include +#include +#include + +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= 100) + return value[n-1]; + + double index = ((n + 1.) * q / 100.) - 1; + + vector 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 &values) +{ + for (int i=0; i= 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 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 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 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(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 +ResultDatabase::GetResultsForTest(const string &test) +{ + // get only the given test results + vector retval; + for (int i=0; i & +ResultDatabase::GetResults() const +{ + return results; +} diff --git a/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h new file mode 100644 index 0000000000..4b63a02a1f --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/ResultDatabase.h @@ -0,0 +1,100 @@ +#ifndef RESULT_DATABASE_H +#define RESULT_DATABASE_H + +#include +#include +#include +#include +#include +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 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= FLT_MAX) + return true; + } + return false; + } + }; + + protected: + vector 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 &values); + vector GetResultsForTest(const string &test); + const vector &GetResults() const; + void ClearAllResults(); + void DumpDetailed(ostream&); + void DumpSummary(ostream&); + void DumpCsv(string fileName); + + private: + bool IsFileEmpty(string fileName); + +}; + + +#endif diff --git a/hipamd/samples/1_Utils/hipCommander/TODO b/hipamd/samples/1_Utils/hipCommander/TODO new file mode 100644 index 0000000000..4c835cfced --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/TODO @@ -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: + + diff --git a/hipamd/samples/1_Utils/hipCommander/c.cmd b/hipamd/samples/1_Utils/hipCommander/c.cmd new file mode 100644 index 0000000000..db11071203 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/c.cmd @@ -0,0 +1,3 @@ +loop,1000; H2D; NullKernel; D2H; endloop; +streamsync; +printTiming, 1000 diff --git a/hipamd/samples/1_Utils/hipCommander/classic.cmd b/hipamd/samples/1_Utils/hipCommander/classic.cmd new file mode 100644 index 0000000000..c149eec5f7 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/classic.cmd @@ -0,0 +1 @@ +H2D; NullKernel, D2H, streamsync diff --git a/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp b/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp new file mode 100644 index 0000000000..9c07d066ba --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/hipCommander.cpp @@ -0,0 +1,1096 @@ +#include +#include +#include +#include +#include +#include +#include + +#include +#ifdef __HIP_PLATFORM_HCC__ +#include +#include +#include +#include +#endif + +#include + +#include "ResultDatabase.h" +#include "nullkernel.hip.cpp" + +bool g_printedTiming = false; + +// Cmdline parms: +int p_device = 0; +const char* p_command = "H2D; NullKernel; D2H"; +const char* p_file = nullptr; +unsigned p_verbose = 0x0; +unsigned p_db = 0x0; +unsigned p_blockingSync = 0x0; + +//--- +int p_iterations = 1; + +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" +#define KGRN "\x1B[32m" + + +#define failed(...) \ + printf ("error: ");\ + printf (__VA_ARGS__);\ + printf ("\n");\ + abort(); + + +#define HIPCHECK(error) \ +{\ + hipError_t localError = error; \ + if (localError != hipSuccess) { \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \ + KRED,hipGetErrorString(localError), localError,\ + #error,\ + __FILE__, __LINE__,KNRM); \ + failed("API returned error code.");\ + }\ +} +#define HIPASSERT(condition, msg) \ + if (! (condition) ) { \ + failed("%sassertion %s at %s:%d: %s%s\n", \ + KRED, #condition,\ + __FILE__, __LINE__,msg, KNRM); \ + } + + + + + + + +int parseInt(const char *str, int *output) +{ + char *next; + *output = strtol(str, &next, 0); + return !strlen(next); +} + + +void printConfig() { + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, p_device)); + + printf ("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz\n", props.name, props.totalGlobalMem/1024.0/1024.0/1024.0, props.multiProcessorCount, props.clockRate/1000.0); +} + + + + +void help() { + printf ("Usage: hipBusBandwidth [OPTIONS]\n"); + printf (" --file, -f : Read string of commands from file\n"); + printf (" --command, -c : String specifying commands to run.\n"); + printf (" --iterations, -i : Number of copy iterations to run.\n"); + printf (" --device, -d : Device ID to use (0..numDevices).\n"); + printf (" --verbose, -v : Verbose printing of status. Fore more info, combine with HIP_TRACE_API on ROCm\n"); +}; + + + +int parseStandardArguments(int argc, char *argv[]) +{ + for (int i = 1; i < argc; i++) { + const char *arg = argv[i]; + + if (!strcmp(arg, " ")) { + // skip NULL args. + } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { + if (++i >= argc || !parseInt(argv[i], &p_iterations)) { + failed("Bad --iterations argument"); + } + + } else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) { + if (++i >= argc || !parseInt(argv[i], &p_device)) { + failed("Bad --device argument"); + } + + } else if (!strcmp(arg, "--file") || (!strcmp(arg, "-f"))) { + if (++i >= argc) { + failed("Bad --file argument"); + } else { + p_file = argv[i]; + } + + } else if (!strcmp(arg, "--commands") || (!strcmp(arg, "-c"))) { + if (++i >= argc) { + failed("Bad --commands argument"); + } else { + p_command = argv[i]; + } + + } else if (!strcmp(arg, "--verbose") || (!strcmp(arg, "-v"))) { + p_verbose = 1; + + } else if (!strcmp(arg, "--blockingSync") || (!strcmp(arg, "-B"))) { + p_blockingSync = 1; + + + } else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) { + help(); + exit(EXIT_SUCCESS); + + } else { + failed("Bad argument '%s'", arg); + } + } + + return 0; +}; + +// Returns the current system time in microseconds +inline long long get_time() +{ + struct timeval tv; + gettimeofday(&tv, 0); + return (tv.tv_sec * 1000000) + tv.tv_usec; +} + + +class Command; + + +//================================================================================================= +// A stream of commands , specified as a string. +class CommandStream { +public: + // State that is inherited by sub-blocks: + struct CommandStreamState { + hipStream_t _currentStream; + std::vector _streams; + vector _subBlocks; + }; +public: + CommandStream(std::string commandStreamString, int iterations); + ~CommandStream(); + + hipStream_t currentStream() const { return _state._currentStream; }; + + void print(const std::string &indent="") const; + void printBrief(std::ostream &s=std::cout) const ; + void run(); + void recordTime(); + void printTiming(int iterations=0); + + CommandStream *currentCommandStream() { + return _parseInSubBlock ? _state._subBlocks.back() : this; + }; + + void enterSubBlock(CommandStream *commandStream) { + _parseInSubBlock = true; + _state._subBlocks.push_back(commandStream); + }; + + void exitSubBlock() { + _parseInSubBlock = false; + }; + + + void setParent(CommandStream *parentCmdStream) + { + _parentCommandStream = parentCmdStream; + _state = parentCmdStream->_state; + }; + CommandStream * getParent() { return _parentCommandStream; }; + + void setStream(int streamIndex); + + CommandStreamState &getState() { return _state; }; + +private: + static void tokenize(const std::string &s, char delim, std::vector &tokens); + void parse(const std::string fullCmd); + +protected: + CommandStreamState _state; +private: + + + // List of commands to run in this stream: + std::vector _commands; + + + + // Number of iterations to run the command loop + int _iterations; + + + + + // Us to run the the command-stream. Only valid after run is called. + long long _startTime; + double _elapsedUs; + + // Track nested loop of command streams: + CommandStream *_parentCommandStream; + + // Track if we are parsing commands in the subblock. + bool _parseInSubBlock; + +}; + + +//================================================================================================= +class Command { +public: + + // @p minArgs : Minimum arguments for command. -1 = don't check. + // @p maxArgs : Minimum arguments for command. 0 means min=max, ie exact #arguments expected. -1 = don't check max. + Command(CommandStream *cmdStream, const std::vector &args, int minArgs=0, int maxArgs=0) + : _commandStream(cmdStream), + _args(args) + { + int numArgs = args.size() - 1; + + if ((minArgs != -1 ) && (numArgs < minArgs)) { + // TODO - print full command here. + failed ("Not enough arguments for command %s. (Expected %d, got %d)", args[0].c_str(), minArgs, numArgs); + } + + // Check for an exact number of arguments: + if (maxArgs == 0) { + maxArgs = minArgs; + } + if ((maxArgs != -1 ) && (numArgs > maxArgs)) { + failed ("Too many arguments for command %s. (Expected %d, got %d)", args[0].c_str(), maxArgs, numArgs); + } + }; + + void printBrief(std::ostream &s=std::cout) const + { + s << _args[0]; + } + + virtual ~Command() {}; + + virtual void print(const std::string &indent = "") const { + std::cout << indent << "["; + std::for_each(_args.begin(), _args.end(), [] (const std::string &s) { + std::cout << s; + }); + std::cout << "]"; + }; + + + virtual void run() = 0; + +protected: + int readIntArg(int argIndex, const std::string &argName) + { + // TODO - catch references to non-existant arguments here. + int argVal; + try { + argVal = std::stoi(_args[argIndex]); + } catch (std::invalid_argument) { + failed ("Command %s has bad %s argument ('%s')", _args[0].c_str(), argName.c_str(), _args[argIndex].c_str()); + } + return argVal; + } +protected: + CommandStream *_commandStream; + std::vector _args; +}; + + +#define FILENAME "nullkernel.hsaco" +#define KERNEL_NAME "NullKernel" + + +#ifdef __HIP_PLATFORM_HCC__ +//================================================================================================= +// Use Aql to launch the NULL kernel. +class AqlKernelCommand : public Command +{ +public: + AqlKernelCommand(CommandStream *cmdStream, const std::vector args) : + Command(cmdStream, args) + { + hc::accelerator_view *av; + HIPCHECK(hipHccGetAcceleratorView(cmdStream->currentStream(), &av)); + + hc::accelerator acc = av->get_accelerator(); + + hsa_region_t systemRegion = *(hsa_region_t*)acc.get_hsa_am_system_region(); + + _hsaAgent = *(hsa_agent_t*) acc.get_hsa_agent(); + + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) + { + uint64_t elfSize = ElfSize(&buffer[0]); + + assert(fsize == elfSize); + + //TODO - replace module load code with explicit module load and unload. + + hipModule_t module; + HIPCHECK(hipModuleLoadData(&module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); + + } else { + failed("could not open code object '%s'\n", FILENAME); + } + }; + + ~AqlKernelCommand() {}; + + void run() override { +#define LEN 64 + uint32_t len = LEN; + uint32_t one = 1; + + float *Ad = NULL; + + size_t argSize = 36; + char argBuffer[argSize]; + *(uint32_t*) (&argBuffer[0]) = len; + *(uint32_t*) (&argBuffer[4]) = one; + *(uint32_t*) (&argBuffer[8]) = one; + *(uint32_t*) (&argBuffer[12]) = len; + *(uint32_t*) (&argBuffer[16]) = one; + *(uint32_t*) (&argBuffer[20]) = one; + *(float**) (&argBuffer[24]) = Ad; // Ad pointer argument + + + void *config[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, + HIP_LAUNCH_PARAM_END + }; + + hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + }; + + +public: + hsa_queue_t _hsaQueue; + hsa_agent_t _hsaAgent; + + hipFunction_t _function; + +private: + static uint64_t ElfSize(const void *emi){ + const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; + const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); + + uint64_t max_offset = ehdr->e_shoff; + uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; + + for(uint16_t i=0;i < ehdr->e_shnum;++i){ + uint64_t cur_offset = static_cast(shdr[i].sh_offset); + if(max_offset < cur_offset){ + max_offset = cur_offset; + total_size = max_offset; + if(SHT_NOBITS != shdr[i].sh_type){ + total_size += static_cast(shdr[i].sh_size); + } + } + } + return total_size; + } +}; +#endif + +//================================================================================================= +// HCC optimizes away fully NULL kernel calls, so run one that is nearly null: +class ModuleKernelCommand : public Command +{ +public: + ModuleKernelCommand(CommandStream *cmdStream, const std::vector args) : + Command(cmdStream, args), + _stream (cmdStream->currentStream()) { + + hipModule_t module; + HIPCHECK(hipModuleLoad(&module, FILENAME)); + HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME)); + }; + ~ModuleKernelCommand() {}; + + void run() override { +#define LEN 64 + uint32_t len = LEN; + uint32_t one = 1; + + float *Ad = NULL; + + size_t argSize = 36; + char argBuffer[argSize]; + *(uint32_t*) (&argBuffer[0]) = len; + *(uint32_t*) (&argBuffer[4]) = one; + *(uint32_t*) (&argBuffer[8]) = one; + *(uint32_t*) (&argBuffer[12]) = len; + *(uint32_t*) (&argBuffer[16]) = one; + *(uint32_t*) (&argBuffer[20]) = one; + *(float**) (&argBuffer[24]) = Ad; // Ad pointer argument + + + void *config[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], + HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, + HIP_LAUNCH_PARAM_END + }; + + hipModuleLaunchKernel(_function, len, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); + }; + + +public: + hipFunction_t _function; + hipStream_t _stream; +}; + + +class KernelCommand : public Command +{ +public: + enum Type {Null, VectorAdd}; + KernelCommand(CommandStream *cmdStream, const std::vector args, Type kind) : + Command(cmdStream, args), + _kind(kind), + _stream (cmdStream->currentStream()) { + }; + ~KernelCommand() {}; + + + void run() override { + static const int gridX = 64; + static const int groupX = 64; + + switch (_kind) { + case Null: + hipLaunchKernel(NullKernel, dim3(gridX/groupX), dim3(gridX), 0, _stream, nullptr); + break; + case VectorAdd: + assert(0); // TODO + break; + }; + } +private: + Type _kind; + hipStream_t _stream; +}; + + +#ifdef __HIP_PLATFORM_HCC__ +//================================================================================================= +class PfeCommand : public Command +{ +public: + PfeCommand(CommandStream *cmdStream, const std::vector args, hipStream_t stream = 0) : + Command(cmdStream, args) + { + HIPCHECK(hipHccGetAcceleratorView(stream, &_av)); + } + + ~PfeCommand() { } + + + void run() override { + static const int gridX = 64; + static const int groupX = 64; + auto cf = hc::parallel_for_each(*_av, hc::extent<1>(gridX).tile(groupX), + [=](hc::index<1>& idx) __HC__ { + }); + } +private: + hc::accelerator_view *_av; +}; +#endif + + + +//================================================================================================= +class CopyCommand : public Command +{ +enum MemType {PinnedHost, UnpinnedHost, Device} ; + +public: + CopyCommand(CommandStream *cmdStream, const std::vector &args, hipMemcpyKind kind, bool isAsync, bool isPinnedHost) ; + + ~CopyCommand() + { + if (_dst) { + dealloc(_dst, _dstType); + _dst = NULL; + }; + + if (_src) { + dealloc(_src, _srcType); + _src = NULL; + } + } + + + void run() override { + if (_isAsync) { + HIPCHECK(hipMemcpyAsync(_dst, _src, _sizeBytes, _kind, _stream)); + } else { + HIPCHECK(hipMemcpy(_dst, _src, _sizeBytes, _kind)); + } + }; + +private: + void * alloc(size_t size, MemType memType) { + void * p; + if (memType == Device) { + HIPCHECK(hipMalloc(&p, size)); + + } else if (memType == PinnedHost) { + HIPCHECK(hipHostMalloc(&p, size)); + + } else if (memType == UnpinnedHost) { + p = (char*)malloc(size); + HIPASSERT(p, "malloc failed"); + + } else { + HIPASSERT(0, "unsupported memType"); + } + + return p; + }; + + + void dealloc(void *p, MemType memType) { + if (memType == Device) { + HIPCHECK(hipFree(p)); + } else if (memType == PinnedHost) { + HIPCHECK(hipHostFree(p)); + } else if (memType == UnpinnedHost) { + free(p); + } else { + HIPASSERT(0, "unsupported memType"); + } + } + + +private: + bool _isAsync; + hipStream_t _stream; + hipMemcpyKind _kind; + + size_t _sizeBytes; + void *_dst; + MemType _dstType; + + void *_src; + MemType _srcType; +}; + + +//================================================================================================= +class DeviceSyncCommand : public Command +{ +public: + DeviceSyncCommand(CommandStream *cmdStream, const std::vector &args) : + Command(cmdStream, args) {}; + + void run() override { + HIPCHECK(hipDeviceSynchronize()); + }; +}; + + +//================================================================================================= +class StreamSyncCommand : public Command +{ +public: + StreamSyncCommand(CommandStream *cmdStream, const std::vector &args) : + Command(cmdStream, args), + _stream(cmdStream->currentStream()) + {}; + + const char *help() { + return "synchronizes the current stream"; + }; + + + + void run() override { + HIPCHECK(hipStreamSynchronize(_stream)); + }; + +private: + hipStream_t _stream; +}; + + +//================================================================================================= + +//================================================================================================= +class LoopCommand : public Command +{ +public: + LoopCommand(CommandStream *parentCmdStream, const std::vector &args) : + Command(parentCmdStream, args, 1) + { + int loopCnt; + try { + loopCnt = std::stoi(args[1]); + } catch (std::invalid_argument) { + failed ("bad LOOP_CNT=%s", args[1].c_str()); + } + + _commandStream = new CommandStream("", loopCnt); + _commandStream->setParent(parentCmdStream); + parentCmdStream->enterSubBlock(_commandStream); + + }; + + + void print(const std::string &indent = "") const override { + Command::print(); + _commandStream->print (indent + " "); + }; + + void run() override { + _commandStream->run(); + }; +}; + + +//================================================================================================= +class EndBlockCommand : public Command +{ +public: + EndBlockCommand(CommandStream *blockCmdStream, CommandStream *parentCmdStream, const std::vector &args) : + Command(parentCmdStream, args, 0, 1), + _blockCmdStream(blockCmdStream), + _printTiming(0) + { + int argCnt = args.size()-1; + if (argCnt >= 1 ) { + _printTiming = readIntArg(1, "PRINT_TIMING"); + } + + if (parentCmdStream == nullptr) { + failed ("%s without corresponding command to start block", args[0].c_str()); + } + parentCmdStream->exitSubBlock(); + }; + + void run() override { + if (_printTiming) { + _blockCmdStream->printTiming(); + } + + }; +private: + + CommandStream *_blockCmdStream; + + // print the stream when loop exits. + int _printTiming; +}; + + +//================================================================================================= +class SetStreamCommand : public Command +{ +public: + SetStreamCommand(CommandStream *cmdStream, const std::vector &args) : + Command(cmdStream, args, 1) + { + int streamIndex = readIntArg(1, "STREAM_INDEX"); + + cmdStream->setStream(streamIndex); + + }; + + void run() override { + }; +}; + + +//================================================================================================= +class PrintTimingCommand : public Command +{ +public: + PrintTimingCommand(CommandStream *cmdStream, const std::vector &args) + : Command(cmdStream, args, 1) + { + _iterations = readIntArg(1, "ITERATIONS"); + }; + + void run() override { + _commandStream->printTiming(_iterations); + }; + +private: + int _iterations; +}; + + +//================================================================================================= +CopyCommand::CopyCommand(CommandStream *cmdStream, const std::vector &args, + hipMemcpyKind kind, bool isAsync, bool isPinnedHost) : + Command(cmdStream, args) , + _isAsync(isAsync), + _kind(kind), + _stream(cmdStream->currentStream()) + { + switch (kind) { + case hipMemcpyDeviceToHost: + _srcType = Device; + _dstType = isPinnedHost ? PinnedHost : UnpinnedHost; + break; + case hipMemcpyHostToDevice: + _srcType = isPinnedHost ? PinnedHost : UnpinnedHost; + _dstType = Device; + break; + default: + HIPASSERT(0, "Unknown hipMemcpyKind"); + }; + + _sizeBytes = 64; //TODO, support reading from arg. + + _dst = alloc(_sizeBytes, _dstType); + _src = alloc(_sizeBytes, _srcType); + }; + + +//================================================================================================= +//================================================================================================= +// Implementations: +//================================================================================================= + +//================================================================================================= +CommandStream::CommandStream(std::string commandStreamString, int iterations) + : _iterations(iterations), + _startTime(0), + _elapsedUs(0.0), + _parentCommandStream(nullptr), + _parseInSubBlock(false) +{ + std::vector tokens; + tokenize(commandStreamString, ';', tokens); + + + std::for_each(tokens.begin(), tokens.end(), [&] (const std::string s) { + this->parse(s); + }); + + setStream(0); +} + + +CommandStream::~CommandStream() +{ + std::for_each(_state._streams.begin(), _state._streams.end(), [&] (hipStream_t s) { + if (s) { + HIPCHECK(hipStreamDestroy(s)); + } + }); + + std::for_each(_commands.begin(), _commands.end(), [&] (Command *c) { + delete c; + }); + + +} + + +void CommandStream::setStream(int streamIndex) +{ + + if (streamIndex >= _state._streams.size()) { + _state._streams.resize(streamIndex+1); + } + + if (streamIndex && (_state._streams[streamIndex] == nullptr)) { + // Create new stream: + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + _state._streams[streamIndex] = stream; + _state._currentStream = stream; + } else { + // Use existing stream: + + _state._currentStream = _state._streams[streamIndex]; + } + +} + + +void CommandStream::tokenize(const std::string &s, char delim, std::vector &tokens) +{ + std::stringstream ss; + ss.str(s); + std::string item; + while (getline(ss, item, delim)) { + item.erase (std::remove (item.begin(), item.end(), ' '), item.end()); // remove whitespace. + tokens.push_back(item); + } +} + +void trim(std::string *s) +{ + // trim whitespace from begin and end: + const char *t = "\t\n\r\f\v"; + s->erase(0, s->find_first_not_of(t)); + s->erase(s->find_last_not_of(t)+1); +} + +void ltrim(std::string *s) +{ + // trim whitespace from begin and end: + const char *t = "\t\n\r\f\v"; + s->erase(0, s->find_first_not_of(t)); +} + +void CommandStream::parse(std::string fullCmd) +{ + //convert to lower-case: + std::transform(fullCmd.begin(), fullCmd.end(), fullCmd.begin(), ::tolower); + trim(&fullCmd); + + if (p_db) { + printf ("parse: <%s>\n", fullCmd.c_str()); + } + + + + std::string c; + std::vector args; + size_t leftParenZ = fullCmd.find_first_of('('); + if (leftParenZ == string::npos) { + c = fullCmd; + args.push_back(c); + } else { + c = fullCmd.substr(0, leftParenZ); + args.push_back(c); + size_t rightParenZ = fullCmd.find_first_of(')', leftParenZ); + std::string argStr = fullCmd.substr(leftParenZ+1, rightParenZ-leftParenZ-1); + //printf ("c=%s argstr='%s' leftParenZ=%zu rightParenZ=%zu\n", c.c_str(), argStr.c_str(), leftParenZ, rightParenZ); + tokenize(argStr, ',', args); + + } + + + + + + if ((args.size()==0) || (fullCmd.c_str()[0] == '#') ) { + if (p_db) { + printf (" skip comment\n"); + } + return; + } + + + + Command *cmd = NULL; + CommandStream *cmdStream = currentCommandStream(); + + if (c == "h2d") { + cmd = new CopyCommand(cmdStream, args, hipMemcpyHostToDevice, true/*isAsync*/, true/*isPinned*/); + //= h2d + //= Performs an async host-to-device copy of array A_h to A_d. + //= The size of these arrays may be set with the datasize command. + + } else if (c == "d2h") { + cmd = new CopyCommand(cmdStream, args, hipMemcpyDeviceToHost, true/*isAsync*/, true/*isPinned*/); + //= d2h + //= Performs an async device-to-host copy of array A_d to A_h. + //= The size of these arrays may be set with the datasize command. + + } else if (c == "modulekernel") { + cmd = new ModuleKernelCommand(cmdStream, args); + + } else if (c == "nullkernel") { + cmd = new KernelCommand(cmdStream, args, KernelCommand::Null); + //= nullkernel + //= Dispatches a null kernel to the device. + + } else if (c == "vectoraddkernel") { + cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd); + +#ifdef __HIP_PLATFORM_HCC__ + } else if (c == "nullpfe") { + cmd = new PfeCommand(cmdStream, args); + + } else if (c == "aqlkernel") { + cmd = new AqlKernelCommand(cmdStream, args); +#endif + + } else if (c == "devicesync") { + cmd = new DeviceSyncCommand(cmdStream, args); + + } else if (c == "streamsync") { + //= streamsync + //= Execute hipStreamSynchronize. + //= This will cause the host thread to wait until the current stream + //= completes all pending operations. + cmd = new StreamSyncCommand(cmdStream, args); + + } else if (c == "setstream") { + //= setstream(STREAM_INDEX); + //= Set current stream used by subsequent commands. + //= STREAM_INDEX is index starting from 0...N. + //= This function will create new stream on first call to setstream or re-use previous + //= stream if setstream has already been called with STREAM_INDEX. + //= STREAM_INDEX=0 will use the default "null" stream associated with the device, and will not create a new stream. + //= The default stream has special, conservative synchronization properties. + + cmd = new SetStreamCommand(cmdStream, args); + + } else if (c == "printtiming") { + cmd = new PrintTimingCommand(cmdStream, args); + + } else if (c == "loop") { + //= loop(LOOP_CNT) + //= Loop over next set of commands (until 'endloop' command) for LOOP_CNT iterations. + //= Loops can be nested. + + cmd = new LoopCommand(cmdStream, args); + + } else if (c == "endloop") { + //= endloop + //= End a looped sequence. Must be paired with a preceding loop command. + //= Command between the `loop` and `endloop` must be executed + + CommandStream * parentCmdStream = cmdStream->getParent() ; + cmd = new EndBlockCommand(cmdStream, parentCmdStream, args); + cmdStream = parentCmdStream; + + } else { + std::cerr << "error: Bad command '" << fullCmd << "\n"; + HIPASSERT(0, "bad command in command-stream"); + } + + if (cmd) { + cmdStream->_commands.push_back(cmd); + } +} + + + + +void CommandStream::print(const std::string &indent) const +{ + for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { + (*cmdI)->print(indent); + }; +} + + +void CommandStream::printBrief(std::ostream &s) const +{ + for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { + (*cmdI)->printBrief(s); + s << ";"; + }; +} + +void CommandStream::run() +{ + _startTime = get_time(); + for (int i=0; i<_iterations; i++) { + for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) { + if (p_verbose) { + (*cmdI)->print(); + } + (*cmdI)->run(); + } + } + + // Record time, if not already stored. (an earlier printTime command will also store the time) + recordTime(); +}; + +void CommandStream::recordTime() +{ + if (_elapsedUs == 0.0) { + auto stopTime = get_time(); + _elapsedUs = stopTime - _startTime; + } +} + + +void CommandStream::printTiming(int iterations) +{ + + if ((_state._subBlocks.size() == 1) && (_commands.size()==1)) { + //printf ("print just the loop\n"); + _state._subBlocks.front()->printTiming(iterations); + } else { + g_printedTiming = true; + + recordTime(); + if (iterations == 0) { + iterations = _iterations; + } + std::cout << "command<"; printBrief(std::cout); + std::cout << ">," ; + printf (" iterations,%d, total_time,%6.3f, time/iteration,%6.3f\n", iterations, _elapsedUs, _elapsedUs/iterations); + } +}; + + + + + +//================================================================================================= +int main(int argc, char *argv[]) +{ + parseStandardArguments(argc, argv); + + printConfig(); + + CommandStream *cs; + + if (p_blockingSync) { +#ifdef __HIP_PLATFORM_HCC__ + printf ("setting BlockingSync for AMD\n"); + setenv("HIP_BLOCKING_SYNC", "1", 1); + +#endif +#ifdef __HIP_PLATFORM_NVCC__ + printf ("setting cudaDeviceBlockingSync\n"); + HIPCHECK(hipSetDeviceFlags(cudaDeviceBlockingSync)); +#endif + }; + + + if (p_file) { + // TODO - catch exception on file IO here: + std::ifstream file(p_file); + std::string str; + std::string file_contents; + while (std::getline(file, str)) + { + file_contents += str; + } + + cs = new CommandStream(file_contents, p_iterations); + + } else { + cs = new CommandStream(p_command, p_iterations); + } + + cs->print(); + printf ("------\n"); + + cs->run(); + if (!g_printedTiming) { + cs->printTiming(); + } + + delete cs; +} + + + +// TODO - add error checking for arguments. diff --git a/hipamd/samples/1_Utils/hipCommander/l2.hcm b/hipamd/samples/1_Utils/hipCommander/l2.hcm new file mode 100644 index 0000000000..b541bd6a66 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/l2.hcm @@ -0,0 +1,3 @@ +setstream,1; +NullKernel; streamsync; +loop,10000; H2D; NullKernel; streamsync; endloop,1; diff --git a/hipamd/samples/1_Utils/hipCommander/loop.hcm b/hipamd/samples/1_Utils/hipCommander/loop.hcm new file mode 100644 index 0000000000..db11071203 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/loop.hcm @@ -0,0 +1,3 @@ +loop,1000; H2D; NullKernel; D2H; endloop; +streamsync; +printTiming, 1000 diff --git a/hipamd/samples/1_Utils/hipCommander/loop2.hcm b/hipamd/samples/1_Utils/hipCommander/loop2.hcm new file mode 100644 index 0000000000..b8a14aa156 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/loop2.hcm @@ -0,0 +1,2 @@ +setstream,1; +loop,1000; NullKernel; syncstream; endloop,1, diff --git a/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp b/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp new file mode 100644 index 0000000000..890e9bdc1e --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/nullkernel.hip.cpp @@ -0,0 +1,7 @@ +#include "hip/hip_runtime.h" + +extern "C" __global__ void NullKernel(hipLaunchParm lp, float* Ad){ + if (Ad) { + Ad[0] = 42; + } +} diff --git a/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco b/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco new file mode 100755 index 0000000000..585b55cce5 Binary files /dev/null and b/hipamd/samples/1_Utils/hipCommander/nullkernel.hsaco differ diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm new file mode 100644 index 0000000000..e43960dc5a --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/perf/latency2.hcm @@ -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); diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm new file mode 100644 index 0000000000..511ab355d5 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/perf/latency_hostsync.hcm @@ -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; + diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm new file mode 100644 index 0000000000..c89d738be9 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/perf/latency_nosync.hcm @@ -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; diff --git a/hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm b/hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm new file mode 100644 index 0000000000..69345b23b0 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/perf/latency_nullstream.hcm @@ -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; diff --git a/hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm b/hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm new file mode 100644 index 0000000000..d1d4091fad --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm @@ -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); diff --git a/hipamd/samples/1_Utils/hipCommander/setstream.hcm b/hipamd/samples/1_Utils/hipCommander/setstream.hcm new file mode 100644 index 0000000000..a7bdd093b8 --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/setstream.hcm @@ -0,0 +1,3 @@ +setstream,1; +setstream,2; H2D; NullKernel; D2H; +streamsync diff --git a/hipamd/samples/1_Utils/hipCommander/testcase.cpp b/hipamd/samples/1_Utils/hipCommander/testcase.cpp new file mode 100644 index 0000000000..0eeae7c50f --- /dev/null +++ b/hipamd/samples/1_Utils/hipCommander/testcase.cpp @@ -0,0 +1,21 @@ +#include + +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); +} diff --git a/hipamd/samples/2_Cookbook/8_peer2peer/Makefile b/hipamd/samples/2_Cookbook/8_peer2peer/Makefile index a1dad7d1da..5cb7473921 100644 --- a/hipamd/samples/2_Cookbook/8_peer2peer/Makefile +++ b/hipamd/samples/2_Cookbook/8_peer2peer/Makefile @@ -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 + diff --git a/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp b/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp index 624de56cb0..990599e1cb 100644 --- a/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp +++ b/hipamd/samples/2_Cookbook/8_peer2peer/peer2peer.cpp @@ -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, diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index b267c8cbd7..203e7a7826 100644 --- a/hipamd/src/device_util.cpp +++ b/hipamd/src/device_util.cpp @@ -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;igetDevice()->_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++; diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index d066079fc0..36790271de 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -109,114 +109,6 @@ std::atomic g_lastShortTid(1); std::vector g_dbStartTriggers; std::vector 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 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 ? ©Device->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, ©Device, &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. diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 672b9f2ee2..089d1a626e 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -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 -} diff --git a/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp b/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp index b901f660ca..fbe26814e7 100644 --- a/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp +++ b/hipamd/tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp @@ -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); } diff --git a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 0dab2d7106..330c4fc799 100644 --- a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -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) diff --git a/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/hipamd/tests/src/deviceLib/hipFloatMath.cpp new file mode 100644 index 0000000000..f137ca2602 --- /dev/null +++ b/hipamd/tests/src/deviceLib/hipFloatMath.cpp @@ -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(); +} diff --git a/hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp b/hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp new file mode 100644 index 0000000000..4f6c2cd44a --- /dev/null +++ b/hipamd/tests/src/deviceLib/hipFloatMathPrecise.cpp @@ -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); +} diff --git a/hipamd/tests/src/deviceLib/hip_trig.cpp b/hipamd/tests/src/deviceLib/hip_trig.cpp new file mode 100644 index 0000000000..7f9b5d60b0 --- /dev/null +++ b/hipamd/tests/src/deviceLib/hip_trig.cpp @@ -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 +#include +#include +#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 +#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)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp new file mode 100644 index 0000000000..a7e1814468 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetAttribute.cpp @@ -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 +#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)); + +} diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp new file mode 100644 index 0000000000..17f7ed7fbd --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetCacheConfig.cpp @@ -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 +#include"test_common.h" + +int main() +{ + hipFuncCache_t pCacheConfig; + HIP_PRINT_STATUS(hipDeviceGetCacheConfig(&pCacheConfig)); + HIP_PRINT_STATUS(hipDeviceGetCacheConfig(NULL)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp new file mode 100644 index 0000000000..57b8594527 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetLimit.cpp @@ -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 +#include "test_common.h" + +int main() +{ + hipLimit_t lim = hipLimitMallocHeapSize; + HIP_PRINT_STATUS(hipDeviceGetLimit(NULL, lim)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp b/hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp new file mode 100644 index 0000000000..ce447aa828 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipDeviceGetSharedMemConfig.cpp @@ -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 +#include"test_common.h" + +int main() +{ + hipSharedMemConfig_t config; + HIP_PRINT_STATUS(hipDeviceGetSharedMemConfig(NULL)); + HIP_PRINT_STATUS(hipDeviceGetSharedMemConfig(&config)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp b/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp new file mode 100644 index 0000000000..0e9e5d96f2 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipGetDevice.cpp @@ -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 +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipGetDevice(NULL)); + HIP_PRINT_STATUS(hipGetDevice(0)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp b/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp new file mode 100644 index 0000000000..6565741b2f --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipGetDeviceCount.cpp @@ -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 +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipGetDeviceCount(NULL)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp b/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp new file mode 100644 index 0000000000..325d9dd410 --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipGetDeviceProperties.cpp @@ -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 +#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)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp b/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp new file mode 100644 index 0000000000..99c62a03be --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipSetDevice.cpp @@ -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 +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipSetDevice(-1)); + int count; + hipGetDeviceCount(&count); + HIP_PRINT_STATUS(hipSetDevice(count+1)); +} diff --git a/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp b/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp new file mode 100644 index 0000000000..e933ab56fa --- /dev/null +++ b/hipamd/tests/src/nvcc/Device/hipSetDeviceFlags.cpp @@ -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 +#include "test_common.h" + +int main() +{ + HIP_PRINT_STATUS(hipSetDeviceFlags(-1)); + HIP_PRINT_STATUS(hipSetDeviceFlags(11)); +} diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h index 2b30b9d48f..1250de4801 100644 --- a/hipamd/tests/src/test_common.h +++ b/hipamd/tests/src/test_common.h @@ -23,6 +23,8 @@ THE SOFTWARE. #include #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<