diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 1ba58496f4..94ed2a7562 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -63,6 +63,7 @@ if(HIP_PLATFORM STREQUAL "hcc") set(HCC_HOME $ENV{HCC_HOME} CACHE PATH "Path to which HCC has been installed") endif() endif() + if(DEFINED ENV{HIP_DEVELOPER}) add_to_config(_buildInfo HCC_HOME) endif() @@ -178,6 +179,7 @@ if(HIP_PLATFORM STREQUAL "hcc") src/hip_peer.cpp src/hip_stream.cpp src/hip_module.cpp + src/grid_launch.cpp src/env.cpp) set(SOURCE_FILES_DEVICE @@ -189,9 +191,11 @@ if(HIP_PLATFORM STREQUAL "hcc") execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") + #find_package(LLVM HINTS ${HCC_HOME}/compiler/lib/cmake) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803 --amdgpu-target=gfx900") add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME}) target_link_libraries(hip_hcc PRIVATE hc_am) + #target_link_libraries(hip_hcc PUBLIC LLVMAMDGPUUtils) add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME}) target_link_libraries(hip_hcc_static PRIVATE hc_am) add_dependencies(hip_hcc_static hip_hcc) diff --git a/projects/clr/hipamd/bin/hccgenco.sh b/projects/clr/hipamd/bin/hccgenco.sh index a7b4177624..9d7f602911 100755 --- a/projects/clr/hipamd/bin/hccgenco.sh +++ b/projects/clr/hipamd/bin/hccgenco.sh @@ -10,7 +10,7 @@ if [ $# = 0 ]; then fi : ${ROCM_PATH:=/opt/rocm} -: ${ROCM_TARGET:=fiji} +: ${ROCM_TARGET:=gfx803} INPUT_FILES="" OUTPUT_FILE="" @@ -45,16 +45,17 @@ for inputfile in $INPUT_FILES; do done printf "\nint main(){}\n" >> $hipgenisa_main -$HIP_PATH/bin/hipcc $hipgenisa_files -o $hipgenisa_dir/a.out +$HIP_PATH/bin/hipcc -DGENERIC_GRID_LAUNCH=0 $hipgenisa_files -o $hipgenisa_dir/a.out mv dump* $hipgenisa_dir +hsaco_file="dump-$ROCM_TARGET.hsaco" map_sym="" -kernels=$(objdump -t $hipgenisa_dir/dump-fiji.hsaco | grep grid_launch_parm | sed 's/ \+/ /g; s/\t/ /g' | cut -d" " -f6) +kernels=$(objdump -t $hipgenisa_dir/$hsaco_file | grep grid_launch_parm | sed 's/ \+/ /g; s/\t/ /g' | cut -d" " -f6) for mangled_sym in $kernels; do real_sym=$(c++filt $(c++filt $mangled_sym | cut -d: -f3 | sed 's/_functor//g') | cut -d\( -f1) map_sym="--redefine-sym $mangled_sym=$real_sym $map_sym" done -objcopy -F elf64-little $map_sym $hipgenisa_dir/dump-fiji.hsaco $OUTPUT_FILE +objcopy -F elf64-little $map_sym $hipgenisa_dir/$hsaco_file $OUTPUT_FILE rm $hipgenisa_files rm -r $hipgenisa_dir diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 381c774c94..bcd3e3a591 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -92,7 +92,7 @@ if ($HIP_PLATFORM eq "hcc") { $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'} // 1; $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger"; - $ROCM_TARGET=$ENV{'ROCM_TARGET'} // "fiji"; + $ROCM_TARGET=$ENV{'ROCM_TARGET'} // "gfx803"; # HCC* may be used to compile src/hip_hcc.o (and also feed the HIPCXXFLAGS below) $HCC = "$HCC_HOME/bin/hcc"; @@ -103,6 +103,7 @@ if ($HIP_PLATFORM eq "hcc") { $HIPLDFLAGS = `${HCC_HOME}/bin/hcc-config --ldflags`; + #### GCC system includes workaround #### $HCC_WA_FLAGS = " "; if ($HCC_VERSION_MAJOR eq 1) { @@ -122,12 +123,21 @@ if ($HIP_PLATFORM eq "hcc") { } } + # Force -stdlib=libc++ on UB14.04 + $HOST_OSNAME= `cat /etc/os-release | grep "^ID\=" | cut -d= -f2 | tr -d '\n'`; + $HOST_OSVER= `cat /etc/os-release | grep "^VERSION_ID\=" | cut -d= -f2 | tr -d '\n'`; + if ($HOST_OSNAME eq "ubuntu" and $HOST_OSVER eq "\"14.04\"") { + $HIPCXXFLAGS .= " -stdlib=libc++"; + $setStdLib = 1; + } + $HIPCXXFLAGS .= " -I$HIP_PATH/include/hip/hcc_detail/cuda"; $HIPCXXFLAGS .= " -I$HSA_PATH/include"; $HIPCXXFLAGS .= " -Wno-deprecated-register"; $HIPLDFLAGS .= " -lsupc++"; - $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt"; + $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt "; +# $HIPLDFLAGS .= " -L$HCC_HOME/compiler/lib -lLLVMAMDGPUDesc -lLLVMAMDGPUUtils -lLLVMMC -lLLVMCore -lLLVMSupport "; # Add trace marker library: # TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library. @@ -402,9 +412,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc') if ($needHipHcc) { if ($linkType eq 0) { - substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ; + substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ; } else { - substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a "; + substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a "; } } diff --git a/projects/clr/hipamd/bin/hipconvertinplace-perl.sh b/projects/clr/hipamd/bin/hipconvertinplace-perl.sh new file mode 100755 index 0000000000..a8c8d6d9e8 --- /dev/null +++ b/projects/clr/hipamd/bin/hipconvertinplace-perl.sh @@ -0,0 +1,18 @@ +#!/bin/bash + +#usage : hipconvertinplace.sh [DIRNAME] [HIPIFY_OPTIONS] + +#hipify "inplace" all code files in specified directory. +# This can be quite handy when dealing with an existing CUDA code base since the script +# preserves the existing directory structure. + +# For each code file, this script will: +# - If ".prehip file does not exist, copy the original code to a new file with extension ".prehip". Then Hipify the code file. +# - If ".prehip" file exists, this is used as input to hipify. +# (this is useful for testing improvements to the hipify toolset). + + +SCRIPT_DIR=`dirname $0` +SEARCH_DIR=$1 +shift +$SCRIPT_DIR/hipify -inplace -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` diff --git a/projects/clr/hipamd/bin/hipconvertinplace.sh b/projects/clr/hipamd/bin/hipconvertinplace.sh index a8c8d6d9e8..a765ab39fa 100755 --- a/projects/clr/hipamd/bin/hipconvertinplace.sh +++ b/projects/clr/hipamd/bin/hipconvertinplace.sh @@ -1,18 +1,24 @@ #!/bin/bash -#usage : hipconvertinplace.sh [DIRNAME] [HIPIFY_OPTIONS] +#usage : hipconvertinplace.sh DIRNAME [hipify options] [--] [clang options] -#hipify "inplace" all code files in specified directory. +#hipify "inplace" all code files in specified directory. # This can be quite handy when dealing with an existing CUDA code base since the script # preserves the existing directory structure. -# For each code file, this script will: -# - If ".prehip file does not exist, copy the original code to a new file with extension ".prehip". Then Hipify the code file. -# - If ".prehip" file exists, this is used as input to hipify. -# (this is useful for testing improvements to the hipify toolset). - - SCRIPT_DIR=`dirname $0` SEARCH_DIR=$1 -shift -$SCRIPT_DIR/hipify -inplace -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` + +hipify_args='' +while (( "$#" )); do + shift + if [ "$1" != "--" ]; then + hipify_args="$hipify_args $1" + else + shift + break + fi +done +clang_args="$@" + +$SCRIPT_DIR/hipify-clang -inplace -print-stats $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args diff --git a/projects/clr/hipamd/bin/hipconvertinplace2.sh b/projects/clr/hipamd/bin/hipconvertinplace2.sh deleted file mode 100644 index a765ab39fa..0000000000 --- a/projects/clr/hipamd/bin/hipconvertinplace2.sh +++ /dev/null @@ -1,24 +0,0 @@ -#!/bin/bash - -#usage : hipconvertinplace.sh DIRNAME [hipify options] [--] [clang options] - -#hipify "inplace" all code files in specified directory. -# This can be quite handy when dealing with an existing CUDA code base since the script -# preserves the existing directory structure. - -SCRIPT_DIR=`dirname $0` -SEARCH_DIR=$1 - -hipify_args='' -while (( "$#" )); do - shift - if [ "$1" != "--" ]; then - hipify_args="$hipify_args $1" - else - shift - break - fi -done -clang_args="$@" - -$SCRIPT_DIR/hipify-clang -inplace -print-stats $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args diff --git a/projects/clr/hipamd/bin/hipexamine-perl.sh b/projects/clr/hipamd/bin/hipexamine-perl.sh new file mode 100755 index 0000000000..40c1bf466d --- /dev/null +++ b/projects/clr/hipamd/bin/hipexamine-perl.sh @@ -0,0 +1,12 @@ +#!/bin/bash + +#usage : hipexamine.sh DIRNAME [hipify.pl options] + +# Generate HIP stats (LOC, CUDA->API conversions, missing functionality) for all the code files +# in the specified directory. + + +SCRIPT_DIR=`dirname $0` +SEARCH_DIR=$1 +shift +$SCRIPT_DIR/hipify -no-output -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` diff --git a/projects/clr/hipamd/bin/hipexamine.sh b/projects/clr/hipamd/bin/hipexamine.sh index 40c1bf466d..2a6fab7110 100755 --- a/projects/clr/hipamd/bin/hipexamine.sh +++ b/projects/clr/hipamd/bin/hipexamine.sh @@ -1,12 +1,22 @@ #!/bin/bash -#usage : hipexamine.sh DIRNAME [hipify.pl options] - -# Generate HIP stats (LOC, CUDA->API conversions, missing functionality) for all the code files -# in the specified directory. +#usage : hipexamine2.sh DIRNAME [hipify options] [--] [clang options] +# Generate CUDA->HIP conversion statistics for all the code files in the specified directory. SCRIPT_DIR=`dirname $0` SEARCH_DIR=$1 -shift -$SCRIPT_DIR/hipify -no-output -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` + +hipify_args='' +while (( "$#" )); do + shift + if [ "$1" != "--" ]; then + hipify_args="$hipify_args $1" + else + shift + break + fi +done +clang_args="$@" + +$SCRIPT_DIR/hipify-clang -examine $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args diff --git a/projects/clr/hipamd/bin/hipexamine2.sh b/projects/clr/hipamd/bin/hipexamine2.sh deleted file mode 100644 index 2a6fab7110..0000000000 --- a/projects/clr/hipamd/bin/hipexamine2.sh +++ /dev/null @@ -1,22 +0,0 @@ -#!/bin/bash - -#usage : hipexamine2.sh DIRNAME [hipify options] [--] [clang options] - -# Generate CUDA->HIP conversion statistics for all the code files in the specified directory. - -SCRIPT_DIR=`dirname $0` -SEARCH_DIR=$1 - -hipify_args='' -while (( "$#" )); do - shift - if [ "$1" != "--" ]; then - hipify_args="$hipify_args $1" - else - shift - break - fi -done -clang_args="$@" - -$SCRIPT_DIR/hipify-clang -examine $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args diff --git a/projects/clr/hipamd/docs/markdown/hip_bugs.md b/projects/clr/hipamd/docs/markdown/hip_bugs.md index 3274ed5b5c..e15c37fc54 100644 --- a/projects/clr/hipamd/docs/markdown/hip_bugs.md +++ b/projects/clr/hipamd/docs/markdown/hip_bugs.md @@ -4,6 +4,7 @@ - [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**](#errors-related-to-undefined-reference-to-hclaunchkernel__grid_launch_parm) - [Application hangs after a hipLaunchKernel call](#what-if-i-see-application-hangs-after-a-hiplaunchkernel-call) +- [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method) @@ -23,8 +24,6 @@ namespace { __global__ MyKernel … - Avoid calling member functions - -### What if I see application hangs after a hipLaunchKernel call? If hipLaunchKernel takes parameters that request explicitly memcpy, then it will cause application hang. Reason is that the hipLaunchKernel macro locks the stream. If kernel paramters are actually function calls which invoke other hip apis (i.e. memcpy) to the same stream, then deadlock occurs. @@ -43,4 +42,8 @@ auto bot_gpu_data = bottom[0]->gpu_data(); hipLaunchKernel( LRNComputeDiff, dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads, bot_gpu_data); -``` \ No newline at end of file +``` + +### What is the current limitation of HIP Generic Grid Launch method? +1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak))); +2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch. \ No newline at end of file diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index d7235c4ebc..8ccb458103 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -26,6 +26,7 @@ - [How do I trace HIP application flow?](#how-do-i-trace-hip-application-flow) * [Using CodeXL markers for HIP Functions](#using-codexl-markers-for-hip-functions) * [Using HIP_TRACE_API](#using-hip_trace_api) +- [How do I enable HIP Generic Grid Launch option?](#how-do-i-enable-hip-generic-grid-launch-option) @@ -235,3 +236,9 @@ Unlike CUDA, in HCC, for functions defined in the header files, the keyword of " Thus, if failed to define "static" keyword, you might see a lot of "symbol multiply defined!" errors at compilation. The workaround is to explicitly add the keyword of "static" before any functions that were defined as "__forceinline__". +### How do I disable HIP Generic Grid Launch option? +Generic Grid Launch(GGL) is currently the default method for hip kernel launch. +To disable it and use the legancy grid launch method, please either change the default value of GENERIC_GRID_LAUNCH to 0 in the following to header files and rebuild HIP: +$HIP/include/hip/hcc_detail/hip_runtime_api.h +$HIP/include/hip/hcc_detail/host_defines.h +Or pass "-DGENERIC_GRID_LAUNCH=0" to hipcc at application compilation time. diff --git a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md index 3fd4ea9aca..0c7f3c8d25 100644 --- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md +++ b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md @@ -680,7 +680,7 @@ The user can specify the target for which the binary can be generated. HIP/HCC d The file format for binary is `.co` which means Code Object. The following command builds the code object using `hipcc`. `hipcc --genco --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]` -```[TARGET GPU] = fiji/hawaii +```[TARGET GPU] = gfx803/gfx701 [INPUT FILE] = Name of the file containing kernels [OUTPUT FILE] = Name of the generated code object file``` diff --git a/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md b/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md index f51f53a092..dd3b9c3e86 100644 --- a/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md +++ b/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md @@ -1,151 +1,119 @@ # Porting CUDA Driver API ## Introduction to the CUDA Driver and Runtime APIs -CUDA provides a separate CUDA Driver and Runtime APIs. The two APIs have significant overlap in functionality: +CUDA provides a separate CUDA Driver and Runtime APIs. The two APIs have significant overlap in functionality: - Both APIs support events, streams, memory management, memory copy, and error handling. -- Both APIs deliver similar performance. -- Driver APIs calls begin with the prefix `cu` while Runtime APIs begin with the prefix `cuda`. For example, the Driver API API contains `cuEventCreate` while the Runtime API contains `cudaEventCreate`, with similar functionality. -- The Driver API defines a different but largely overlapping error code code space than the Runtime API, and uses a different coding convention. For example, Driver API defines `CUDA_ERROR_INVALID_VALUE` while the Runtime API defines `cudaErrorInvalidValue` +- Both APIs deliver similar performance. +- Driver APIs calls begin with the prefix `cu` while Runtime APIs begin with the prefix `cuda`. For example, the Driver API API contains `cuEventCreate` while the Runtime API contains `cudaEventCreate`, with similar functionality. +- The Driver API defines a different but largely overlapping error code space than the Runtime API, and uses a different coding convention. For example, Driver API defines `CUDA_ERROR_INVALID_VALUE` while the Runtime API defines `cudaErrorInvalidValue` The Driver API offers two additional pieces of functionality not provided by the Runtime API: cuModule and cuCtx APIs. ### cuModule API -The Module section of the Driver API provides additional control over how and when accelerator - code objects are loaded. +The Module section of the Driver API provides additional control over how and when accelerator code objects are loaded. For example, the driver API allows code objects to be loaded from files or memory pointers. Symbols for kernels or global data can be extracted from the loaded code objects. -In contrast, the Runtime API automatically loads and (if necessary) compiles all of the -kernels from a executable binary when run. +In contrast, the Runtime API automatically loads and (if necessary) compiles all of the kernels from an executable binary when run. In this mode, NVCC must be used to compile kernel code so the automatic loading can function correctly. Both Driver and Runtime APIs define a function for launching kernels (called `cuLaunchKernel` or `cudaLaunchKernel`. The kernel arguments and the execution configuration (grid dimensions, group dimensions, dynamic shared memory, and stream) are passed as arguments to the launch function. -The Runtime additionally provides the `<<< >>>` syntax for launching kernels, which resembles -a special function call and is easier to use than explicit launch API (in particular with respect to -handling of kernel arguments). -However, this syntax is not standard C++ and is available only when NVCC is used to compile the host code. +The Runtime additionally provides the `<<< >>>` syntax for launching kernels, which resembles a special function call and is easier to use than explicit launch API (in particular with respect to handling of kernel arguments). +However, this syntax is not standard C++ and is available only when NVCC is used to compile the host code. -The Module features are useful in an environment which generate the code objects directly, such as a new -accelerator language front-end. Here, NVCC is not used. Instead, the environment may have a -different kernel language or different compilation flow. +The Module features are useful in an environment which generates the code objects directly, such as a new accelerator language front-end. +Here, NVCC is not used. Instead, the environment may have a different kernel language or different compilation flow. Other environments have many kernels and do not want them to be all loaded automatically. The Module functions can be used to load the generated code objects and launch kernels. -As we will see below, HIP defines a Module API which provides similar explicit control over code -object management. +As we will see below, HIP defines a Module API which provides similar explicit control over code object management. ### cuCtx API -The Driver API defines "Context" and "Devices" as separate entities. -Contexts contain a single device, and a device can theoretically have multiple contexts. -Each context contains a set of streams and events specific to the context. -Historically contexts also defined a unique address space for the GPU, though this may not longer be the case in Unified Memory platforms (since the CPU and all the devices in the same process share a single unified address space). -The Context APIs also provide a mechanism to switch between devices, which allowed -a single CPU thread to send commands to different GPUs. HIP as well as a recent versions -of CUDA Runtime provide other mechanisms to accomplish this feat - for example using streams or -`cudaSetDevice`. +The Driver API defines "Context" and "Devices" as separate entities. +Contexts contain a single device, and a device can theoretically have multiple contexts. +Each context contains a set of streams and events specific to the context. +Historically contexts also defined a unique address space for the GPU, though this may no longer be the case in Unified Memory platforms (since the CPU and all the devices in the same process share a single unified address space). +The Context APIs also provide a mechanism to switch between devices, which allowed a single CPU thread to send commands to different GPUs. +HIP as well as a recent versions of CUDA Runtime provide other mechanisms to accomplish this feat - for example using streams or `cudaSetDevice`. -The CUDA Runtime API unifies the Context API with the Device API. This simplifies the APIs -and has little loss of functionality since each Context can contain a single device, -and the benefits of multiple contexts has been replaced with other interfaces. HIP provides -a context API to facilitate easy porting from existing Driver codes. -In HIP, the Ctx functions largely provide an alternate syntax for changing the active device. +The CUDA Runtime API unifies the Context API with the Device API. This simplifies the APIs and has little loss of functionality since each Context can contain a single device, and the benefits of multiple contexts has been replaced with other interfaces. +HIP provides a context API to facilitate easy porting from existing Driver codes. +In HIP, the Ctx functions largely provide an alternate syntax for changing the active device. Most new applications will prefer to use `hipSetDevice` or the stream APIs. ## HIP Module and Ctx APIs -Rather than present two separate APIs, HIP extends the HIP API with new APIs for Modules -and Ctx control. +Rather than present two separate APIs, HIP extends the HIP API with new APIs for Modules and Ctx control. ### hipModule API -Like the CUDA Driver API, the Module API provides additional control -over how code is loaded, including options to load code from files or from in-memory pointers. -NVCC and HCC target different architectures and use different code object formats : NVCC -is `cubin` or `ptx` files, while the HCC path is the `hsaco` format. The external -compilers which generate these code objects are responsible for generating and loading -the correct code object for each platform. Notably, there is not a fat binary format that -can contain code for both NVCC and HCC platforms. The following table summarizes the -formats used on each platform: +Like the CUDA Driver API, the Module API provides additional control over how code is loaded, including options to load code from files or from in-memory pointers. +NVCC and HCC target different architectures and use different code object formats: NVCC is `cubin` or `ptx` files, while the HCC path is the `hsaco` format. +The external compilers which generate these code objects are responsible for generating and loading the correct code object for each platform. +Notably, there is not a fat binary format that can contain code for both NVCC and HCC platforms. The following table summarizes the formats used on each platform: -| Format | APIs | NVCC | HCC | -| --- | --- | --- | --- | +| Format | APIs | NVCC | HCC | +| --- | --- | --- | --- | | Code Object | hipModuleLoad, hipModuleLoadData | .cubin or PTX text | .hsaco | -| Fat Binary | hipModuleLoadFatBin | .fatbin | Under Development | +| Fat Binary | hipModuleLoadFatBin | .fatbin | Under Development | -hipcc uses NVCC and HCC to compile host codes. Both of these may embed code objects -into the final executable, and these code objects will be automatically loaded when -the application starts. -The hipModule API can be used to load additional code objects, and in this way -provides an extended capability to the automatically loaded code objects. HCC allows -both of these capabilities to be used together, if desired. Of course it is possible -to create a program with no kernels and thus no automatic loading. +`hipcc` uses NVCC and HCC to compile host codes. Both of these may embed code objects into the final executable, and these code objects will be automatically loaded when the application starts. +The hipModule API can be used to load additional code objects, and in this way provides an extended capability to the automatically loaded code objects. +HCC allows both of these capabilities to be used together, if desired. Of course it is possible to create a program with no kernels and thus no automatic loading. ### hipCtx API -HIP provides a `Ctx` API as a thin layer over the existing Device functions. This Ctx API -can be used to set the current context, or to query properties of the device associated with -the context. The current context is implicitly used by other APIs such as `hipStreamCreate`. +HIP provides a `Ctx` API as a thin layer over the existing Device functions. This Ctx API can be used to set the current context, or to query properties of the device associated with the context. +The current context is implicitly used by other APIs such as `hipStreamCreate`. ### hipify translation of CUDA Driver API -The hipify tool will convert CUDA Driver APIs for streams, events, memory management to -the equivalent HIP driver calls. For example, `cuEventCreate` will be translated to -`hipEventCreate`. Hipify also converts error code from the Driver namespace and coding -convention to the equivalent HIP error code. Thus, HIP unifies the APIs for these common functions. -[hipify support for translating driver API is Under Development] +The hipify tool converts CUDA Driver APIs for streams, events, modules, devices, memory management, context, profiler to the equivalent HIP driver calls. For example, `cuEventCreate` will be translated to `hipEventCreate`. +Hipify also converts error code from the Driver namespace and coding convention to the equivalent HIP error code. Thus, HIP unifies the APIs for these common functions. -The memory copy APIs require additional explanation. The CUDA driver includes the memory -direction in the name of the API (ie `cuMemcpyH2D`) while the CUDA driver API provides -a single memory copy API with a parameter that specifies the direction and additionally -supports a "default" direction where the runtime determines the direction automatically. +The memory copy API requires additional explanation. The CUDA driver includes the memory direction in the name of the API (ie `cuMemcpyH2D`) while the CUDA driver API provides a single memory copy API with a parameter that specifies the direction and additionally supports a "default" direction where the runtime determines the direction automatically. HIP provides APIs with both styles: for example, `hipMemcpyH2D` as well as `hipMemcpy`. -The first flavor may be faster in some cases since they avoid host overhead to detect the -different memory directions. +The first flavor may be faster in some cases since they avoid host overhead to detect the different memory directions. -HIP defines a single error space, and uses camel-case for all errors (i.e. `hipErrorInvalidValue`). +HIP defines a single error space, and uses camel-case for all errors (i.e. `hipErrorInvalidValue`). ### HCC Implementation Notes #### .hsaco The .hsaco format used by HCC is described in more detail [here](https://github.com/RadeonOpenCompute/ROCm-ComputeABI-Doc). -An example and blog that show how to use the format is [here](http://gpuopen.com/rocm-with-harmony-combining-opencl-hcc-hsa-in-a-single-program). hsaco can be generated by hcc + extractkernel tool, -cloc, the GCN assembler, or other tools. +An example and blog that show how to use the format is [here](http://gpuopen.com/rocm-with-harmony-combining-opencl-hcc-hsa-in-a-single-program). hsaco can be generated by hcc + extractkernel tool, cloc, the GCN assembler, or other tools. #### Address Spaces -HCC defines a process-wide address space where the CPU and all devices allocate -addresses from a single unified pool. Thus addresses may be shared between contexts, and -unlike the original CUDA definition a new context does not create a new address space for -the device. +HCC defines a process-wide address space where the CPU and all devices allocate addresses from a single unified pool. +Thus addresses may be shared between contexts, and unlike the original CUDA definition a new context does not create a new address space for the device. #### Using hipModuleLaunchKernel `hipModuleLaunchKernel` is `cuLaunchKernel` in HIP world. It takes the same arguments as `cuLaunchKernel`. The argument `kernelParams` is not fully implemented for HCC. The workaround for it is, to use platform specific macros for each target. Or, `extra` argument can be used which works on both the platforms. #### Additional Information - HCC allocates staging buffers (used for unpinned copies) on a per-device basis. -- HCC creates a primary context when the HIP API is called. So in a pure driver API code, HIP/HCC will create a primary context while HIP/NVCC will have empty context stack. HIP/HCC will push primary context to context stack when it is empty. This can have subtle differences on applications which mix the runtime and driver APIs. +- HCC creates a primary context when the HIP API is called. So in a pure driver API code, HIP/HCC will create a primary context while HIP/NVCC will have empty context stack. +HIP/HCC will push primary context to context stack when it is empty. This can have subtle differences on applications which mix the runtime and driver APIs. ### NVCC Implementation Notes #### Interoperation between HIP and CUDA Driver -CUDA applications may want to mix CUDA driver code with HIP code (see example below). This -table shows the type equivalence to enable this interaction. - -|**HIP Type** |**CU Driver Type**| **CUDA Runtime Type** | -| ---- | ---- | ---- | -| hipModule | CUmodule | | -| hipFunction | CUfunction | | -| hipCtx_t | CUcontext | | -| hipDevice_t | CUdevice | | -| hipStream_t | CUstream | cudaStream_t | -| hipEvent_t | CUevent | cudaEvent_t | -| hipArray | CUarray | cudaArray | - +CUDA applications may want to mix CUDA driver code with HIP code (see example below). This table shows the type equivalence to enable this interaction. +|**HIP Type** |**CU Driver Type**|**CUDA Runtime Type**| +| ---- | ---- | ---- | +| hipModule | CUmodule | | +| hipFunction | CUfunction | | +| hipCtx_t | CUcontext | | +| hipDevice_t | CUdevice | | +| hipStream_t | CUstream | cudaStream_t | +| hipEvent_t | CUevent | cudaEvent_t | +| hipArray | CUarray | cudaArray | #### Compilation Flags -The hipModule interface does not support the hipModuleLoadEx function, which is used to control PTX compilaton options. +The hipModule interface does not support the `cuModuleLoadDataEx` function, which is used to control PTX compilation options. HCC does not use PTX and does not support the same compilation options. In fact, HCC code objects always contain fully compiled ISA and do not require additional compilation as part of the load step. -Code which requires this functionally should use platform-specific coding, calling `cuModuleLoadEx` -on the NVCC path and hipModuleLoad on the hcc path. For example: +Code which requires this functionally should use platform-specific coding, calling `cuModuleLoadDataEx` on the NVCC path and `hipModuleLoadData` on the hcc path. +For example: ``` hipModule module; @@ -240,7 +208,7 @@ int main(){ HIP_LAUNCH_PARAM_END }; - hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&con fig); + hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); hipMemcpyDtoH(B, Bd, SIZE); for(uint32_t i=0;isecond.hipName; DEBUG(dbgs() << "Identifier " << name << " found as an actual argument in expansion of macro " - << macroName << "\n" + << MacroNameTok.getIdentifierInfo()->getName() << "\n" << "will be replaced with: " << repName << "\n"); size_t length = name.size(); SourceLocation sl = tok.getLocation(); diff --git a/projects/clr/hipamd/include/hip/channel_descriptor.h b/projects/clr/hipamd/include/hip/channel_descriptor.h index af8875e256..b8e750b079 100644 --- a/projects/clr/hipamd/include/hip/channel_descriptor.h +++ b/projects/clr/hipamd/include/hip/channel_descriptor.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H +#define HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H // Some standard header files, these are included by hc.hpp and so want to make them avail on both // paths to provide a consistent include env and avoid "missing symbol" errors that only appears @@ -34,3 +35,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/device_functions.h b/projects/clr/hipamd/include/hip/device_functions.h index 24211b7d2d..aae6775d48 100644 --- a/projects/clr/hipamd/include/hip/device_functions.h +++ b/projects/clr/hipamd/include/hip/device_functions.h @@ -1,13 +1,16 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -17,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_DEVICE_FUNCTIONS_H -#define HIP_DEVICE_FUNCTIONS_H +#ifndef HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H +#define HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H #include diff --git a/projects/clr/hipamd/include/hip/driver_types.h b/projects/clr/hipamd/include/hip/driver_types.h index a4010d6b4e..5d06457dd5 100644 --- a/projects/clr/hipamd/include/hip/driver_types.h +++ b/projects/clr/hipamd/include/hip/driver_types.h @@ -1,22 +1,27 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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, +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. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_DRIVER_TYPES_H +#define HIP_INCLUDE_HIP_DRIVER_TYPES_H #include @@ -27,3 +32,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/hcc_detail/channel_descriptor.h b/projects/clr/hipamd/include/hip/hcc_detail/channel_descriptor.h index 85689438e2..4be023f6ca 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/channel_descriptor.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/channel_descriptor.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H -#define HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H #include #include diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hcc_acc.h b/projects/clr/hipamd/include/hip/hcc_detail/concepts.hpp similarity index 64% rename from projects/clr/hipamd/include/hip/hcc_detail/hcc_acc.h rename to projects/clr/hipamd/include/hip/hcc_detail/concepts.hpp index c36acc52f5..5c50f5d577 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hcc_acc.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/concepts.hpp @@ -20,27 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HCC_ACC_H -#define HCC_ACC_H -#include "hip/hip_runtime_api.h" +#pragma once -#if __cplusplus -#ifdef __HCC__ -#include -/** - * @brief Return hc::accelerator associated with the specified deviceId - * @return #hipSuccess, #hipErrorInvalidDevice - */ -hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc); +namespace hip_impl // Documentation only. +{ + #define requires(...) -/** - * @brief Return hc::accelerator_view associated with the specified stream - * - * If stream is 0, the accelerator_view for the default stream is returned. - * @return #hipSuccess - */ -hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av); -#endif -#endif - -#endif + #define FunctionalProcedure typename +} diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index 212f22d5dc..8073503364 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -1,13 +1,16 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -17,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H -#define HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H #include #include diff --git a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h index 1fe72b0507..3578ddc609 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h @@ -1,13 +1,16 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -17,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_DETAIL_DRIVER_TYPES_H -#define HIP_HCC_DETAIL_DRIVER_TYPES_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H enum hipChannelFormatKind { diff --git a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_GGL.hpp b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_GGL.hpp new file mode 100644 index 0000000000..4fd7c3ff3a --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_GGL.hpp @@ -0,0 +1,853 @@ +/* +Copyright (c) 2015 - present 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. +*/ + +#pragma once + +#include "concepts.hpp" +#include "helpers.hpp" + +#include "hc.hpp" +#include "hip_hcc.h" +#include "hip_runtime.h" + +#include +#include +#include +#include +#include + +namespace hip_impl +{ + namespace + { + struct New_grid_launch_tag {}; + struct Old_grid_launch_tag {}; + + template + class RAII_guard { + D dtor_; + public: + RAII_guard() = default; + + RAII_guard(const C& ctor, D dtor) : dtor_{std::move(dtor)} + { + ctor(); + } + + RAII_guard(const RAII_guard&) = default; + RAII_guard(RAII_guard&&) = default; + + RAII_guard& operator=(const RAII_guard&) = default; + RAII_guard& operator=(RAII_guard&&) = default; + + ~RAII_guard() { dtor_(); } + }; + + template + RAII_guard make_RAII_guard(const C& ctor, D dtor) + { + return RAII_guard{ctor, std::move(dtor)}; + } + + template + using is_new_grid_launch_t = typename std::conditional< + std::is_callable{}, + New_grid_launch_tag, + Old_grid_launch_tag>::type; + } + + // TODO: - dispatch rank should be derived from the domain dimensions passed + // in, and not always assumed to be 3; + + template + requires(Domain == {Ts...}) + inline + void grid_launch_hip_impl_( + New_grid_launch_tag, + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + const hc::accelerator_view& acc_v, + K k, + Ts&&... args) + { + const auto d = hc::extent<3>{ + num_blocks.z * dim_blocks.z, + num_blocks.y * dim_blocks.y, + num_blocks.x * dim_blocks.x}.tile_with_dynamic( + dim_blocks.z, + dim_blocks.y, + dim_blocks.x, + group_mem_bytes); + + try { + hc::parallel_for_each( + acc_v, + d, + [=](const hc::tiled_index<3>& idx) [[hc]] { + k(args...); + }); + } + catch (std::exception& ex) { + std::cerr << "Failed in " << __FUNCTION__ << ", with exception: " + << ex.what() << std::endl; + throw; + } + } + + // TODO: these are workarounds, they should be removed. + + hc::accelerator_view lock_stream_hip_(hipStream_t&, void*&); + void unlock_stream_hip_( + hipStream_t, void*, const char*, hc::accelerator_view*); + + template + requires(Domain == {Ts...}) + inline + void grid_launch_hip_impl_( + New_grid_launch_tag, + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + const char* kernel_name, + K k, + Ts&&... args) + { + void* lck_stream = nullptr; + auto acc_v = lock_stream_hip_(stream, lck_stream); + auto stream_guard = make_RAII_guard( + [](){ /* perhaps use a slimmed down ihipPrintKernelLaunch here */ }, + std::bind( + unlock_stream_hip_, stream, lck_stream, kernel_name, &acc_v)); + + try { + grid_launch_hip_impl_( + New_grid_launch_tag{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + acc_v, + std::move(k), + std::forward(args)...); + } + catch (std::exception& ex) { + std::cerr << "Failed in " << __FUNCTION__ << ", with exception: " + << ex.what() << std::endl; + throw; + } + } + + template + requires(Domain == {hipLaunchParm, Ts...}) + inline + void grid_launch_hip_impl_( + Old_grid_launch_tag, + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + K k, + Ts&&... args) + { + grid_launch_hip_impl_( + New_grid_launch_tag{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + std::move(k), + hipLaunchParm{}, + std::forward(args)...); + } + + template + requires(Domain == {hipLaunchParm, Ts...}) + inline + void grid_launch_hip_impl_( + Old_grid_launch_tag, + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + const char* kernel_name, + K k, + Ts&&... args) + { + grid_launch_hip_impl_( + New_grid_launch_tag{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + kernel_name, + std::move(k), + hipLaunchParm{}, + std::forward(args)...); + } + + template + requires(Domain == {Ts...}) + inline + std::enable_if_t::value> grid_launch_hip_( + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + const char* kernel_name, + K k, + Ts&& ... args) + { + grid_launch_hip_impl_( + is_new_grid_launch_t{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + kernel_name, + std::move(k), + std::forward(args)...); + } + + template + requires(Domain == {Ts...}) + inline + std::enable_if_t::value> grid_launch_hip_( + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + K k, + Ts&& ... args) + { + grid_launch_hip_impl_( + is_new_grid_launch_t{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + std::move(k), + std::forward(args)...); + } + + namespace + { + template + constexpr + inline + T&& forward_(std::remove_reference_t& x) [[hc]] + { + return static_cast(x); + } + + template + struct Forwarder { + template + void operator()(Ts&&...args) const [[hc]] + { + k(forward_(args)...); + } + }; + } + + template + requires(Domain == {Ts...}) + inline + void grid_launch( + New_grid_launch_tag, + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + Ts&&... args) + { + grid_launch_hip_impl_( + New_grid_launch_tag{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + Forwarder{}, + std::forward(args)...); + } + + template + requires(Domain == {Ts...}) + inline + void grid_launch( + Old_grid_launch_tag, + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + Ts&&... args) + { + grid_launch_hip_( + New_grid_launch_tag{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + hipLaunchParm{}, + std::forward(args)...); + } + + template + requires(Domain == {Ts...}) + inline + std::enable_if_t::value> grid_launch_hip_( + dim3 num_blocks, + dim3 dim_blocks, + int group_mem_bytes, + hipStream_t stream, + Ts&&... args) + { + grid_launch_hip_( + is_new_grid_launch_t{}, + std::move(num_blocks), + std::move(dim_blocks), + group_mem_bytes, + std::move(stream), + std::forward(args)...); + } + + // TODO: these are temporary, they need to be completely removed once we + // enable C++14 support and can have proper generic, variadic lambdas. + #define make_kernel_lambda_hip_26(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19, p20, p21, p22, p23, p24)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_,\ + const std::decay_t& _p20_,\ + const std::decay_t& _p21_,\ + const std::decay_t& _p22_,\ + const std::decay_t& _p23_,\ + const std::decay_t& _p24_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_, _p20_, _p21_, _p22_, _p23_, _p24_);\ + } + #define make_kernel_lambda_hip_25(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19, p20, p21, p22, p23)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_,\ + const std::decay_t& _p20_,\ + const std::decay_t& _p21_,\ + const std::decay_t& _p22_,\ + const std::decay_t& _p23_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_, _p20_, _p21_, _p22_, _p23_);\ + } + #define make_kernel_lambda_hip_24(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19, p20, p21, p22)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_,\ + const std::decay_t& _p20_,\ + const std::decay_t& _p21_,\ + const std::decay_t& _p22_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_, _p20_, _p21_, _p22_);\ + } + #define make_kernel_lambda_hip_23(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19, p20, p21)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_,\ + const std::decay_t& _p20_,\ + const std::decay_t& _p21_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_, _p20_, _p21_);\ + } + #define make_kernel_lambda_hip_22(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19, p20)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_,\ + const std::decay_t& _p20_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_, _p20_);\ + } + #define make_kernel_lambda_hip_21(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_);\ + } + #define make_kernel_lambda_hip_20(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_);\ + } + #define make_kernel_lambda_hip_19(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_);\ + } + #define make_kernel_lambda_hip_18(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_);\ + } + #define make_kernel_lambda_hip_17(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_);\ + } + #define make_kernel_lambda_hip_16(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_);\ + } + #define make_kernel_lambda_hip_15(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_);\ + } + #define make_kernel_lambda_hip_14(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_);\ + } + #define make_kernel_lambda_hip_13(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_);\ + } + #define make_kernel_lambda_hip_12(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_,\ + _p10_);\ + } + #define make_kernel_lambda_hip_11(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_);\ + } + #define make_kernel_lambda_hip_10(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_);\ + } + #define make_kernel_lambda_hip_9(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_);\ + } + #define make_kernel_lambda_hip_8(kernel_name, p0, p1, p2, p3, p4, p5, p6)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_);\ + } + #define make_kernel_lambda_hip_7(kernel_name, p0, p1, p2, p3, p4, p5)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_);\ + } + #define make_kernel_lambda_hip_6(kernel_name, p0, p1, p2, p3, p4)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_);\ + } + #define make_kernel_lambda_hip_5(kernel_name, p0, p1, p2, p3)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_);\ + } + #define make_kernel_lambda_hip_4(kernel_name, p0, p1, p2)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_);\ + } + #define make_kernel_lambda_hip_3(kernel_name, p0, p1)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_) [[hc]] {\ + kernel_name(_p0_, _p1_);\ + } + #define make_kernel_lambda_hip_2(kernel_name, p0)\ + [](const std::decay_t& _p0_) [[hc]] {\ + kernel_name(_p0_);\ + } + #define make_kernel_lambda_hip_1(kernel_name)\ + []() [[hc]] { return kernel_name(hipLaunchParm{}); } + + #define make_kernel_lambda_hip_(...)\ + overload_macro_hip_(make_kernel_lambda_hip_, __VA_ARGS__) + + #define hipLaunchKernelGGL(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + ...)\ + do {\ + hip_impl::grid_launch_hip_(\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + #kernel_name,\ + make_kernel_lambda_hip_(kernel_name, __VA_ARGS__),\ + ##__VA_ARGS__);\ + } while(0) + + #define hipLaunchKernel(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + ...)\ + {\ + hipLaunchKernelGGL(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + hipLaunchParm{},\ + ##__VA_ARGS__);\ + } +} diff --git a/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp b/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp new file mode 100644 index 0000000000..e5a84a4678 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp @@ -0,0 +1,144 @@ +/* +Copyright (c) 2015 - present 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. +*/ + +#pragma once + +#include // For std::conditional, std::decay, std::enable_if, + // std::false_type, std result_of and std::true_type. +#include // For std::declval. + +namespace std +{ // TODO: these should be removed as soon as possible. + #if (__cplusplus < 201406L) + template + using void_t = void; + + #if (__cplusplus < 201402L) + template + using enable_if_t = typename enable_if::type; + template + using conditional_t = typename conditional::type; + template + using decay_t = typename decay::type; + template + using result_of_t = typename result_of::type; + template + using remove_reference_t = typename remove_reference::type; + template< + FunctionalProcedure F, + unsigned int n = 0u, + typename = void> + struct is_callable_impl : is_callable_impl {}; + + // Pointer to member function, call through non-pointer. + template + struct is_callable_impl< + F(C, Ts...), + 0u, + void_t().*declval())(declval()...))> + > : true_type { + }; + + // Pointer to member function, call through pointer. + template + struct is_callable_impl< + F(C, Ts...), + 1u, + void_t()).*declval())(declval()...))> + > : std::true_type { + }; + + // Pointer to member data, call through non-pointer, no args. + template + struct is_callable_impl< + F(C), + 2u, + void_t().*declval())> + > : true_type { + }; + + // Pointer to member data, call through pointer, no args. + template + struct is_callable_impl< + F(C), + 3u, + void_t().*declval())> + > : true_type { + }; + + // General call, n args. + template + struct is_callable_impl< + F(Ts...), + 4u, + void_t()(declval()...))> + > : true_type { + }; + + // Not callable. + template + struct is_callable_impl : false_type {}; + + template + struct is_callable : is_callable_impl {}; + #else + template + struct is_callable_impl : false_type {}; + + template + struct is_callable_impl< + F(Ts...), + void_t>> : true_type {}; + + template + struct is_callable : is_callable_impl {}; + #endif + template + struct disjunction : false_type {}; + template + struct disjunction : B1 {}; + template + struct disjunction + : conditional_t> + {}; + #endif +} + +namespace hip_impl // Only for documentation, macros ignore namespaces. +{ + #define count_macro_args_impl_hip_(\ + _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15,\ + _16, _17, _18, _19, _20, _21, _22, _23, _24, _25, _26, _27, _28, _29,\ + _30, _31, _n, ...)\ + _n + #define count_macro_args_hip_(...)\ + count_macro_args_impl_hip_(\ + , ##__VA_ARGS__, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19,\ + 18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0) + + #define overloaded_macro_expand_hip_(macro, arg_cnt) macro##arg_cnt + #define overload_macro_impl_hip_(macro, arg_cnt)\ + overloaded_macro_expand_hip_(macro, arg_cnt) + #define overload_macro_hip_(macro, ...)\ + overload_macro_impl_hip_(macro, count_macro_args_hip_(__VA_ARGS__))\ + (__VA_ARGS__) +} diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_complex.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_complex.h index f50a601b90..9ff75d381a 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_complex.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_complex.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H -#define INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #include "./hip_fp16.h" #include "./hip_vector_types.h" diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h index b0276d51b3..febc1b4fce 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_DETAIL_FP16_H -#define HIP_HCC_DETAIL_FP16_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H #include "hip/hip_runtime.h" @@ -464,9 +464,6 @@ __device__ static inline __half2 h2trunc(const __half2 h) { return a; } - - - #endif #if __clang_major__ == 3 diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_hcc.h new file mode 100644 index 0000000000..645e980376 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_hcc.h @@ -0,0 +1,62 @@ +/* +Copyright (c) 2015 - present 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_INCLUDE_HIP_HCC_DETAIL_HIP_HCC_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_HCC_H + +#include "hip/hip_runtime_api.h" + +#if __cplusplus +#ifdef __HCC__ +#include +/** + * @brief Return hc::accelerator associated with the specified deviceId + * @return #hipSuccess, #hipErrorInvalidDevice + */ +hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc); + +/** + * @brief Return hc::accelerator_view associated with the specified stream + * + * If stream is 0, the accelerator_view for the default stream is returned. + * @return #hipSuccess + */ +hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av); + + +#endif // #ifdef __HCC__ + +hipError_t hipHccModuleLaunchKernel(hipFunction_t f, + uint32_t globalWorkSizeX, + uint32_t globalWorkSizeY, + uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, + uint32_t localWorkSizeY, + uint32_t localWorkSizeZ, + size_t sharedMemBytes, + hipStream_t hStream, + void **kernelParams, + void **extra); + +#endif // #if __cplusplus + +#endif // diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_ldg.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_ldg.h index 65292951f0..473e70b4cb 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_ldg.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_ldg.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,15 +20,14 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_LDG_H -#define HIP_LDG_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H -#if __HCC__ +#if defined __HCC__ #if __hcc_workweek__ >= 16164 #include "hip_vector_types.h" #include "host_defines.h" - __device__ char __ldg(const char* ); __device__ char2 __ldg(const char2* ); __device__ char4 __ldg(const char4* ); diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 332e9bab46..06ce65bc9a 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -26,13 +26,12 @@ THE SOFTWARE. */ //#pragma once -#ifndef HIP_HCC_DETAIL_RUNTIME_H -#define HIP_HCC_DETAIL_RUNTIME_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H //--- // Top part of file can be compiled with any compiler - //#include #if __cplusplus #include @@ -40,30 +39,53 @@ THE SOFTWARE. #include #include #include -#endif +#endif//__cplusplus + // Define NVCC_COMPAT for CUDA compatibility #define NVCC_COMPAT #define CUDA_SUCCESS hipSuccess #include -//#include "hip/hcc_detail/hip_hcc.h" + +#if USE_PROMOTE_FREE_HCC == 1 +#define ADDRESS_SPACE_1 +#define ADDRESS_SPACE_3 +#else +#define ADDRESS_SPACE_1 __attribute__((address_space(1))) +#define ADDRESS_SPACE_3 __attribute__((address_space(3))) +#endif + //--- // Remainder of this file only compiles with HCC -#ifdef __HCC__ +#if defined __HCC__ #include +//TODO-HCC-GL - change this to typedef. +//typedef grid_launch_parm hipLaunchParm ; -#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) -// Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0. +#if GENERIC_GRID_LAUNCH == 0 + #define hipLaunchParm grid_launch_parm #else +namespace hip_impl +{ + struct Empty_launch_parm {}; +} +#define hipLaunchParm hip_impl::Empty_launch_parm +#endif //GENERIC_GRID_LAUNCH + +#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) || GENERIC_GRID_LAUNCH == 1 +#else // Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0. #error (HCC must support GRID_LAUNCH_20) -#endif +#endif //GRID_LAUNCH_VERSION + +#endif //HCC + +#if GENERIC_GRID_LAUNCH==1 && defined __HCC__ +#include "grid_launch_GGL.hpp" +#endif//GENERIC_GRID_LAUNCH 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 @@ -77,6 +99,17 @@ extern int HIP_TRACE_API; #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__ #endif + + + +// TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call. +#if __HIP_DEVICE_COMPILE__ == 1 + #undef assert + #define assert(COND) { if (COND) {} } +#endif + + + // Feature tests: #if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0) // Device compile and not host compile: @@ -126,8 +159,7 @@ extern int HIP_TRACE_API; // TODO - hipify-clang - change to use the function call. //#define warpSize hc::__wavesize() -extern const int warpSize; - +static constexpr int warpSize = 64; #define clock_t long long int __device__ long long int clock64(); @@ -247,7 +279,7 @@ __device__ float __shfl(float input, int lane, int width); __device__ float __shfl_up(float input, unsigned int lane_delta, int width); __device__ float __shfl_down(float input, unsigned int lane_delta, int width); __device__ float __shfl_xor(float input, int lane_mask, int width); -#endif +#endif //__cplusplus __device__ unsigned __hip_ds_bpermute(int index, unsigned src); __device__ float __hip_ds_bpermutef(int index, float src); @@ -259,12 +291,12 @@ __device__ float __hip_ds_swizzlef(float src, int pattern); __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl); -#endif +#endif //__HIP_ARCH_GFX803__ == 1 __host__ __device__ int min(int arg1, int arg2); __host__ __device__ int max(int arg1, int arg2); -__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr(); +__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr(); /** @@ -387,17 +419,18 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) -#define HIP_KERNEL_NAME(...) __VA_ARGS__ +#define HIP_KERNEL_NAME(...) (__VA_ARGS__) #define HIP_SYMBOL(X) #X -#ifdef __HCC_CPP__ +#if defined __HCC_CPP__ extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr); extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr); extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr); extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr); extern void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp); - +#if GENERIC_GRID_LAUNCH == 0 +//#warning "Original hipLaunchKernel defined" // Due to multiple overloaded versions of ihipPreLaunchKernel, the numBlocks3D and blockDim3D can be either size_t or dim3 types #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ do {\ @@ -407,13 +440,13 @@ do {\ _kernelName (lp, ##__VA_ARGS__);\ ihipPostLaunchKernel(#_kernelName, trueStream, lp);\ } while(0) - +#endif //GENERIC_GRID_LAUNCH #elif defined (__HCC_C__) //TODO - develop C interface. -#endif +#endif //__HCC_CPP__ /** * extern __shared__ @@ -422,12 +455,11 @@ do {\ // Macro to replace extern __shared__ declarations // to local variable definitions #define HIP_DYNAMIC_SHARED(type, var) \ - __attribute__((address_space(3))) type* var = \ - (__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \ + ADDRESS_SPACE_3 type* var = \ + (ADDRESS_SPACE_3 type*)__get_dynamicgroupbaseptr(); \ -#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3))) +#define HIP_DYNAMIC_SHARED_ATTRIBUTE ADDRESS_SPACE_3 -#endif // __HCC__ /** @@ -451,4 +483,4 @@ do {\ -#endif +#endif//HIP_HCC_DETAIL_RUNTIME_H diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 0d3ecc6613..0daca7a53b 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -21,17 +21,20 @@ THE SOFTWARE. */ //#pragma once -#ifndef HIP_HCC_DETAIL_HIP_RUNTIME_API_H -#define HIP_HCC_DETAIL_HIP_RUNTIME_API_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_API_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_API_H /** * @file hcc_detail/hip_runtime_api.h * @brief Contains C function APIs for HIP runtime. This file does not use any HCC builtin or special language extensions (-hc mode) ; those functions in hip_runtime.h. */ - #include #include #include +#ifndef GENERIC_GRID_LAUNCH +#define GENERIC_GRID_LAUNCH 1 +#endif + #include #include #include @@ -1153,7 +1156,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz * * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync */ -hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind); +hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind); /** @@ -1173,11 +1176,11 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz * * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync */ -hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); +hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); -hipError_t hipMemcpyFromSymbol(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind); +hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind); -hipError_t hipMemcpyFromSymbolAsync(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); +hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); /** * @brief Copy data from src to dst asynchronously. @@ -1693,6 +1696,66 @@ hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags); */ hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx); +/** + * @brief Get the state of the primary context. + * + * @param [in] Device to get primary context flags for + * @param [out] Pointer to store flags + * @param [out] Pointer to store context state; 0 = inactive, 1 = active + * + * @returns #hipSuccess + * + * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice + */ +hipError_t hipDevicePrimaryCtxGetState ( hipDevice_t dev, unsigned int* flags, int* active ); + +/** + * @brief Release the primary context on the GPU. + * + * @param [in] Device which primary context is released + * + * @returns #hipSuccess + * + * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice + * @warning This function return #hipSuccess though doesn't release the primaryCtx by design on HIP/HCC path. + */ +hipError_t hipDevicePrimaryCtxRelease ( hipDevice_t dev); + +/** + * @brief Retain the primary context on the GPU. + * + * @param [out] Returned context handle of the new context + * @param [in] Device which primary context is released + * + * @returns #hipSuccess + * + * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice + */ +hipError_t hipDevicePrimaryCtxRetain ( hipCtx_t* pctx, hipDevice_t dev ); + +/** + * @brief Resets the primary context on the GPU. + * + * @param [in] Device which primary context is reset + * + * @returns #hipSuccess + * + * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice + */ +hipError_t hipDevicePrimaryCtxReset ( hipDevice_t dev ); + +/** + * @brief Set flags for the primary context. + * + * @param [in] Device for which the primary context flags are set + * @param [in] New flags for the device + * + * @returns #hipSuccess, #hipErrorContextAlreadyInUse + * + * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice + */ +hipError_t hipDevicePrimaryCtxSetFlags ( hipDevice_t dev, unsigned int flags ); + // doxygen end Context Management /** * @} diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_texture.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_texture.h index d60f66ddc2..c6f5a1cfb2 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_texture.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_texture.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -22,8 +22,8 @@ THE SOFTWARE. //#pragma once -#ifndef HIP_HCC_DETAIL_TEXTURE_H -#define HIP_HCC_DETAIL_TEXTURE_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_TEXTURE_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_TEXTURE_H /** * @file hcc_detail/hip_texture.h diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index 8e6ec49511..42e1d6663c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -25,8 +25,8 @@ THE SOFTWARE. * @brief Defines the different newt vector types for HIP runtime. */ -#ifndef HIP_VECTOR_TYPES_H -#define HIP_VECTOR_TYPES_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H #if defined (__HCC__) && (__hcc_workweek__ < 16032) #error("This version of HIP requires a newer version of HCC."); diff --git a/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h b/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h index e401cb24f3..b0a7421d18 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -25,8 +25,15 @@ THE SOFTWARE. * @brief TODO-doc */ -#ifndef HOST_DEFINES_H -#define HOST_DEFINES_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HOST_DEFINES_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HOST_DEFINES_H + +#define USE_PROMOTE_FREE_HCC 1 + +// Add guard to Generic Grid Launch method +#ifndef GENERIC_GRID_LAUNCH +#define GENERIC_GRID_LAUNCH 1 +#endif #ifdef __HCC__ /** @@ -35,7 +42,14 @@ THE SOFTWARE. #define __host__ __attribute__((cpu)) #define __device__ __attribute__((hc)) +//#warning "HOST DEFINE header included" +#if GENERIC_GRID_LAUNCH == 0 +//#warning "original global define reached" #define __global__ __attribute__((hc_grid_launch)) __attribute__((used)) +#else +//#warning "GGL global define reached" +#define __global__ [[hc]] __attribute__((weak)) +#endif //GENERIC_GRID_LAUNCH #define __noinline__ __attribute__((noinline)) #define __forceinline__ __attribute__((always_inline)) @@ -47,7 +61,11 @@ THE SOFTWARE. */ // _restrict is supported by the compiler #define __shared__ tile_static -#define __constant__ __attribute__((address_space(1))) +#if USE_PROMOTE_FREE_HCC==1 +#define __constant__ __attribute__((hc)) +#else +#define __constant__ ADDRESS_SPACE_1 +#endif #else // Non-HCC compiler diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index 2cc4ef81bd..c3b8186fd3 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h @@ -1,13 +1,16 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -17,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_DETAIL_MATH_FUNCTIONS_H -#define HIP_HCC_DETAIL_MATH_FUNCTIONS_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H #include #include @@ -66,7 +69,7 @@ __device__ float j0f(float x); __device__ float j1f(float x); __device__ float jnf(int n, float x); __device__ float ldexpf(float x, int exp); -//__device__ float lgammaf(float x); +__device__ float lgammaf(float x); __device__ long long int llrintf(float x); __device__ long long int llroundf(float x); __device__ float log10f(float x); diff --git a/projects/clr/hipamd/include/hip/hcc_detail/texture_types.h b/projects/clr/hipamd/include/hip/hcc_detail/texture_types.h index 107b5d26c1..74680bbc76 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/texture_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/texture_types.h @@ -1,13 +1,16 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -18,8 +21,8 @@ THE SOFTWARE. */ -#ifndef HIP_HCC_DETAIL_TEXTURE_TYPES_H -#define HIP_HCC_DETAIL_TEXTURE_TYPES_H +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_TYPES_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_TEXTURE_TYPES_H #include diff --git a/projects/clr/hipamd/include/hip/hip_common.h b/projects/clr/hipamd/include/hip/hip_common.h index 6317a792ee..da8ec4a55d 100644 --- a/projects/clr/hipamd/include/hip/hip_common.h +++ b/projects/clr/hipamd/include/hip/hip_common.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_HIP_COMMON_H +#define HIP_INCLUDE_HIP_HIP_COMMON_H // Common code included at start of every hip file. // Auto enable __HIP_PLATFORM_HCC__ if compiling with HCC @@ -73,3 +74,5 @@ THE SOFTWARE. #define __HIP_ARCH_HAS_3DGRID__ (0) #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0) #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/hip_complex.h b/projects/clr/hipamd/include/hip/hip_complex.h index ea15137894..dc691be480 100644 --- a/projects/clr/hipamd/include/hip/hip_complex.h +++ b/projects/clr/hipamd/include/hip/hip_complex.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_HIP_COMPLEX_H +#define HIP_INCLUDE_HIP_HIP_COMPLEX_H #include @@ -31,3 +32,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/hip_fp16.h b/projects/clr/hipamd/include/hip/hip_fp16.h index 2f64c1a143..0e002d9396 100644 --- a/projects/clr/hipamd/include/hip/hip_fp16.h +++ b/projects/clr/hipamd/include/hip/hip_fp16.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifdef HIP_INCLUDE_HIP_HIP_FP16_H +#define HIP_INCLUDE_HIP_HIP_FP16_H #include @@ -31,3 +32,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/hcc.h b/projects/clr/hipamd/include/hip/hip_hcc.h similarity index 85% rename from projects/clr/hipamd/include/hip/hcc.h rename to projects/clr/hipamd/include/hip/hip_hcc.h index 9b8a649412..3407a311bd 100644 --- a/projects/clr/hipamd/include/hip/hcc.h +++ b/projects/clr/hipamd/include/hip/hip_hcc.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,11 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_H -#define HIP_HCC_H +#ifndef HIP_INCLUDE_HIP_HIP_HCC_H +#define HIP_INCLUDE_HIP_HIP_HCC_H #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) -#include "hip/hcc_detail/hcc_acc.h" +#include "hip/hcc_detail/hip_hcc.h" #endif #endif diff --git a/projects/clr/hipamd/include/hip/hip_profile.h b/projects/clr/hipamd/include/hip/hip_profile.h index e621ae8c79..389f334c74 100644 --- a/projects/clr/hipamd/include/hip/hip_profile.h +++ b/projects/clr/hipamd/include/hip/hip_profile.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_HIP_PROFILE_H +#define HIP_INCLUDE_HIP_HIP_PROFILE_H #if not defined (ENABLE_HIP_PROFILE) #define ENABLE_HIP_PROFILE 1 @@ -36,3 +37,5 @@ THE SOFTWARE. #define HIP_BEGIN_MARKER(markerName, group) #define HIP_END_MARKER() #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/hip_runtime.h b/projects/clr/hipamd/include/hip/hip_runtime.h index 9bc45f300d..fba4d46d8f 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hip_runtime.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -33,7 +33,8 @@ THE SOFTWARE. //! hip_runtime.h : includes everything in hip_api.h, plus math builtins and kernel launch macros. //! hip_runtime_api.h : Defines HIP API. This is a C header file and does not use any C++ features. -#pragma once +#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_H +#define HIP_INCLUDE_HIP_HIP_RUNTIME_H // Some standard header files, these are included by hc.hpp and so want to make them avail on both // paths to provide a consistent include env and avoid "missing symbol" errors that only appears @@ -61,3 +62,5 @@ THE SOFTWARE. #include #include + +#endif diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index 818c0b7c34..8eae1d6a3a 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -27,7 +27,8 @@ THE SOFTWARE. * This file can be compiled with a standard compiler. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_API_H +#define HIP_INCLUDE_HIP_HIP_RUNTIME_API_H #include // for getDeviceProp @@ -283,3 +284,5 @@ static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags return hipHostMalloc((void**)ptr, size, flags); } #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/hip_texture.h b/projects/clr/hipamd/include/hip/hip_texture.h index 66ec4a6ca1..a15c5a1016 100644 --- a/projects/clr/hipamd/include/hip/hip_texture.h +++ b/projects/clr/hipamd/include/hip/hip_texture.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,10 +20,8 @@ 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 +#ifndef HIP_INCLUDE_HIP_HIP_TEXTURE_H +#define HIP_INCLUDE_HIP_HIP_TEXTURE_H #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) #include diff --git a/projects/clr/hipamd/include/hip/hip_vector_types.h b/projects/clr/hipamd/include/hip/hip_vector_types.h index 33827e4d96..1d3d6b92f6 100644 --- a/projects/clr/hipamd/include/hip/hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/hip_vector_types.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -22,7 +22,8 @@ THE SOFTWARE. //! hip_vector_types.h : Defines the HIP vector types. -#pragma once +#ifndef HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H +#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H #include @@ -36,3 +37,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/math_functions.h b/projects/clr/hipamd/include/hip/math_functions.h index ebcdc26749..6f47b5e0e2 100644 --- a/projects/clr/hipamd/include/hip/math_functions.h +++ b/projects/clr/hipamd/include/hip/math_functions.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_MATH_FUNCTIONS_H +#define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H // Some standard header files, these are included by hc.hpp and so want to make them avail on both // paths to provide a consistent include env and avoid "missing symbol" errors that only appears @@ -34,3 +35,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/channel_descriptor.h b/projects/clr/hipamd/include/hip/nvcc_detail/channel_descriptor.h index 8502745968..2e88c56268 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/channel_descriptor.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/channel_descriptor.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,6 +20,9 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_CHANNEL_DESCRIPTOR_H +#define HIP_INCLUDE_HIP_NVCC_DETAIL_CHANNEL_DESCRIPTOR_H #include"channel_descriptor.h" + +#endif diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h index 174cabc12c..84afb13e50 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIPCOMPLEX_H -#define HIPCOMPLEX_H +#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COMPLEX_H +#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COMPLEX_H #include"cuComplex.h" diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h index 2c774bfb7d..b4fa13f48c 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_H +#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_H #include @@ -105,3 +106,5 @@ kernelName<<>>(0, ##__VA_ARGS__);\ #define HIP_DYNAMIC_SHARED_ATTRIBUTE #endif + +#endif diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 8c3e0da639..7e881df3ab 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H +#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H #include #include @@ -333,6 +334,16 @@ inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* return hipCUDAErrorTohipError(cudaMemcpyToSymbolAsync(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType))); } +inline static hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind) +{ + return hipCUDAErrorTohipError(cudaMemcpyFromSymbol(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind))); +} + +inline static hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream) +{ + return hipCUDAErrorTohipError(cudaMemcpyFromSymbolAsync(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind), stream)); +} + inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind){ return hipCUDAErrorTohipError(cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind))); } @@ -670,6 +681,31 @@ inline static hipError_t hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned in return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags)); } +inline static hipError_t hipDevicePrimaryCtxGetState ( hipDevice_t dev, unsigned int* flags, int* active ) +{ + return hipCUResultTohipError(cuDevicePrimaryCtxGetState(dev, flags, active)); +} + +inline static hipError_t hipDevicePrimaryCtxRelease ( hipDevice_t dev) +{ + return hipCUResultTohipError(cuDevicePrimaryCtxRelease(dev)); +} + +inline static hipError_t hipDevicePrimaryCtxRetain ( hipCtx_t* pctx, hipDevice_t dev ) +{ + return hipCUResultTohipError(cuDevicePrimaryCtxRetain(pctx, dev)); +} + +inline static hipError_t hipDevicePrimaryCtxReset ( hipDevice_t dev ) +{ + return hipCUResultTohipError(cuDevicePrimaryCtxReset(dev)); +} + +inline static hipError_t hipDevicePrimaryCtxSetFlags ( hipDevice_t dev, unsigned int flags ) +{ + return hipCUResultTohipError(cuDevicePrimaryCtxSetFlags(dev, flags)); +} + inline static hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr ) { return hipCUResultTohipError(cuMemGetAddressRange( pbase , psize , dptr)); diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_texture.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_texture.h index 388733e492..c669d62192 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_texture.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_texture.h @@ -1,5 +1,27 @@ -#ifndef HIP_TEXTURE_H -#define HIP_TEXTURE_H +/* +Copyright (c) 2015 - present 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_INCLUDE_HIP_NVCC_DETAIL_HIP_TEXTURE_H +#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_TEXTURE_H #include diff --git a/projects/clr/hipamd/include/hip/texture_types.h b/projects/clr/hipamd/include/hip/texture_types.h index 2561e12eb5..ca6101cf79 100644 --- a/projects/clr/hipamd/include/hip/texture_types.h +++ b/projects/clr/hipamd/include/hip/texture_types.h @@ -1,13 +1,16 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -17,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#pragma once +#ifndef HIP_INCLUDE_HIP_TEXTURE_TYPES_H +#define HIP_INCLUDE_HIP_TEXTURE_TYPES_H #include @@ -28,3 +32,5 @@ THE SOFTWARE. #else #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); #endif + +#endif diff --git a/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile b/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile index 0965ae7296..78f6a2faa8 100644 --- a/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile +++ b/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile @@ -24,4 +24,3 @@ $(EXE): bit_extract.cpp clean: rm -f *.o $(EXE) - diff --git a/projects/clr/hipamd/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/clr/hipamd/samples/0_Intro/bit_extract/bit_extract.cpp index 1535d2bd98..a30f2d052d 100644 --- a/projects/clr/hipamd/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/projects/clr/hipamd/samples/0_Intro/bit_extract/bit_extract.cpp @@ -37,7 +37,7 @@ THE SOFTWARE. }\ } -void __global__ +__global__ void bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N) { size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); @@ -45,7 +45,7 @@ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t for (size_t i=offset; i> 8); #endif @@ -73,7 +73,7 @@ int main(int argc, char *argv[]) C_h = (uint32_t*)malloc(Nbytes); CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); - for (size_t i=0; i +#include +#include + +#define LEN 64 +#define SIZE LEN<<2 + +#define fileName "test.co" +#define kernel_name "vadd" + +int main(){ + float *A, *B, *C; + hipDeviceptr_t Ad, Bd, Cd; + A = new float[LEN]; + B = new float[LEN]; + C = new float[LEN]; + + for(uint32_t i=0;i +#include +#include + +#ifdef __HIP_PLATFORM_HCC__ +#include +#endif + +#define LEN 64 +#define SIZE LEN<<2 + +#define fileName "vcpy_kernel.code" +#define kernel_name "hello_world" + +#define HIP_CHECK(status) \ +if(status != hipSuccess) {std::cout<<"Got Status: "< -#include -#include +#include +#include +#include +#include #define LEN 64 #define SIZE LEN<<2 @@ -32,6 +33,9 @@ THE SOFTWARE. #define fileName "vcpy_kernel.code" #define kernel_name "hello_world" +#define HIP_CHECK(status) \ +if(status != hipSuccess) {std::cout<<"Got Status: "< +#include +#include + +namespace AMDGPU { +namespace RuntimeMD { + + // Version and revision of runtime metadata + const unsigned char MDVersion = 2; + const unsigned char MDRevision = 1; + + // Name of keys for runtime metadata. + namespace KeyName { + + // Runtime metadata version + const char MDVersion[] = "amd.MDVersion"; + + // Instruction set architecture information + const char IsaInfo[] = "amd.IsaInfo"; + // Wavefront size + const char IsaInfoWavefrontSize[] = "amd.IsaInfoWavefrontSize"; + // Local memory size in bytes + const char IsaInfoLocalMemorySize[] = "amd.IsaInfoLocalMemorySize"; + // Number of execution units per compute unit + const char IsaInfoEUsPerCU[] = "amd.IsaInfoEUsPerCU"; + // Maximum number of waves per execution unit + const char IsaInfoMaxWavesPerEU[] = "amd.IsaInfoMaxWavesPerEU"; + // Maximum flat work group size + const char IsaInfoMaxFlatWorkGroupSize[] = "amd.IsaInfoMaxFlatWorkGroupSize"; + // SGPR allocation granularity + const char IsaInfoSGPRAllocGranule[] = "amd.IsaInfoSGPRAllocGranule"; + // Total number of SGPRs + const char IsaInfoTotalNumSGPRs[] = "amd.IsaInfoTotalNumSGPRs"; + // Addressable number of SGPRs + const char IsaInfoAddressableNumSGPRs[] = "amd.IsaInfoAddressableNumSGPRs"; + // VGPR allocation granularity + const char IsaInfoVGPRAllocGranule[] = "amd.IsaInfoVGPRAllocGranule"; + // Total number of VGPRs + const char IsaInfoTotalNumVGPRs[] = "amd.IsaInfoTotalNumVGPRs"; + // Addressable number of VGPRs + const char IsaInfoAddressableNumVGPRs[] = "amd.IsaInfoAddressableNumVGPRs"; + + // Language + const char Language[] = "amd.Language"; + // Language version + const char LanguageVersion[] = "amd.LanguageVersion"; + + // Kernels + const char Kernels[] = "amd.Kernels"; + // Kernel name + const char KernelName[] = "amd.KernelName"; + // Kernel arguments + const char Args[] = "amd.Args"; + // Kernel argument size in bytes + const char ArgSize[] = "amd.ArgSize"; + // Kernel argument alignment + const char ArgAlign[] = "amd.ArgAlign"; + // Kernel argument type name + const char ArgTypeName[] = "amd.ArgTypeName"; + // Kernel argument name + const char ArgName[] = "amd.ArgName"; + // Kernel argument kind + const char ArgKind[] = "amd.ArgKind"; + // Kernel argument value type + const char ArgValueType[] = "amd.ArgValueType"; + // Kernel argument address qualifier + const char ArgAddrQual[] = "amd.ArgAddrQual"; + // Kernel argument access qualifier + const char ArgAccQual[] = "amd.ArgAccQual"; + // Kernel argument is const qualified + const char ArgIsConst[] = "amd.ArgIsConst"; + // Kernel argument is restrict qualified + const char ArgIsRestrict[] = "amd.ArgIsRestrict"; + // Kernel argument is volatile qualified + const char ArgIsVolatile[] = "amd.ArgIsVolatile"; + // Kernel argument is pipe qualified + const char ArgIsPipe[] = "amd.ArgIsPipe"; + // Required work group size + const char ReqdWorkGroupSize[] = "amd.ReqdWorkGroupSize"; + // Work group size hint + const char WorkGroupSizeHint[] = "amd.WorkGroupSizeHint"; + // Vector type hint + const char VecTypeHint[] = "amd.VecTypeHint"; + // Kernel index for device enqueue + const char KernelIndex[] = "amd.KernelIndex"; + // No partial work groups + const char NoPartialWorkGroups[] = "amd.NoPartialWorkGroups"; + // Prinf function call information + const char PrintfInfo[] = "amd.PrintfInfo"; + // The actual kernel argument access qualifier + const char ArgActualAcc[] = "amd.ArgActualAcc"; + // Alignment of pointee type + const char ArgPointeeAlign[] = "amd.ArgPointeeAlign"; + + } // end namespace KeyName + + namespace KernelArg { + + enum Kind : uint8_t { + ByValue = 0, + GlobalBuffer = 1, + DynamicSharedPointer = 2, + Sampler = 3, + Image = 4, + Pipe = 5, + Queue = 6, + HiddenGlobalOffsetX = 7, + HiddenGlobalOffsetY = 8, + HiddenGlobalOffsetZ = 9, + HiddenNone = 10, + HiddenPrintfBuffer = 11, + HiddenDefaultQueue = 12, + HiddenCompletionAction = 13, + }; + + enum ValueType : uint16_t { + Struct = 0, + I8 = 1, + U8 = 2, + I16 = 3, + U16 = 4, + F16 = 5, + I32 = 6, + U32 = 7, + F32 = 8, + I64 = 9, + U64 = 10, + F64 = 11, + }; + + // Avoid using 'None' since it conflicts with a macro in X11 header file. + enum AccessQualifer : uint8_t { + AccNone = 0, + ReadOnly = 1, + WriteOnly = 2, + ReadWrite = 3, + }; + + enum AddressSpaceQualifer : uint8_t { + Private = 0, + Global = 1, + Constant = 2, + Local = 3, + Generic = 4, + Region = 5, + }; + + } // end namespace KernelArg + + // Invalid values are used to indicate an optional key should not be emitted. + const uint8_t INVALID_ADDR_QUAL = 0xff; + const uint8_t INVALID_ACC_QUAL = 0xff; + const uint32_t INVALID_KERNEL_INDEX = ~0U; + + namespace KernelArg { + + // In-memory representation of kernel argument information. + struct Metadata { + uint32_t Size = 0; + uint32_t Align = 0; + uint32_t PointeeAlign = 0; + uint8_t Kind = 0; + uint16_t ValueType = 0; + std::string TypeName; + std::string Name; + uint8_t AddrQual = INVALID_ADDR_QUAL; + uint8_t AccQual = INVALID_ACC_QUAL; + uint8_t IsVolatile = 0; + uint8_t IsConst = 0; + uint8_t IsRestrict = 0; + uint8_t IsPipe = 0; + + Metadata() = default; + }; + + } // end namespace KernelArg + + namespace Kernel { + + // In-memory representation of kernel information. + struct Metadata { + std::string Name; + std::string Language; + std::vector LanguageVersion; + std::vector ReqdWorkGroupSize; + std::vector WorkGroupSizeHint; + std::string VecTypeHint; + uint32_t KernelIndex = INVALID_KERNEL_INDEX; + uint8_t NoPartialWorkGroups = 0; + std::vector Args; + + Metadata() = default; + }; + + } // end namespace Kernel + + namespace IsaInfo { + + /// \brief In-memory representation of instruction set architecture + /// information. + struct Metadata { + /// \brief Wavefront size. + unsigned WavefrontSize = 0; + /// \brief Local memory size in bytes. + unsigned LocalMemorySize = 0; + /// \brief Number of execution units per compute unit. + unsigned EUsPerCU = 0; + /// \brief Maximum number of waves per execution unit. + unsigned MaxWavesPerEU = 0; + /// \brief Maximum flat work group size. + unsigned MaxFlatWorkGroupSize = 0; + /// \brief SGPR allocation granularity. + unsigned SGPRAllocGranule = 0; + /// \brief Total number of SGPRs. + unsigned TotalNumSGPRs = 0; + /// \brief Addressable number of SGPRs. + unsigned AddressableNumSGPRs = 0; + /// \brief VGPR allocation granularity. + unsigned VGPRAllocGranule = 0; + /// \brief Total number of VGPRs. + unsigned TotalNumVGPRs = 0; + /// \brief Addressable number of VGPRs. + unsigned AddressableNumVGPRs = 0; + + Metadata() = default; + }; + + } // end namespace IsaInfo + + namespace Program { + + // In-memory representation of program information. + struct Metadata { + std::vector MDVersionSeq; + IsaInfo::Metadata IsaInfo; + std::vector PrintfInfo; + std::vector Kernels; + + explicit Metadata() = default; + + // Construct from an YAML string. + explicit Metadata(const std::string &YAML); + + // Convert to YAML string. + std::string toYAML(); + + // Convert from YAML string. + static Metadata fromYAML(const std::string &S); + }; + + } //end namespace Program + +} // end namespace RuntimeMD +} // end namespace AMDGPU + +#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPURUNTIMEMETADATA_H diff --git a/projects/clr/hipamd/src/device_functions.cpp b/projects/clr/hipamd/src/device_functions.cpp index abc9db570e..10d8d3ab89 100644 --- a/projects/clr/hipamd/src/device_functions.cpp +++ b/projects/clr/hipamd/src/device_functions.cpp @@ -41,8 +41,8 @@ struct holder32Bit { }; } __attribute__((aligned(4))); -struct holder64Bit hold64; -struct holder32Bit hold32; +__device__ struct holder64Bit hold64; +__device__ struct holder32Bit hold32; __device__ float __double2float_rd(double x) { diff --git a/projects/clr/hipamd/src/device_util.cpp b/projects/clr/hipamd/src/device_util.cpp index 4b0e7efefd..8ce53765b5 100644 --- a/projects/clr/hipamd/src/device_util.cpp +++ b/projects/clr/hipamd/src/device_util.cpp @@ -34,8 +34,8 @@ THE SOFTWARE. This is the best place to put them because the device global variables need to be initialized at the start. */ -__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP]; -__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES]; +__device__ ADDRESS_SPACE_1 char gpuHeap[SIZE_OF_HEAP]; +__device__ ADDRESS_SPACE_1 uint32_t gpuFlags[NUM_PAGES]; __device__ void *__hip_hc_malloc(size_t size) { @@ -835,10 +835,6 @@ __device__ float __hip_ynf(int n, float x) return b0; } - - -const int warpSize = 64; - __device__ long long int clock64() { return (long long int)hc::__cycle_u64(); }; __device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); }; @@ -1083,7 +1079,7 @@ __host__ __device__ int max(int arg1, int arg2) return (int)(hc::precise_math::fmax((float)arg1, (float)arg2)); } -__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr() +__device__ ADDRESS_SPACE_3 void* __get_dynamicgroupbaseptr() { return hc::get_dynamic_group_segment_base_pointer(); } diff --git a/projects/clr/hipamd/src/env.cpp b/projects/clr/hipamd/src/env.cpp index 90d00feac0..0899784592 100644 --- a/projects/clr/hipamd/src/env.cpp +++ b/projects/clr/hipamd/src/env.cpp @@ -1,4 +1,26 @@ -#include "hip_hcc.h" +/* +Copyright (c) 2015 - present 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_hcc_internal.h" #include "trace_helper.h" #include "env.h" diff --git a/projects/clr/hipamd/src/grid_launch.cpp b/projects/clr/hipamd/src/grid_launch.cpp new file mode 100644 index 0000000000..7739995600 --- /dev/null +++ b/projects/clr/hipamd/src/grid_launch.cpp @@ -0,0 +1,60 @@ +/* +Copyright (c) 2015 - present 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/hcc_detail/grid_launch_GGL.hpp" + +// Internal header, do not percolate upwards. +#include "hip_hcc_internal.h" +#include "hc.hpp" +#include "trace_helper.h" + +namespace hip_impl +{ + hc::accelerator_view lock_stream_hip_( + hipStream_t& stream, void*& locked_stream) + { // This allocated but does not take ownership of locked_stream. If it is + // not deleted elsewhere it will leak. + using L = decltype(stream->lockopen_preKernelCommand()); + + HIP_INIT(); + + stream = ihipSyncAndResolveStream(stream); + locked_stream = new L{stream->lockopen_preKernelCommand()}; + return (*static_cast(locked_stream))->_av; + } + + void unlock_stream_hip_( + hipStream_t stream, + void* locked_stream, + const char* kernel_name, + hc::accelerator_view* acc_v) + { // Precondition: acc_v is the accelerator_view associated with stream + // which is guarded by locked_stream; + // locked_stream is deletable. + using L = decltype(stream->lockopen_preKernelCommand()); + + stream->lockclose_postKernelCommand(kernel_name, acc_v); + + delete static_cast(locked_stream); + locked_stream = nullptr; + } +} diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index 835cf6e795..11ef6d6da5 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -26,7 +26,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" // Stack of contexts @@ -283,3 +283,79 @@ hipError_t hipCtxGetFlags ( unsigned int* flags ) *flags = tempCtx->_ctxFlags; return ihipLogStatus(e); } + +hipError_t hipDevicePrimaryCtxGetState ( hipDevice_t dev, unsigned int* flags, int* active ) +{ + HIP_INIT_API(dev, flags, active); + hipError_t e = hipSuccess; + auto deviceHandle = ihipGetDevice(dev); + + if (deviceHandle == NULL) { + e = hipErrorInvalidDevice; + } + + ihipCtx_t* tempCtx; + tempCtx = ihipGetTlsDefaultCtx(); + ihipCtx_t* primaryCtx = deviceHandle->_primaryCtx; + if(tempCtx == primaryCtx) { + *active = 1; + *flags = tempCtx->_ctxFlags; + } else { + *active = 0; + *flags = primaryCtx->_ctxFlags; + } + return ihipLogStatus(e); +} + +hipError_t hipDevicePrimaryCtxRelease ( hipDevice_t dev) +{ + HIP_INIT_API(dev); + hipError_t e = hipSuccess; + auto deviceHandle = ihipGetDevice(dev); + + if (deviceHandle == NULL) { + e = hipErrorInvalidDevice; + } + return ihipLogStatus(e); +} + +hipError_t hipDevicePrimaryCtxRetain ( hipCtx_t* pctx, hipDevice_t dev ) +{ + HIP_INIT_API(pctx, dev); + hipError_t e = hipSuccess; + auto deviceHandle = ihipGetDevice(dev); + + if (deviceHandle == NULL) { + e = hipErrorInvalidDevice; + } + *pctx = deviceHandle->_primaryCtx; + return ihipLogStatus(e); +} + +hipError_t hipDevicePrimaryCtxReset ( hipDevice_t dev ) +{ + HIP_INIT_API(dev); + hipError_t e = hipSuccess; + auto deviceHandle = ihipGetDevice(dev); + + if (deviceHandle == NULL) { + e = hipErrorInvalidDevice; + } + ihipCtx_t* primaryCtx = deviceHandle->_primaryCtx; + primaryCtx->locked_reset(); + return ihipLogStatus(e); +} + +hipError_t hipDevicePrimaryCtxSetFlags ( hipDevice_t dev, unsigned int flags ) +{ + HIP_INIT_API(dev, flags); + hipError_t e = hipSuccess; + auto deviceHandle = ihipGetDevice(dev); + + if (deviceHandle == NULL) { + e = hipErrorInvalidDevice; + } else { + e = hipErrorContextAlreadyInUse; + } + return ihipLogStatus(e); +} diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index d3c68e6fdf..88d94411e8 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" #include "device_util.h" diff --git a/projects/clr/hipamd/src/hip_error.cpp b/projects/clr/hipamd/src/hip_error.cpp index 4c14ba4156..21d5b6aa85 100644 --- a/projects/clr/hipamd/src/hip_error.cpp +++ b/projects/clr/hipamd/src/hip_error.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" //------------------------------------------------------------------------------------------------- diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 5a0ed9d8f8..d44f201db5 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" //------------------------------------------------------------------------------------------------- diff --git a/projects/clr/hipamd/src/hip_fp16.cpp b/projects/clr/hipamd/src/hip_fp16.cpp index b306a9d3de..c2b7b47597 100644 --- a/projects/clr/hipamd/src/hip_fp16.cpp +++ b/projects/clr/hipamd/src/hip_fp16.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include"hip/hip_fp16.h" +#include"hip/hcc_detail/hip_fp16.h" struct hipHalfHolder{ union { @@ -31,7 +31,7 @@ struct hipHalfHolder{ #define HINF 65504 -static struct hipHalfHolder __hInfValue = {HINF}; +__device__ static struct hipHalfHolder __hInfValue = {HINF}; __device__ __half __hadd(__half a, __half b) { return a + b; diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 760f46046a..374840f91f 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -43,7 +43,7 @@ THE SOFTWARE. #include "hsa/hsa_ext_amd.h" #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" #include "env.h" @@ -762,7 +762,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_PRODUCT_NAME, &(prop->name)); char archName[256]; err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_NAME, &archName); - + if(strcmp(archName,"gfx701")==0){ prop->gcnArch = 701; } @@ -1556,7 +1556,6 @@ void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launc MARKER_END(); } - //================================================================================================= // HIP API Implementation // @@ -1806,7 +1805,7 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind, void printPointerInfo(unsigned dbFlag, const char *tag, const void *ptr, const hc::AmPointerInfo &ptrInfo) { tprintf (dbFlag, " %s=%p baseHost=%p baseDev=%p sz=%zu home_dev=%d tracked=%d isDevMem=%d registered=%d\n", - tag, ptr, + tag, ptr, ptrInfo._hostPointer, ptrInfo._devicePointer, ptrInfo._sizeBytes, ptrInfo._appId, ptrInfo._sizeBytes != 0, ptrInfo._isInDeviceMem, !ptrInfo._isAmManaged); } @@ -1873,31 +1872,47 @@ void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t hc::am_memtracker_add(ptr, ptrInfo); } -void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind) { if(kind == hipMemcpyHostToHost){ - acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost); + acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset, Kalmar::hcMemcpyHostToHost); } if(kind == hipMemcpyHostToDevice){ - acc.memcpy_symbol(dst, (void*)src, sizeBytes); + acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset); } if(kind == hipMemcpyDeviceToDevice){ - acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToDevice); + acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset, Kalmar::hcMemcpyDeviceToDevice); } if(kind == hipMemcpyDeviceToHost){ - acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToHost); + acc.memcpy_symbol((void*) src, (void*)dst, sizeBytes, offset, Kalmar::hcMemcpyDeviceToHost); } } -void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind) { if(kind == hipMemcpyHostToDevice) { - addSymbolPtrToTracker(acc, dst, sizeBytes); - locked_getAv()->copy_async((void*)src, dst, sizeBytes); + hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); + bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); + if(srcTracked) { + addSymbolPtrToTracker(acc, dst, sizeBytes); + locked_getAv()->copy_async((void*)src, dst, sizeBytes); + } else { + LockedAccessor_StreamCrit_t crit(_criticalData); + this->wait(crit); + acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset); + } } if(kind == hipMemcpyDeviceToHost) { - addSymbolPtrToTracker(acc, src, sizeBytes); - locked_getAv()->copy_async((void*)src, dst, sizeBytes); + hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); + bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); + if(dstTracked) { + addSymbolPtrToTracker(acc, src, sizeBytes); + locked_getAv()->copy_async((void*)src, dst, sizeBytes); + } else { + LockedAccessor_StreamCrit_t crit(_criticalData); + this->wait(crit); + acc.memcpy_symbol((void*)src, (void*)dst, sizeBytes, offset, Kalmar::hcMemcpyDeviceToHost); + } } } diff --git a/projects/clr/hipamd/src/hip_hcc.h b/projects/clr/hipamd/src/hip_hcc_internal.h similarity index 99% rename from projects/clr/hipamd/src/hip_hcc.h rename to projects/clr/hipamd/src/hip_hcc_internal.h index b23aead072..4b960e2820 100644 --- a/projects/clr/hipamd/src/hip_hcc.h +++ b/projects/clr/hipamd/src/hip_hcc_internal.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_HCC_H -#define HIP_HCC_H +#ifndef HIP_SRC_HIP_HCC_INTERNAL_H +#define HIP_SRC_HIP_HCC_INTERNAL_H #include #include @@ -503,8 +503,8 @@ public: void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true); void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); - void lockedSymbolCopySync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind); - void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind); + void lockedSymbolCopySync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, size_t offset, unsigned kind); + void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, size_t offset, unsigned kind); //--- // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex. diff --git a/projects/clr/hipamd/src/hip_ldg.cpp b/projects/clr/hipamd/src/hip_ldg.cpp index d91f54a807..549d3ae085 100644 --- a/projects/clr/hipamd/src/hip_ldg.cpp +++ b/projects/clr/hipamd/src/hip_ldg.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index a92d11b847..805fc9efc0 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -25,7 +25,7 @@ THE SOFTWARE. #include "hsa/hsa_ext_amd.h" #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" #include "hip/hcc_detail/hip_texture.h" #include @@ -261,11 +261,11 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) auto device = ctx->getWriteableDevice(); unsigned amFlags = HIP_COHERENT_HOST_ALLOC ? amHostCoherent : amHostPinned; - *ptr = hip_internal::allocAndSharePtr(HIP_COHERENT_HOST_ALLOC ? "finegrained_host":"pinned_host", + *ptr = hip_internal::allocAndSharePtr(HIP_COHERENT_HOST_ALLOC ? "finegrained_host":"pinned_host", sizeBytes, ctx, amFlags, flags); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; - } + } } } @@ -314,7 +314,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; - } + } } else { hip_status = hipErrorMemoryAllocation; } @@ -372,7 +372,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, am_flags, 0); if (size && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; - } + } } else { hip_status = hipErrorMemoryAllocation; @@ -472,7 +472,7 @@ hipError_t hipHostUnregister(void *hostPtr) return ihipLogStatus(hip_status); } -hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind) +hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind) { HIP_INIT_CMD_API(symbolName, src, count, offset, kind); @@ -485,7 +485,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou hc::accelerator acc = ctx->getDevice()->_acc; - void *dst = acc.get_symbol_address(symbolName); + void *dst = acc.get_symbol_address((const char*) symbolName); tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst); if(dst == nullptr) @@ -497,7 +497,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) { - stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind); + stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind); // acc.memcpy_symbol(dst, (void*)src, count+offset); } else { return ihipLogStatus(hipErrorInvalidValue); @@ -507,7 +507,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou } -hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind) +hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, size_t offset, hipMemcpyKind kind) { HIP_INIT_CMD_API(symbolName, dst, count, offset, kind); @@ -520,7 +520,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count, hc::accelerator acc = ctx->getDevice()->_acc; - void *src = acc.get_symbol_address(symbolName); + void *src = acc.get_symbol_address((const char*) symbolName); tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst); if(dst == nullptr) @@ -532,7 +532,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count, if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) { - stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind); + stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind); } else { return ihipLogStatus(hipErrorInvalidValue); @@ -542,7 +542,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count, } -hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) +hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_CMD_API(symbolName, src, count, offset, kind, stream); @@ -557,7 +557,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ hc::accelerator acc = ctx->getDevice()->_acc; - void *dst = acc.get_symbol_address(symbolName); + void *dst = acc.get_symbol_address((const char*) symbolName); tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst); if(dst == nullptr) @@ -567,7 +567,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ if (stream) { try { - stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count + offset, kind); + stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count, offset, kind); } catch (ihipException ex) { e = ex._code; @@ -580,7 +580,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ } -hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) +hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_CMD_API(symbolName, dst, count, offset, kind, stream); @@ -595,7 +595,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t co hc::accelerator acc = ctx->getDevice()->_acc; - void *src = acc.get_symbol_address(symbolName); + void *src = acc.get_symbol_address((const char*) symbolName); tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, src); if(src == nullptr || dst == nullptr) @@ -603,9 +603,10 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t co return ihipLogStatus(hipErrorInvalidSymbol); } + stream = ihipSyncAndResolveStream(stream); if (stream) { try { - stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind); + stream->lockedSymbolCopyAsync(acc, dst, src, count, offset, kind); } catch (ihipException ex) { e = ex._code; diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 1f20a47c13..67bba5f935 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -23,18 +23,45 @@ THE SOFTWARE. #include #include #include +#include #include +#include +#include #include "hsa/hsa.h" #include "hsa/hsa_ext_amd.h" #include "hsa/amd_hsa_kernel_code.h" #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" //TODO Use Pool APIs from HCC to get memory regions. +#include +inline uint64_t alignTo(uint64_t Value, uint64_t Align, uint64_t Skew = 0) { + assert(Align != 0u && "Align can't be 0."); + Skew %= Align; + return (Value + Align - 1 - Skew) / Align * Align + Skew; +} + +struct ihipKernArgInfo{ + std::vector Size; + std::vector Align; + std::vector ArgType; + std::vector ArgName; + uint32_t totalSize; +}; + +std::map kernelArguments; + +struct MyElfNote { + uint32_t n_namesz = 0; + uint32_t n_descsz = 0; + uint32_t n_type = 0; + + MyElfNote() = default; +}; struct ihipModuleSymbol_t{ uint64_t _object; // The kernel object. @@ -172,9 +199,57 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ (*module)->ptr = p; (*module)->size = size; in.seekg(0, std::ios::beg); - std::copy(std::istreambuf_iterator(in), - std::istreambuf_iterator(), ptr); + std::copy(std::istreambuf_iterator(in), std::istreambuf_iterator(), ptr); +/* + Elf *e = elf_memory((char*)p, size); + if(elf_kind(e) != ELF_K_ELF){ + return ihipLogStatus(hipErrorInvalidValue); + } + size_t numpHdrs; + if(elf_getphdrnum(e, &numpHdrs) != 0){ + return ihipLogStatus(hipErrorInvalidValue); + } + for(size_t i=0;i= sizeof(int)){ + char* ptr = (char*) p + pHdr.p_offset; + char* segmentEnd = ptr + pHdr.p_filesz; + while(ptr < segmentEnd){ + MyElfNote *note = (MyElfNote*) ptr; + char *name = (char*) ¬e[1]; + char *desc = name + alignTo(note->n_namesz, sizeof(int)); + if (note->n_type == 8) { + std::string metadatastr((const char*)desc, (size_t)note->n_descsz); + AMDGPU::RuntimeMD::Program::Metadata meta; + meta = meta.fromYAML(metadatastr); + for(int i=0;in_namesz, sizeof(int)) + + alignTo(note->n_descsz, sizeof(int)); + } + } + } +*/ status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object); if(status != HSA_STATUS_SUCCESS){ @@ -288,16 +363,12 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, } -hipError_t hipModuleLaunchKernel(hipFunction_t f, - uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, - uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, - uint32_t sharedMemBytes, hipStream_t hStream, +hipError_t ihipModuleLaunchKernel(hipFunction_t f, + uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, + size_t sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra) { - HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, - blockDimX, blockDimY, blockDimZ, - sharedMemBytes, hStream, - kernelParams, extra); auto ctx = ihipGetTlsDefaultCtx(); hipError_t ret = hipSuccess; @@ -313,13 +384,26 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, void *config[5] = {0}; size_t kernArgSize; - if(extra != NULL){ + if(kernelParams != NULL){ + std::string name = f->_name; + struct ihipKernArgInfo pl = kernelArguments[name]; + char* argBuf = (char*)malloc(pl.totalSize); + memset(argBuf, 0, pl.totalSize); + int index = 0; + for(int i=0;i_name.c_str()); + hStream = ihipPreLaunchKernel(hStream, dim3(globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ), dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str()); hsa_kernel_dispatch_packet_t aql; @@ -341,12 +425,12 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, //aql.completion_signal._handle = 0; //aql.kernarg_address = 0; - aql.workgroup_size_x = blockDimX; - aql.workgroup_size_y = blockDimY; - aql.workgroup_size_z = blockDimZ; - aql.grid_size_x = blockDimX * gridDimX; - aql.grid_size_y = blockDimY * gridDimY; - aql.grid_size_z = blockDimZ * gridDimZ; + aql.workgroup_size_x = localWorkSizeX; + aql.workgroup_size_y = localWorkSizeY; + aql.workgroup_size_z = localWorkSizeZ; + aql.grid_size_x = globalWorkSizeX; + aql.grid_size_y = globalWorkSizeY; + aql.grid_size_z = globalWorkSizeZ; aql.group_segment_size = f->_groupSegmentSize + sharedMemBytes; aql.private_segment_size = f->_privateSegmentSize; aql.kernel_object = f->_object; @@ -364,13 +448,46 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, nullptr/*completion_future*/); - + if(kernelParams != NULL){ + free(config[1]); + } ihipPostLaunchKernel(f->_name.c_str(), hStream, lp); } - return ihipLogStatus(ret); + return ret; } +hipError_t hipModuleLaunchKernel(hipFunction_t f, + uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, + uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, + uint32_t sharedMemBytes, hipStream_t hStream, + void **kernelParams, void **extra) +{ + HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, + blockDimX, blockDimY, blockDimZ, + sharedMemBytes, hStream, + kernelParams, extra); + return ihipLogStatus(ihipModuleLaunchKernel(f, + blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, + blockDimX, blockDimY, blockDimZ, + sharedMemBytes, hStream, kernelParams, extra)); +} + + +hipError_t hipHccModuleLaunchKernel(hipFunction_t f, + uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, + size_t sharedMemBytes, hipStream_t hStream, + void **kernelParams, void **extra) +{ + HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, + localWorkSizeX, localWorkSizeY, localWorkSizeZ, + sharedMemBytes, hStream, + kernelParams, extra); + return ihipLogStatus(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, + localWorkSizeX, localWorkSizeY, localWorkSizeZ, + sharedMemBytes, hStream, kernelParams, extra)); +} hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char* name) diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index e57665be0c..984110a6b5 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 594fb6e860..d7f8717725 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip_hcc.h" +#include "hip_hcc_internal.h" #include "trace_helper.h" diff --git a/projects/clr/hipamd/src/hip_util.h b/projects/clr/hipamd/src/hip_util.h index f6817ffccb..8c4d19bb40 100644 --- a/projects/clr/hipamd/src/hip_util.h +++ b/projects/clr/hipamd/src/hip_util.h @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - present 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 @@ -20,8 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#ifndef HIP_UTIL_H -#define HIP_UTIL_H +#ifndef HIP_INCLUDE_HCC_DETAIL_HIP_UTIL_H +#define HIP_INCLUDE_HCC_DETAIL_HIP_UTIL_H #include #include diff --git a/projects/clr/hipamd/src/math_functions.cpp b/projects/clr/hipamd/src/math_functions.cpp index 230eb2aacc..92cc8689fc 100644 --- a/projects/clr/hipamd/src/math_functions.cpp +++ b/projects/clr/hipamd/src/math_functions.cpp @@ -188,8 +188,12 @@ __device__ float ldexpf(float x, int exp) } __device__ float lgammaf(float x) { - int sign; - return hc::precise_math::lgammaf(x, &sign); + float val = 0.0f; + float y = x - 1; + while(y > 0){ + val += logf(y--); + } + return val; } __device__ long long int llrintf(float x) { @@ -202,7 +206,8 @@ __device__ long long int llroundf(float x) int y = hc::precise_math::roundf(x); long long int z = y; return z; -}__device__ float log10f(float x) +} +__device__ float log10f(float x) { return hc::precise_math::log10f(x); } @@ -569,8 +574,12 @@ __device__ double ldexp(double x, int exp) } __device__ double lgamma(double x) { - int sign; - return hc::precise_math::lgamma(x, &sign); + double val = 0.0; + double y = x - 1; + while(y > 0){ + val += log(y--); + } + return val; } __device__ long long int llrint(double x) { diff --git a/projects/clr/hipamd/src/trace_helper.h b/projects/clr/hipamd/src/trace_helper.h index 3bf2857c3a..d49cb67be0 100644 --- a/projects/clr/hipamd/src/trace_helper.h +++ b/projects/clr/hipamd/src/trace_helper.h @@ -28,7 +28,6 @@ THE SOFTWARE. #include #include #include - //--- // Helper functions to convert HIP function arguments into strings. // Handles POD data types as well as enumerations (ie hipMemcpyKind). @@ -71,7 +70,6 @@ inline std::string ToString(hipEvent_t v) ss << v; return ss.str(); }; - // hipStream_t template <> inline std::string ToString(hipStream_t v) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index 6bf13a0809..27af03ced3 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp @@ -59,6 +59,8 @@ __device__ void integer_intrinsics() __umulhi((unsigned int)1, (unsigned int)2); __urhadd((unsigned int)1, (unsigned int)2); __usad((unsigned int)1, (unsigned int)2, 0); + + assert(1); } __global__ void compileIntegerIntrinsics(hipLaunchParm lp, int ignored) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 9e188e9f17..c2ffb5ce7d 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -31,9 +31,10 @@ THE SOFTWARE. #define NUM 1024 #define SIZE 1024*4 +// TODO - collapse: #ifdef __HIP_PLATFORM_HCC__ -__attribute__((address_space(1))) int globalIn[NUM]; -__attribute__((address_space(1))) int globalOut[NUM]; +__device__ ADDRESS_SPACE_1 int globalIn[NUM]; +__device__ ADDRESS_SPACE_1 int globalOut[NUM]; #endif #ifdef __HIP_PLATFORM_NVCC__ @@ -92,8 +93,29 @@ int main() hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost); for(unsigned i=0;i correct{ptr, hipFree}; + hipLaunchKernel( + CheckVectorTypes, dim3(1,1,1), dim3(1,1,1), 0, 0, correct.get()); + bool passed = false; + if (hipMemcpyDtoH(&passed, correct.get(), sizeof(bool)) != HIP_SUCCESS) { + return EXIT_FAILURE; + } + + if (passed == true){ + std::cout << "PASSED" << std::endl; + return 0; + } + else + return EXIT_FAILURE; } + diff --git a/projects/clr/hipamd/tests/src/hipHcc.cpp b/projects/clr/hipamd/tests/src/hipHcc.cpp index 9357e5211a..92d9e3e88d 100644 --- a/projects/clr/hipamd/tests/src/hipHcc.cpp +++ b/projects/clr/hipamd/tests/src/hipHcc.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. #include #include #include "hip/hip_runtime.h" -#include "hip/hcc.h" +#include "hip/hip_hcc.h" #include "test_common.h" #define CHECK(error) \ @@ -61,4 +61,3 @@ int main(int argc, char *argv[]) passed(); }; -