Merge branch 'amd-develop' into amd-master
Change-Id: Id5a5d8417ff3de9d55bc42a4ee6c6b367e1ebeff
[ROCm/hip commit: 783159ad73]
Этот коммит содержится в:
Maneesh Gupta
@@ -217,13 +217,14 @@ endif()
|
||||
# Install .version
|
||||
install(FILES ${PROJECT_BINARY_DIR}/.version DESTINATION bin)
|
||||
|
||||
# Install src, bin, include if necessary
|
||||
# Install src, bin, include & cmake if necessary
|
||||
execute_process(COMMAND test ${CMAKE_INSTALL_PREFIX} -ef ${CMAKE_CURRENT_SOURCE_DIR}
|
||||
RESULT_VARIABLE INSTALL_SOURCE)
|
||||
if(NOT ${INSTALL_SOURCE} EQUAL 0)
|
||||
install(DIRECTORY src DESTINATION .)
|
||||
install(DIRECTORY bin DESTINATION . USE_SOURCE_PERMISSIONS)
|
||||
install(DIRECTORY include DESTINATION .)
|
||||
install(DIRECTORY cmake DESTINATION .)
|
||||
endif()
|
||||
|
||||
#############################
|
||||
|
||||
@@ -94,7 +94,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
|
||||
- Code Indentation:
|
||||
- Tabs should be expanded to spaces.
|
||||
- Use 4 spaces indendation.
|
||||
- Capitaliziation and Naming
|
||||
- Capitalization and Naming
|
||||
- Prefer camelCase for HIP interfaces and internal symbols. Note HCC uses _ for separator.
|
||||
This guideline is not yet consistently followed in HIP code - eventual compliance is aspirational.
|
||||
- Member variables should begin with a leading "_". This allows them to be easily distinguished from other variables or functions.
|
||||
@@ -110,6 +110,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
|
||||
doFooElse();
|
||||
}
|
||||
'''
|
||||
- namespace should be on same line as { and separated by a space.
|
||||
- Single-line if statement should still use {/} pair (even though C++ does not require).
|
||||
- Miscellaneous
|
||||
- All references in function parameter lists should be const.
|
||||
@@ -120,6 +121,8 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
|
||||
- HIP_INIT_API() should be placed at the start of each top-level HIP API. This function will make sure the HIP runtime is initialized,
|
||||
and also constructs an appropriate API string for tracing and CodeXL marker tracing. The arguments to HIP_INIT_API should match
|
||||
those of the parent fucntion.
|
||||
- ihipLogStatus should only be called from top-level HIP APIs,and should be called to log and return the error code. The error code
|
||||
is used by the GetLastError and PeekLastError functions - if a HIP API simply returns, then the error will not be logged correctly.
|
||||
|
||||
|
||||
#### Presubmit Testing:
|
||||
|
||||
@@ -168,6 +168,8 @@ my $hasCU = 0; # options contain a cu-style file (HCC must force recogni
|
||||
my $needHipHcc = ($HIP_PLATFORM eq 'hcc'); # set if we need to link hip_hcc.o from src tree. (some builds, ie cmake, provide their own)
|
||||
my $printHipVersion = 0; # print HIP version
|
||||
my $runCmd = 1;
|
||||
my $buildDeps = 0;
|
||||
|
||||
|
||||
my @options = ();
|
||||
my @inputs = ();
|
||||
@@ -181,7 +183,7 @@ my $ISACMD="";
|
||||
if($HIP_PLATFORM eq "hcc"){
|
||||
$ISACMD .= "$HIP_PATH/bin/hipgenisa.sh ";
|
||||
$ISACMD .= $ROCM_PATH;
|
||||
if($ARGV[0] eq "--gencodeobject"){
|
||||
if($ARGV[0] eq "--genco"){
|
||||
foreach $isaarg (@ARGV[1..$#ARGV]){
|
||||
$ISACMD .= " ";
|
||||
$ISACMD .= $isaarg;
|
||||
@@ -196,7 +198,7 @@ if($HIP_PLATFORM eq "hcc"){
|
||||
|
||||
if($HIP_PLATFORM eq "nvcc"){
|
||||
$ISACMD .= "$HIP_PATH/bin/hipcc -ptx ";
|
||||
if($ARGV[0] eq "--gencodeobject"){
|
||||
if($ARGV[0] eq "--genco"){
|
||||
foreach $isaarg (@ARGV[1..$#ARGV]){
|
||||
$ISACMD .= " ";
|
||||
$ISACMD .= $isaarg;
|
||||
@@ -240,6 +242,10 @@ foreach $arg (@ARGV)
|
||||
$printHipVersion = 1;
|
||||
$runCmd = 0;
|
||||
}
|
||||
if($arg eq '-M') {
|
||||
$compileOnly = 1;
|
||||
$buildDeps = 1;
|
||||
}
|
||||
|
||||
if ($arg =~ m/^-/) {
|
||||
# options start with -
|
||||
@@ -278,6 +284,9 @@ if ($hasC and $HIP_PLATFORM eq 'nvcc') {
|
||||
if ($hasCU and $HIP_PLATFORM eq 'hcc') {
|
||||
$HIPCXXFLAGS .= " -x c++";
|
||||
}
|
||||
if ($buildDeps and $HIP_PLATFORM eq 'nvcc') {
|
||||
$HIPCXXFLAGS .= " -M -D__CUDACC__";
|
||||
}
|
||||
|
||||
if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
|
||||
{
|
||||
|
||||
@@ -21,11 +21,22 @@ HIP_PATH="$( command cd -P "$( dirname "$SOURCE" )/.." && pwd )"
|
||||
|
||||
export KMDUMPISA=1
|
||||
export KMDUMPLLVM=1
|
||||
mkdir /tmp/hipgenisa
|
||||
$HIP_PATH/bin/hipcc $FILE_NAMES -o /tmp/hipgenisa/a.out
|
||||
mv dump.* /tmp/hipgenisa/
|
||||
$ROCM_PATH/hcc-lc/bin/llvm-mc -arch=amdgcn -mcpu=$TARGET -filetype=obj /tmp/hipgenisa/dump.isa -o /tmp/hipgenisa/dump.o
|
||||
$ROCM_PATH/llvm/bin/clang -target amdgcn--amdhsa /tmp/hipgenisa/dump.o -o $OUTPUT_FILE
|
||||
rm -r /tmp/hipgenisa
|
||||
export KMDUMPISA=0
|
||||
export KMDUMPLLVM=0
|
||||
hipgenisa_dir=`mktemp -d --tmpdir=/tmp hip.XXXXXXXX`;
|
||||
sed 's/extern \+"C" \+//g' $FILE_NAMES > $FILE_NAMES.kernel.tmp.cpp
|
||||
echo "
|
||||
int main(){}
|
||||
" >> $FILE_NAMES.kernel.tmp.cpp
|
||||
$HIP_PATH/bin/hipcc $FILE_NAMES.kernel.tmp.cpp -o $hipgenisa_dir/a.out
|
||||
mv dump.* $hipgenisa_dir
|
||||
$ROCM_PATH/hcc-lc/bin/llvm-mc -arch=amdgcn -mcpu=$TARGET -filetype=obj $hipgenisa_dir/dump.isa -o $hipgenisa_dir/dump.o
|
||||
$ROCM_PATH/llvm/bin/clang -target amdgcn--amdhsa $hipgenisa_dir/dump.o -o $hipgenisa_dir/dump.co
|
||||
map_sym=""
|
||||
kernels=$(objdump -t $hipgenisa_dir/dump.co | 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.co $OUTPUT_FILE
|
||||
rm $FILE_NAMES.kernel.tmp.cpp
|
||||
rm -r $hipgenisa_dir
|
||||
|
||||
@@ -117,30 +117,200 @@ struct cuda2hipMap {
|
||||
cuda2hipRename["cublas_v2.h"] = {"hipblas.h", CONV_INCLUDE, API_BLAS};
|
||||
|
||||
// Error codes and return types
|
||||
cuda2hipRename["CUresult"] = {"hipError_t", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["cudaError_t"] = {"hipError_t", CONV_TYPE, API_RUNTIME};
|
||||
cuda2hipRename["cudaError"] = {"hipError", CONV_TYPE, API_RUNTIME};
|
||||
cuda2hipRename["cudaSuccess"] = {"hipSuccess", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorMemoryAllocation"] = {"hipErrorMemoryAllocation", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorMemoryFree"] = {"hipErrorMemoryFree", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorUnknownSymbol"] = {"hipErrorUnknownSymbol", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorOutOfResources"] = {"hipErrorOutOfResources", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInvalidValue"] = {"hipErrorInvalidValue", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInvalidResourceHandle"] = {"hipErrorInvalidResourceHandle", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInvalidDevice"] = {"hipErrorInvalidDevice", CONV_ERR, API_RUNTIME};
|
||||
|
||||
// CUDA Driver API error code only
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_CONTEXT"] = {"hipErrorInvalidContext", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_CONTEXT_ALREADY_CURRENT"] = {"hipErrorContextAlreadyCurrent", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_MAP_FAILED"] = {"hipErrorMapFailed", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_UNMAP_FAILED"] = {"hipErrorUnmapFailed", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_ARRAY_IS_MAPPED"] = {"hipErrorArrayIsMapped", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_ALREADY_MAPPED"] = {"hipErrorAlreadyMapped", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_ALREADY_ACQUIRED"] = {"hipErrorAlreadyAcquired", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_NOT_MAPPED"] = {"hipErrorNotMapped", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_NOT_MAPPED_AS_ARRAY"] = {"hipErrorNotMappedAsArray", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_NOT_MAPPED_AS_POINTER"] = {"hipErrorNotMappedAsPointer", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_CONTEXT_ALREADY_IN_USE"] = {"hipErrorContextAlreadyInUse", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_SOURCE"] = {"hipErrorInvalidSource", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_FILE_NOT_FOUND"] = {"hipErrorFileNotFound", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["CUDA_ERROR_NOT_FOUND"] = {"hipErrorNotFound", CONV_ERR, API_DRIVER};
|
||||
|
||||
// CUDA RT API error code only
|
||||
cuda2hipRename["cudaErrorInvalidDeviceFunction"] = {"hipErrorInvalidDeviceFunction", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInvalidConfiguration"] = {"hipErrorInvalidConfiguration", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorPriorLaunchFailure"] = {"hipErrorPriorLaunchFailure", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInvalidMemcpyDirection"] = {"hipErrorInvalidMemcpyDirection", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInvalidDevicePointer"] = {"hipErrorInvalidDevicePointer", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorInitializationError"] = {"hipErrorInvalidDevicePointer", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaErrorMissingConfiguration"] = {"hipErrorMissingConfiguration", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_SUCCESS"] = {"hipSuccess", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaSuccess"] = {"hipSuccess", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_UNKNOWN"] = {"hipErrorUnknown", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_NOT_INITIALIZED"] = {"hipErrorNotInitialized", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInitializationError"] = {"hipErrorNotInitialized", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_DEINITIALIZED"] = {"hipErrorDeinitialized", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorCudartUnloading"] = {"hipErrorDeinitialized", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_OUT_OF_MEMORY"] = {"hipErrorMemoryAllocation", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorMemoryAllocation"] = {"hipErrorMemoryAllocation", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_HANDLE"] = {"hipErrorInvalidResourceHandle", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInvalidResourceHandle"] = {"hipErrorInvalidResourceHandle", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_VALUE"] = {"hipErrorInvalidValue", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInvalidValue"] = {"hipErrorInvalidValue", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_DEVICE"] = {"hipErrorInvalidDevice", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInvalidDevice"] = {"hipErrorInvalidDevice", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_NOT_INITIALIZED"] = {"hipErrorInitializationError", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInitializationError"] = {"hipErrorInitializationError", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_NO_DEVICE"] = {"hipErrorNoDevice", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorNoDevice"] = {"hipErrorNoDevice", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_NOT_READY"] = {"hipErrorNotReady", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorNotReady"] = {"hipErrorNotReady", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PEER_ACCESS_NOT_ENABLED"] = {"hipErrorPeerAccessNotEnabled", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorPeerAccessNotEnabled"] = {"hipErrorPeerAccessNotEnabled", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED"] = {"hipErrorPeerAccessAlreadyEnabled", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorPeerAccessAlreadyEnabled"] = {"hipErrorPeerAccessAlreadyEnabled", CONV_ERR, API_RUNTIME};
|
||||
// NOTE: no corresponding error type in CUDA
|
||||
//cuda2hipRename["cudaErrorRuntimeMemory"] = {"hipErrorRuntimeMemory", CONV_ERR, API_RUNTIME};
|
||||
//cuda2hipRename["cudaErrorRuntimeOther"] = {"hipErrorRuntimeOther", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PEER_ACCESS_UNSUPPORTED"] = {"hipErrorPeerAccessUnsupported", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorPeerAccessUnsupported"] = {"hipErrorPeerAccessUnsupported", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_PTX"] = {"hipErrorInvalidKernelFile", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInvalidPtx"] = {"hipErrorInvalidKernelFile", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_GRAPHICS_CONTEXT"] = {"hipErrorInvalidGraphicsContext", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInvalidGraphicsContext"] = {"hipErrorInvalidGraphicsContext", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND"] = {"hipErrorSharedObjectSymbolNotFound", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorSharedObjectSymbolNotFound"] = {"hipErrorSharedObjectSymbolNotFound", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_SHARED_OBJECT_INIT_FAILED"] = {"hipErrorSharedObjectInitFailed", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorSharedObjectInitFailed"] = {"hipErrorSharedObjectInitFailed", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_OPERATING_SYSTEM"] = {"hipErrorOperatingSystem", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorOperatingSystem"] = {"hipErrorOperatingSystem", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_ILLEGAL_ADDRESS"] = {"hipErrorIllegalAddress", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorIllegalAddress"] = {"hipErrorIllegalAddress", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_LAUNCH_FAILED"] = {"hipErrorLaunchFailure", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorLaunchFailure"] = {"hipErrorLaunchFailure", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_LAUNCH_TIMEOUT"] = {"hipErrorLaunchTimeOut", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorLaunchTimeout"] = {"hipErrorLaunchTimeOut", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES"] = {"hipErrorLaunchOutOfResources", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorLaunchOutOfResources"] = {"hipErrorLaunchOutOfResources", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_ECC_UNCORRECTABLE"] = {"hipErrorECCNotCorrectable", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorECCUncorrectable"] = {"hipErrorECCNotCorrectable", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED"] = {"hipErrorHostMemoryAlreadyRegistered", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorHostMemoryAlreadyRegistered"] = {"hipErrorHostMemoryAlreadyRegistered", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED"] = {"hipErrorHostMemoryNotRegistered", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorHostMemoryNotRegistered"] = {"hipErrorHostMemoryNotRegistered", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_NO_BINARY_FOR_GPU"] = {"hipErrorNoBinaryForGpu", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorNoKernelImageForDevice"] = {"hipErrorNoBinaryForGpu", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_UNSUPPORTED_LIMIT"] = {"hipErrorUnsupportedLimit", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorUnsupportedLimit"] = {"hipErrorUnsupportedLimit", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_INVALID_IMAGE"] = {"hipErrorInvalidImage", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorInvalidKernelImage"] = {"hipErrorInvalidImage", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PROFILER_DISABLED"] = {"hipErrorProfilerDisabled", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorProfilerDisabled"] = {"hipErrorProfilerDisabled", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PROFILER_NOT_INITIALIZED"] = {"hipErrorProfilerNotInitialized", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorProfilerNotInitialized"] = {"hipErrorProfilerNotInitialized", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PROFILER_ALREADY_STARTED"] = {"hipErrorProfilerAlreadyStarted", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorProfilerAlreadyStarted"] = {"hipErrorProfilerAlreadyStarted", CONV_ERR, API_RUNTIME};
|
||||
|
||||
cuda2hipRename["CUDA_ERROR_PROFILER_ALREADY_STOPPED"] = {"hipErrorProfilerAlreadyStopped", CONV_ERR, API_DRIVER};
|
||||
cuda2hipRename["cudaErrorProfilerAlreadyStopped"] = {"hipErrorProfilerAlreadyStopped", CONV_ERR, API_RUNTIME};
|
||||
|
||||
///////////////////////////// CUDA DRIVER API /////////////////////////////
|
||||
// Types
|
||||
// NOTE: CUdevice might be changed to typedef int in the future.
|
||||
cuda2hipRename["CUdevice"] = {"hipDevice_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CUdevice_attribute_enum"] = {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUdevice_attribute"] = {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK"] = {"hipDeviceAttributeMaxThreadsPerBlock", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X"] = {"hipDeviceAttributeMaxBlockDimX", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y"] = {"hipDeviceAttributeMaxBlockDimY", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z"] = {"hipDeviceAttributeMaxBlockDimZ", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X"] = {"hipDeviceAttributeMaxGridDimX", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y"] = {"hipDeviceAttributeMaxGridDimY", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z"] = {"hipDeviceAttributeMaxGridDimZ", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK"] = {"hipDeviceAttributeMaxSharedMemoryPerBlock", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY"] = {"hipDeviceAttributeTotalConstantMemory", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_WARP_SIZE"] = {"hipDeviceAttributeWarpSize", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK"] = {"hipDeviceAttributeMaxRegistersPerBlock", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_CLOCK_RATE"] = {"hipDeviceAttributeClockRate", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE"] = {"hipDeviceAttributeMemoryClockRate", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH"] = {"hipDeviceAttributeMemoryBusWidth", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH"] = {"hipDeviceAttributeMultiprocessorCount", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_COMPUTE_MODE"] = {"hipDeviceAttributeComputeMode", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE"] = {"hipDeviceAttributeL2CacheSize", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR"] = {"hipDeviceAttributeMaxThreadsPerMultiProcessor", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR"] = {"hipDeviceAttributeComputeCapabilityMajor", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR"] = {"hipDeviceAttributeComputeCapabilityMinor", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS"] = {"hipDeviceAttributeConcurrentKernels", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_PCI_BUS_ID"] = {"hipDeviceAttributePciBusId", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID"] = {"hipDeviceAttributePciDeviceId", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR"] = {"hipDeviceAttributeMaxSharedMemoryPerMultiprocessor", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD"] = {"hipDeviceAttributeIsMultiGpuBoard", CONV_DEV, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CUdevprop_st"] = {"hipDeviceProp_t", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUdevprop"] = {"hipDeviceProp_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
// TODO: Analogues enum is needed in HIP. Couldn't map enum to struct hipPointerAttribute_t.
|
||||
// TODO: Do for Pointer Attributes the same as for Device Attributes.
|
||||
// cuda2hipRename["CUpointer_attribute_enum"] = {"hipPointerAttribute_t", CONV_TYPE, API_DRIVER};
|
||||
// cuda2hipRename["CUpointer_attribute"] = {"hipPointerAttribute_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CUfunction"] = {"hipFunction_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
// unsupported yet by HIP
|
||||
// cuda2hipRename["CUfunction_attribute_enum"] = {"hipFuncAttribute_t", CONV_TYPE, API_DRIVER};
|
||||
// cuda2hipRename["CUfunction_attribute"] = {"hipFuncAttribute_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CUfunc_cache_enum"] = {"hipFuncCache", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUfunc_cache"] = {"hipFuncCache", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CU_FUNC_CACHE_PREFER_NONE"] = {"hipFuncCachePreferNone", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_FUNC_CACHE_PREFER_SHARED"] = {"hipFuncCachePreferShared", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_FUNC_CACHE_PREFER_L1"] = {"hipFuncCachePreferL1", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_FUNC_CACHE_PREFER_EQUAL"] = {"hipFuncCachePreferEqual", CONV_DEV, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CUsharedconfig_enum"] = {"hipSharedMemConfig", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUsharedconfig"] = {"hipSharedMemConfig", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE"] = {"hipSharedMemBankSizeDefault", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE"] = {"hipSharedMemBankSizeFourByte", CONV_DEV, API_DRIVER};
|
||||
cuda2hipRename["CU_SHARED_MEM_CONFIG_EIGHT_BYTE_BANK_SIZE"] = {"hipSharedMemBankSizeEightByte", CONV_DEV, API_DRIVER};
|
||||
|
||||
cuda2hipRename["CUcontext"] = {"hipCtx_t", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUmodule"] = {"hipModule_t", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUevent"] = {"hipEvent_t", CONV_TYPE, API_DRIVER};
|
||||
cuda2hipRename["CUstream"] = {"hipStream_t", CONV_TYPE, API_DRIVER};
|
||||
|
||||
/////////////////////////////// CUDA RT API ///////////////////////////////
|
||||
// Error API
|
||||
cuda2hipRename["cudaGetLastError"] = {"hipGetLastError", CONV_ERR, API_RUNTIME};
|
||||
cuda2hipRename["cudaPeekAtLastError"] = {"hipPeekAtLastError", CONV_ERR, API_RUNTIME};
|
||||
@@ -279,6 +449,7 @@ struct cuda2hipMap {
|
||||
cuda2hipRename["cudaDevAttrL2CacheSize"] = {"hipDeviceAttributeL2CacheSize", CONV_DEV, API_RUNTIME};
|
||||
cuda2hipRename["cudaDevAttrMaxThreadsPerMultiProcessor"] = {"hipDeviceAttributeMaxThreadsPerMultiProcessor", CONV_DEV, API_RUNTIME};
|
||||
cuda2hipRename["cudaDevAttrComputeCapabilityMajor"] = {"hipDeviceAttributeComputeCapabilityMajor", CONV_DEV, API_RUNTIME};
|
||||
cuda2hipRename["cudaDevAttrComputeCapabilityMinor"] = {"hipDeviceAttributeComputeCapabilityMinor", CONV_DEV, API_RUNTIME};
|
||||
cuda2hipRename["cudaDevAttrConcurrentKernels"] = {"hipDeviceAttributeConcurrentKernels", CONV_DEV, API_RUNTIME};
|
||||
cuda2hipRename["cudaDevAttrPciBusId"] = {"hipDeviceAttributePciBusId", CONV_DEV, API_RUNTIME};
|
||||
cuda2hipRename["cudaDevAttrPciDeviceId"] = {"hipDeviceAttributePciDeviceId", CONV_DEV, API_RUNTIME};
|
||||
@@ -1038,8 +1209,7 @@ static void processString(StringRef s, const cuda2hipMap &map,
|
||||
int64_t countReps[CONV_LAST],
|
||||
int64_t countApiReps[API_LAST]) {
|
||||
size_t begin = 0;
|
||||
while ((begin = s.find("cuda", begin)) != StringRef::npos ||
|
||||
(begin = s.find("cublas", begin)) != StringRef::npos) {
|
||||
while ((begin = s.find("cu", begin)) != StringRef::npos) {
|
||||
const size_t end = s.find_first_of(" ", begin + 4);
|
||||
StringRef name = s.slice(begin, end);
|
||||
const auto found = map.cuda2hipRename.find(name);
|
||||
@@ -1709,7 +1879,7 @@ static cl::opt<bool>
|
||||
|
||||
void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callback) {
|
||||
Finder.addMatcher(callExpr(isExpansionInMainFile(),
|
||||
callee(functionDecl(matchesName("cuda.*|cublas.*"))))
|
||||
callee(functionDecl(matchesName("cu.*"))))
|
||||
.bind("cudaCall"),
|
||||
Callback);
|
||||
Finder.addMatcher(cudaKernelCallExpr().bind("cudaLaunchKernel"), Callback);
|
||||
@@ -1720,7 +1890,7 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
|
||||
Callback);
|
||||
Finder.addMatcher(declRefExpr(isExpansionInMainFile(),
|
||||
to(enumConstantDecl(
|
||||
matchesName("cuda.*|cublas.*|CUDA.*|CUBLAS*"))))
|
||||
matchesName("cu.*|CU.*"))))
|
||||
.bind("cudaEnumConstantRef"),
|
||||
Callback);
|
||||
Finder.addMatcher(varDecl(isExpansionInMainFile(),
|
||||
@@ -1728,36 +1898,36 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
|
||||
.bind("cudaEnumConstantDecl"),
|
||||
Callback);
|
||||
Finder.addMatcher(varDecl(isExpansionInMainFile(),
|
||||
hasType(typedefDecl(matchesName("cuda.*|cublas.*"))))
|
||||
hasType(typedefDecl(matchesName("cu.*|CU.*"))))
|
||||
.bind("cudaTypedefVar"),
|
||||
Callback);
|
||||
// Array of elements of typedef type, Example: cudaStream_t streams[2];
|
||||
Finder.addMatcher(varDecl(isExpansionInMainFile(),
|
||||
hasType(arrayType(hasElementType(typedefType(
|
||||
hasDeclaration(typedefDecl(matchesName("cuda.*|cublas.*"))))))))
|
||||
hasDeclaration(typedefDecl(matchesName("cu.*|CU.*"))))))))
|
||||
.bind("cudaTypedefVar"),
|
||||
Callback);
|
||||
Finder.addMatcher(varDecl(isExpansionInMainFile(),
|
||||
hasType(cxxRecordDecl(matchesName("cuda.*|cublas.*"))))
|
||||
hasType(cxxRecordDecl(matchesName("cu.*|CU.*"))))
|
||||
.bind("cudaStructVar"),
|
||||
Callback);
|
||||
Finder.addMatcher(varDecl(isExpansionInMainFile(),
|
||||
hasType(pointsTo(cxxRecordDecl(
|
||||
matchesName("cuda.*|cublas.*")))))
|
||||
matchesName("cu.*|CU.*")))))
|
||||
.bind("cudaStructVarPtr"),
|
||||
Callback);
|
||||
Finder.addMatcher(parmVarDecl(isExpansionInMainFile(),
|
||||
hasType(namedDecl(matchesName("cuda.*|cublas.*"))))
|
||||
hasType(namedDecl(matchesName("cu.*|CU.*"))))
|
||||
.bind("cudaParamDecl"),
|
||||
Callback);
|
||||
Finder.addMatcher(parmVarDecl(isExpansionInMainFile(),
|
||||
hasType(pointsTo(namedDecl(
|
||||
matchesName("cuda.*|cublas.*")))))
|
||||
matchesName("cu.*|CU.*")))))
|
||||
.bind("cudaParamDeclPtr"),
|
||||
Callback);
|
||||
Finder.addMatcher(expr(isExpansionInMainFile(),
|
||||
sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(
|
||||
cxxRecordDecl(matchesName("cuda.*|cublas.*")))))))
|
||||
cxxRecordDecl(matchesName("cu.*|CU.*")))))))
|
||||
.bind("cudaStructSizeOf"),
|
||||
Callback);
|
||||
Finder.addMatcher(stringLiteral(isExpansionInMainFile()).bind("stringLiteral"),
|
||||
|
||||
@@ -1,12 +0,0 @@
|
||||
|
||||
|
||||
#Our FindHCC is in the same directory as the this CMakeLists.txt, which is CMAKE_SOURCE_DIR
|
||||
# Let cmake know where to find FindHCC.cmake
|
||||
#
|
||||
|
||||
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_SOURCE_DIR})
|
||||
|
||||
find_package(HIP REQUIRED)
|
||||
|
||||
message(STATUS "HIP_PLATFORM = ${HIP_PLATFORM}")
|
||||
|
||||
@@ -1,73 +0,0 @@
|
||||
|
||||
# findHCC does not currently address versioning, i.e.
|
||||
# a rich directory structure where version number is a subdirectory under root
|
||||
# Also, supported only on UNIX 64 bit systems.
|
||||
|
||||
if(UNIX)
|
||||
if(CMAKE_SIZEOF_VOID_P EQUAL 8)
|
||||
|
||||
find_library(HSA_LIBRARY
|
||||
NAMES hsa-runtime64
|
||||
PATHS
|
||||
ENV HSA_PATH
|
||||
/opt/rocm/hsa
|
||||
PATH_SUFFIXES
|
||||
lib)
|
||||
|
||||
if( NOT DEFINED ENV{HSA_PATH} )
|
||||
set( ENV{HSA_PATH} /opt/rocm/hsa)
|
||||
endif()
|
||||
|
||||
find_program(HCC
|
||||
NAMES hcc
|
||||
PATHS
|
||||
ENV HCC_PATH
|
||||
/opt/rocm/hcc
|
||||
PATH_SUFFIXES
|
||||
/bin)
|
||||
|
||||
if( NOT DEFINED ENV{HCC_PATH} )
|
||||
set( ENV{HCC_PATH} /opt/rocm/hcc)
|
||||
endif()
|
||||
|
||||
# this is now dynamic
|
||||
# find_library(AMP_LIBRARY
|
||||
# NAMES mcwamp
|
||||
# PATHS
|
||||
# ENV NCC_PATH
|
||||
# /opt/rocm/hcc
|
||||
# PATH_SUFFIXES
|
||||
# /lib)
|
||||
|
||||
find_path(HCC_INCLUDE_DIR
|
||||
NAMES
|
||||
hc.hpp
|
||||
PATHS
|
||||
ENV NCC_PATH
|
||||
/opt/rocm/hcc
|
||||
PATH_SUFFIXES
|
||||
/include)
|
||||
|
||||
|
||||
|
||||
set(HSA_LIBRARIES ${HSA_LIBRARY})
|
||||
#set(HCC_LIBRARIES ${AMP_LIBRARY})
|
||||
set(HCC_INCLUDE_DIRS ${HCC_INCLUDE_DIR})
|
||||
|
||||
include(FindPackageHandleStandardArgs)
|
||||
find_package_handle_standard_args(
|
||||
HCC
|
||||
FOUND_VAR HCC_FOUND
|
||||
REQUIRED_VARS HSA_LIBRARIES HCC_INCLUDE_DIRS HCC)
|
||||
|
||||
mark_as_advanced(
|
||||
HSA_LIBRARIES
|
||||
HCC_INCLUDE_DIRS
|
||||
)
|
||||
|
||||
else()
|
||||
message(SEND_ERROR "HCC is currently supported only on 64 bit UNIX platforms")
|
||||
endif()
|
||||
else()
|
||||
message(SEND_ERROR "HCC is currently supported on unix platforms")
|
||||
endif()
|
||||
@@ -1,56 +1,463 @@
|
||||
###############################################################################
|
||||
# FindHIP.cmake
|
||||
###############################################################################
|
||||
|
||||
find_package(CUDA)
|
||||
###############################################################################
|
||||
# SET: Variable defaults
|
||||
###############################################################################
|
||||
# User defined flags
|
||||
set(HIP_HIPCC_FLAGS "" CACHE STRING "Semicolon delimited flags for HIPCC")
|
||||
set(HIP_HCC_FLAGS "" CACHE STRING "Semicolon delimited flags for HCC")
|
||||
set(HIP_NVCC_FLAGS "" CACHE STRING "Semicolon delimted flags for NVCC")
|
||||
mark_as_advanced(HIP_HIPCC_FLAGS HIP_HCC_FLAGS HIP_NVCC_FLAGS)
|
||||
set(_hip_configuration_types ${CMAKE_CONFIGURATION_TYPES} ${CMAKE_BUILD_TYPE} Debug MinSizeRel Release RelWithDebInfo)
|
||||
list(REMOVE_DUPLICATES _hip_configuration_types)
|
||||
foreach(config ${_hip_configuration_types})
|
||||
string(TOUPPER ${config} config_upper)
|
||||
set(HIP_HIPCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for HIPCC")
|
||||
set(HIP_HCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for HCC")
|
||||
set(HIP_NVCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for NVCC")
|
||||
mark_as_advanced(HIP_HIPCC_FLAGS_${config_upper} HIP_HCC_FLAGS_${config_upper} HIP_NVCC_FLAGS_${config_upper})
|
||||
endforeach()
|
||||
option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON)
|
||||
mark_as_advanced(HIP_HOST_COMPILATION_CPP)
|
||||
|
||||
if( CUDA_INCLUDE_DIRS AND CUDA_VERSION AND CUDA_NVCC_EXECUTABLE)
|
||||
message(STATUS "CUDA_VERSION = ${CUDA_VERSION}")
|
||||
message(STATUS "CUDA_INCLUDE_DIRS = ${CUDA_INCLUDE_DIRS}")
|
||||
message(STATUS "CUDA_NVCC_EXECUTABLE = ${CUDA_NVCC_EXECUTABLE}")
|
||||
|
||||
set( HIP_PLATFORM "nvcc" )
|
||||
|
||||
#export the environment variable, so that HIPCC can find it.
|
||||
set(ENV{HIP_PLATFORM} nvcc)
|
||||
|
||||
else()
|
||||
find_package(HCC)
|
||||
|
||||
message(STATUS "HCC_FOUND = ${HCC_FOUND}")
|
||||
message(STATUS "HCC = ${HCC}")
|
||||
message(STATUS "HCC_INCLUDE_DIRS = ${HCC_INCLUDE_DIRS}")
|
||||
message(STATUS "HSA_LIBRARIES = ${HSA_LIBRARIES}")
|
||||
###############################################################################
|
||||
# FIND: HIP and associated helper binaries
|
||||
###############################################################################
|
||||
# HIP is supported on Linux only
|
||||
if(UNIX AND NOT APPLE AND NOT CYGWIN)
|
||||
# Search for HIP installation
|
||||
if(NOT HIP_ROOT_DIR)
|
||||
# Search in user specified path first
|
||||
find_path(
|
||||
HIP_ROOT_DIR
|
||||
NAMES hipconfig
|
||||
PATHS
|
||||
ENV ROCM_PATH
|
||||
ENV HIP_PATH
|
||||
PATH_SUFFIXES bin
|
||||
DOC "HIP installed location"
|
||||
NO_DEFAULT_PATH
|
||||
)
|
||||
# Now search in default path
|
||||
find_path(
|
||||
HIP_ROOT_DIR
|
||||
NAMES hipconfig
|
||||
PATHS
|
||||
/opt/rocm
|
||||
/opt/rocm/hip
|
||||
PATH_SUFFIXES bin
|
||||
DOC "HIP installed location"
|
||||
)
|
||||
|
||||
if( ${HCC_FOUND} STREQUAL "TRUE" )
|
||||
|
||||
# This directory is hip/cmake! HIP_PATH should be one directory up!
|
||||
set (HIP_PATH $ENV{HIP_PATH})
|
||||
if (NOT DEFINED HIP_PATH)
|
||||
set (HIP_PATH ${CMAKE_CURRENT_SOURCE_DIR}/..)
|
||||
set( ENV{HIP_PATH} ${HIP_PATH})
|
||||
endif()
|
||||
# Check if we found HIP installation
|
||||
if(HIP_ROOT_DIR)
|
||||
# If so, fix the path
|
||||
string(REGEX REPLACE "[/\\\\]?bin[64]*[/\\\\]?$" "" HIP_ROOT_DIR ${HIP_ROOT_DIR})
|
||||
# And push it back to the cache
|
||||
set(HIP_ROOT_DIR ${HIP_ROOT_DIR} CACHE PATH "HIP installed location" FORCE)
|
||||
endif()
|
||||
if(NOT EXISTS ${HIP_ROOT_DIR})
|
||||
if(HIP_FIND_REQUIRED)
|
||||
message(FATAL_ERROR "Specify HIP_ROOT_DIR")
|
||||
elseif(NOT HIP_FIND_QUIETLY)
|
||||
message("HIP_ROOT_DIR not found or specified")
|
||||
endif()
|
||||
endif()
|
||||
endif()
|
||||
|
||||
message(STATUS "ENV HIP_PATH = $ENV{HIP_PATH}")
|
||||
|
||||
find_program(HIPCC
|
||||
NAMES hipcc
|
||||
PATHS
|
||||
ENV HIP_PATH
|
||||
PATH_SUFFIXES
|
||||
/bin)
|
||||
|
||||
message(STATUS "HIPCC = ${HIPCC}")
|
||||
|
||||
if( DEFINED HIPCC)
|
||||
|
||||
set( HIP_PLATFORM "hcc" )
|
||||
#export the environment variable, so that HIPCC can find it.
|
||||
set(ENV{HIP_PLATFORM} "hcc")
|
||||
set (CMAKE_CXX_COMPILER ${HIPCC})
|
||||
|
||||
else()
|
||||
message(SEND_ERROR "Did not find HIPCC")
|
||||
endif()
|
||||
else()
|
||||
message(SEND_ERROR "hcc not found")
|
||||
endif()
|
||||
|
||||
# Find HIPCC executable
|
||||
find_program(
|
||||
HIP_HIPCC_EXECUTABLE
|
||||
NAMES hipcc
|
||||
PATHS
|
||||
"${HIP_ROOT_DIR}"
|
||||
ENV ROCM_PATH
|
||||
ENV HIP_PATH
|
||||
/opt/rocm
|
||||
/opt/rocm/hip
|
||||
PATH_SUFFIXES bin
|
||||
NO_DEFAULT_PATH
|
||||
)
|
||||
if(NOT HIP_HIPCC_EXECUTABLE)
|
||||
# Now search in default paths
|
||||
find_program(HIP_HIPCC_EXECUTABLE hipcc)
|
||||
endif()
|
||||
mark_as_advanced(HIP_HIPCC_EXECUTABLE)
|
||||
|
||||
# Find HIPCONFIG executable
|
||||
find_program(
|
||||
HIP_HIPCONFIG_EXECUTABLE
|
||||
NAMES hipconfig
|
||||
PATHS
|
||||
"${HIP_ROOT_DIR}"
|
||||
ENV ROCM_PATH
|
||||
ENV HIP_PATH
|
||||
/opt/rocm
|
||||
/opt/rocm/hip
|
||||
PATH_SUFFIXES bin
|
||||
NO_DEFAULT_PATH
|
||||
)
|
||||
if(NOT HIP_HIPCONFIG_EXECUTABLE)
|
||||
# Now search in default paths
|
||||
find_program(HIP_HIPCONFIG_EXECUTABLE hipconfig)
|
||||
endif()
|
||||
mark_as_advanced(HIP_HIPCONFIG_EXECUTABLE)
|
||||
|
||||
if(HIP_HIPCONFIG_EXECUTABLE AND NOT HIP_VERSION)
|
||||
# Compute the version
|
||||
execute_process(
|
||||
COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --version
|
||||
OUTPUT_VARIABLE _hip_version
|
||||
ERROR_VARIABLE _hip_error
|
||||
OUTPUT_STRIP_TRAILING_WHITESPACE
|
||||
ERROR_STRIP_TRAILING_WHITESPACE
|
||||
)
|
||||
if(NOT _hip_error)
|
||||
set(HIP_VERSION ${_hip_version} CACHE STRING "Version of HIP as computed from hipcc")
|
||||
else()
|
||||
set(HIP_VERSION "0.0.0" CACHE STRING "Version of HIP as computed by FindHIP()")
|
||||
endif()
|
||||
mark_as_advanced(HIP_VERSION)
|
||||
endif()
|
||||
if(HIP_VERSION)
|
||||
string(REPLACE "." ";" _hip_version_list "${HIP_VERSION}")
|
||||
list(GET _hip_version_list 0 HIP_VERSION_MAJOR)
|
||||
list(GET _hip_version_list 1 HIP_VERSION_MINOR)
|
||||
list(GET _hip_version_list 2 HIP_VERSION_PATCH)
|
||||
set(HIP_VERSION_STRING "${HIP_VERSION}")
|
||||
endif()
|
||||
|
||||
if(HIP_HIPCONFIG_EXECUTABLE AND NOT HIP_PLATFORM)
|
||||
# Compute the platform
|
||||
execute_process(
|
||||
COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --platform
|
||||
OUTPUT_VARIABLE _hip_platform
|
||||
OUTPUT_STRIP_TRAILING_WHITESPACE
|
||||
)
|
||||
set(HIP_PLATFORM ${_hip_platform} CACHE STRING "HIP platform as computed by hipconfig")
|
||||
mark_as_advanced(HIP_PLATFORM)
|
||||
endif()
|
||||
endif()
|
||||
|
||||
include(FindPackageHandleStandardArgs)
|
||||
find_package_handle_standard_args(
|
||||
HIP
|
||||
REQUIRED_VARS
|
||||
HIP_ROOT_DIR
|
||||
HIP_HIPCC_EXECUTABLE
|
||||
HIP_HIPCONFIG_EXECUTABLE
|
||||
HIP_PLATFORM
|
||||
VERSION_VAR HIP_VERSION
|
||||
)
|
||||
|
||||
###############################################################################
|
||||
# MACRO: Locate helper files
|
||||
###############################################################################
|
||||
macro(HIP_FIND_HELPER_FILE _name _extension)
|
||||
set(_hip_full_name "${_name}.${_extension}")
|
||||
get_filename_component(CMAKE_CURRENT_LIST_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH)
|
||||
set(HIP_${_name} "${CMAKE_CURRENT_LIST_DIR}/FindHIP/${_hip_full_name}")
|
||||
if(NOT EXISTS "${HIP_${_name}}")
|
||||
set(error_message "${_hip_full_name} not found in ${CMAKE_CURRENT_LIST_DIR}/FindHIP")
|
||||
if(HIP_FIND_REQUIRED)
|
||||
message(FATAL_ERROR "${error_message}")
|
||||
else()
|
||||
if(NOT HIP_FIND_QUIETLY)
|
||||
message(STATUS "${error_message}")
|
||||
endif()
|
||||
endif()
|
||||
endif()
|
||||
# Set this variable as internal, so the user isn't bugged with it.
|
||||
set(HIP_${_name} ${HIP_${_name}} CACHE INTERNAL "Location of ${_full_name}" FORCE)
|
||||
endmacro()
|
||||
|
||||
###############################################################################
|
||||
hip_find_helper_file(run_make2cmake cmake)
|
||||
hip_find_helper_file(run_hipcc cmake)
|
||||
###############################################################################
|
||||
|
||||
###############################################################################
|
||||
# MACRO: Seperate the options from the sources
|
||||
###############################################################################
|
||||
macro(HIP_GET_SOURCES_AND_OPTIONS _sources _cmake_options _hipcc_options _hcc_options _nvcc_options)
|
||||
set(${_sources})
|
||||
set(${_cmake_options})
|
||||
set(${_hipcc_options})
|
||||
set(${_hcc_options})
|
||||
set(${_nvcc_options})
|
||||
set(_hipcc_found_options FALSE)
|
||||
set(_hcc_found_options FALSE)
|
||||
set(_nvcc_found_options FALSE)
|
||||
foreach(arg ${ARGN})
|
||||
if("x${arg}" STREQUAL "xHIPCC_OPTIONS")
|
||||
set(_hipcc_found_options TRUE)
|
||||
elseif("x${arg}" STREQUAL "xHCC_OPTIONS")
|
||||
set(_hcc_found_options TRUE)
|
||||
elseif("x${arg}" STREQUAL "xNVCC_OPTIONS")
|
||||
set(_nvcc_found_options TRUE)
|
||||
elseif(
|
||||
"x${arg}" STREQUAL "xEXCLUDE_FROM_ALL" OR
|
||||
"x${arg}" STREQUAL "xSTATIC" OR
|
||||
"x${arg}" STREQUAL "xSHARED" OR
|
||||
"x${arg}" STREQUAL "xMODULE"
|
||||
)
|
||||
list(APPEND ${_cmake_options} ${arg})
|
||||
else()
|
||||
if(_hipcc_found_options)
|
||||
list(APPEND ${_hipcc_options} ${arg})
|
||||
elseif(_hcc_found_options)
|
||||
list(APPEND ${_hcc_options} ${arg})
|
||||
elseif(_nvcc_found_options)
|
||||
list(APPEND ${_nvcc_options} ${arg})
|
||||
else()
|
||||
# Assume this is a file
|
||||
list(APPEND ${_sources} ${arg})
|
||||
endif()
|
||||
endif()
|
||||
endforeach()
|
||||
endmacro()
|
||||
|
||||
###############################################################################
|
||||
# MACRO: Add include directories to pass to the hipcc command
|
||||
###############################################################################
|
||||
set(HIP_HIPCC_INCLUDE_ARGS_USER "")
|
||||
macro(HIP_INCLUDE_DIRECTORIES)
|
||||
foreach(dir ${ARGN})
|
||||
list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER -I${dir})
|
||||
endforeach()
|
||||
endmacro()
|
||||
|
||||
###############################################################################
|
||||
# FUNCTION: Helper to avoid clashes of files with the same basename but different paths
|
||||
###############################################################################
|
||||
function(HIP_COMPUTE_BUILD_PATH path build_path)
|
||||
# Convert to cmake style paths
|
||||
file(TO_CMAKE_PATH "${path}" bpath)
|
||||
if (IS_ABSOLUTE "${bpath}")
|
||||
string(FIND "${bpath}" "${CMAKE_CURRENT_BINARY_DIR}" _binary_dir_pos)
|
||||
if (_binary_dir_pos EQUAL 0)
|
||||
file(RELATIVE_PATH bpath "${CMAKE_CURRENT_BINARY_DIR}" "${bpath}")
|
||||
else()
|
||||
file(RELATIVE_PATH bpath "${CMAKE_CURRENT_SOURCE_DIR}" "${bpath}")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
# Remove leading /
|
||||
string(REGEX REPLACE "^[/]+" "" bpath "${bpath}")
|
||||
# Avoid absolute paths by removing ':'
|
||||
string(REPLACE ":" "_" bpath "${bpath}")
|
||||
# Avoid relative paths that go up the tree
|
||||
string(REPLACE "../" "__/" bpath "${bpath}")
|
||||
# Avoid spaces
|
||||
string(REPLACE " " "_" bpath "${bpath}")
|
||||
# Strip off the filename
|
||||
get_filename_component(bpath "${bpath}" PATH)
|
||||
|
||||
set(${build_path} "${bpath}" PARENT_SCOPE)
|
||||
endfunction()
|
||||
|
||||
###############################################################################
|
||||
# MACRO: Parse OPTIONS from ARGN & set variables prefixed by _option_prefix
|
||||
###############################################################################
|
||||
macro(HIP_PARSE_HIPCC_OPTIONS _option_prefix)
|
||||
set(_hip_found_config)
|
||||
foreach(arg ${ARGN})
|
||||
# Determine if we are dealing with a per-configuration flag
|
||||
foreach(config ${_hip_configuration_types})
|
||||
string(TOUPPER ${config} config_upper)
|
||||
if(arg STREQUAL "${config_upper}")
|
||||
set(_hip_found_config _${arg})
|
||||
# Clear arg to prevent it from being processed anymore
|
||||
set(arg)
|
||||
endif()
|
||||
endforeach()
|
||||
if(arg)
|
||||
list(APPEND ${_option_prefix}${_hip_found_config} "${arg}")
|
||||
endif()
|
||||
endforeach()
|
||||
endmacro()
|
||||
|
||||
###############################################################################
|
||||
# MACRO: Try and include dependency file if it exists
|
||||
###############################################################################
|
||||
macro(HIP_INCLUDE_HIPCC_DEPENDENCIES dependency_file)
|
||||
set(HIP_HIPCC_DEPEND)
|
||||
set(HIP_HIPCC_DEPEND_REGENERATE FALSE)
|
||||
|
||||
# Create the dependency file if it doesn't exist
|
||||
if(NOT EXISTS ${dependency_file})
|
||||
file(WRITE ${dependency_file} "# Generated by: FindHIP.cmake. Do not edit.\n")
|
||||
endif()
|
||||
# Include the dependency file
|
||||
include(${dependency_file})
|
||||
|
||||
# Verify the existence of all the included files
|
||||
if(HIP_HIPCC_DEPEND)
|
||||
foreach(f ${HIP_HIPCC_DEPEND})
|
||||
if(NOT EXISTS ${f})
|
||||
# If they aren't there, regenerate the file again
|
||||
set(HIP_HIPCC_DEPEND_REGENERATE TRUE)
|
||||
endif()
|
||||
endforeach()
|
||||
else()
|
||||
# No dependencies, so regenerate the file
|
||||
set(CUDA_NVCC_DEPEND_REGENERATE TRUE)
|
||||
endif()
|
||||
|
||||
# Regenerate the dependency file if needed
|
||||
if(HIP_HIPCC_DEPEND_REGENERATE)
|
||||
set(HIP_HIPCC_DEPEND ${dependency_file})
|
||||
file(WRITE ${dependency_file} "# Generated by: FindHIP.cmake. Do not edit.\n")
|
||||
endif()
|
||||
endmacro()
|
||||
|
||||
###############################################################################
|
||||
# MACRO: Prepare cmake commands for the target
|
||||
###############################################################################
|
||||
macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files)
|
||||
set(_hip_flags "")
|
||||
set(_hip_build_configuration "${CMAKE_BUILD_TYPE}")
|
||||
if(HIP_HOST_COMPILATION_CPP)
|
||||
set(HIP_C_OR_CXX CXX)
|
||||
else()
|
||||
set(HIP_C_OR_CXX C)
|
||||
endif()
|
||||
set(generated_extension ${CMAKE_${HIP_C_OR_CXX}_OUTPUT_EXTENSION})
|
||||
|
||||
# Initialize list of includes with those specified by the user. Append with
|
||||
# ones specified to cmake directly.
|
||||
set(HIP_HIPCC_INCLUDE_ARGS ${HIP_HIPCC_INCLUDE_ARGS_USER})
|
||||
get_directory_property(_hip_include_directories INCLUDE_DIRECTORIES)
|
||||
list(REMOVE_DUPLICATES _hip_include_directories)
|
||||
if(_hip_include_directories)
|
||||
foreach(dir ${_hip_include_directories})
|
||||
list(APPEND HIP_HIPCC_INCLUDE_ARGS -I${dir})
|
||||
endforeach()
|
||||
endif()
|
||||
|
||||
HIP_GET_SOURCES_AND_OPTIONS(_hip_sources _hip_cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN})
|
||||
HIP_PARSE_HIPCC_OPTIONS(HIP_HIPCC_FLAGS ${_hipcc_options})
|
||||
HIP_PARSE_HIPCC_OPTIONS(HIP_HCC_FLAGS ${_hcc_options})
|
||||
HIP_PARSE_HIPCC_OPTIONS(HIP_NVCC_FLAGS ${_nvcc_options})
|
||||
|
||||
# Set host compiler
|
||||
set(HIP_HOST_COMPILER "${CMAKE_${HIP_C_OR_CXX}_COMPILER}")
|
||||
|
||||
# Set compiler flags
|
||||
set(_HIP_HOST_FLAGS "set(CMAKE_HOST_FLAGS ${CMAKE_${HIP_C_OR_CXX}_FLAGS})")
|
||||
set(_HIP_HIPCC_FLAGS "set(HIP_HIPCC_FLAGS ${HIP_HIPCC_FLAGS})")
|
||||
set(_HIP_HCC_FLAGS "set(HIP_HCC_FLAGS ${HIP_HCC_FLAGS})")
|
||||
set(_HIP_NVCC_FLAGS "set(HIP_NVCC_FLAGS ${HIP_NVCC_FLAGS})")
|
||||
foreach(config ${_hip_configuration_types})
|
||||
string(TOUPPER ${config} config_upper)
|
||||
set(_HIP_HOST_FLAGS "${_HIP_HOST_FLAGS}\nset(CMAKE_HOST_FLAGS_${config_upper} ${CMAKE_${HIP_C_OR_CXX}_FLAGS_${config_upper}})")
|
||||
set(_HIP_HIPCC_FLAGS "${_HIP_HIPCC_FLAGS}\nset(HIP_HIPCC_FLAGS_${config_upper} ${HIP_HIPCC_FLAGS_${config_upper}})")
|
||||
set(_HIP_HCC_FLAGS "${_HIP_HCC_FLAGS}\nset(HIP_HCC_FLAGS_${config_upper} ${HIP_HCC_FLAGS_${config_upper}})")
|
||||
set(_HIP_NVCC_FLAGS "${_HIP_NVCC_FLAGS}\nset(HIP_NVCC_FLAGS_${config_upper} ${HIP_NVCC_FLAGS_${config_upper}})")
|
||||
endforeach()
|
||||
|
||||
# Reset the output variable
|
||||
set(_hip_generated_files "")
|
||||
|
||||
# Iterate over all arguments and create custom commands for all source files
|
||||
foreach(file ${ARGN})
|
||||
# Ignore any file marked as a HEADER_FILE_ONLY
|
||||
get_source_file_property(_is_header ${file} HEADER_FILE_ONLY)
|
||||
# Allow per source file overrides of the format. Also allows compiling non .cu files.
|
||||
get_source_file_property(_hip_source_format ${file} HIP_SOURCE_PROPERTY_FORMAT)
|
||||
if((${file} MATCHES "\\.cu$" OR _hip_source_format) AND NOT _is_header)
|
||||
set(host_flag FALSE)
|
||||
else()
|
||||
set(host_flag TRUE)
|
||||
endif()
|
||||
|
||||
if (NOT host_flag)
|
||||
# Determine output directory
|
||||
HIP_COMPUTE_BUILD_PATH("${file}" hip_build_path)
|
||||
set(hip_compile_output_dir "${CMAKE_CURRENT_BINARY_DIR}/CMakeFiles/${_target}.dir/${hip_build_path}")
|
||||
|
||||
get_filename_component(basename ${file} NAME)
|
||||
set(generated_file_path "${hip_compile_output_dir}/${CMAKE_CFG_INTDIR}")
|
||||
set(generated_file_basename "${_target}_generated_${basename}${generated_extension}")
|
||||
|
||||
# Set file names
|
||||
set(generated_file "${generated_file_path}/${generated_file_basename}")
|
||||
set(cmake_dependency_file "${hip_compile_output_dir}/${generated_file_basename}.depend")
|
||||
set(custom_target_script_pregen "${hip_compile_output_dir}/${generated_file_basename}.cmake.pre-gen")
|
||||
set(custom_target_script "${hip_compile_output_dir}/${generated_file_basename}.cmake")
|
||||
|
||||
# Set properties for object files
|
||||
set_source_files_properties("${generated_file}"
|
||||
PROPERTIES
|
||||
EXTERNAL_OBJECT true # This is an object file not to be compiled, but only be linked
|
||||
)
|
||||
|
||||
# Don't add CMAKE_CURRENT_SOURCE_DIR if the path is already an absolute path
|
||||
get_filename_component(file_path "${file}" PATH)
|
||||
if(IS_ABSOLUTE "${file_path}")
|
||||
set(source_file "${file}")
|
||||
else()
|
||||
set(source_file "${CMAKE_CURRENT_SOURCE_DIR}/${file}")
|
||||
endif()
|
||||
|
||||
# Bring in the dependencies
|
||||
HIP_INCLUDE_HIPCC_DEPENDENCIES(${cmake_dependency_file})
|
||||
|
||||
# Configure the build script
|
||||
configure_file("${HIP_run_hipcc}" "${custom_target_script_pregen}" @ONLY)
|
||||
file(GENERATE
|
||||
OUTPUT "${custom_target_script}"
|
||||
INPUT "${custom_target_script_pregen}"
|
||||
)
|
||||
set(main_dep DEPENDS ${source_file})
|
||||
set(verbose_output "$(VERBOSE)")
|
||||
|
||||
# Create up the comment string
|
||||
file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}")
|
||||
set(hip_build_comment_string "Building HIPCC (${cuda_build_type}) object ${generated_file_relative_path}")
|
||||
|
||||
# Build the generated file and dependency file
|
||||
add_custom_command(
|
||||
OUTPUT ${generated_file}
|
||||
# These output files depend on the source_file and the contents of cmake_dependency_file
|
||||
${main_dep}
|
||||
DEPENDS ${HIP_HIPCC_DEPEND}
|
||||
DEPENDS ${custom_target_script}
|
||||
# Make sure the output directory exists before trying to write to it.
|
||||
COMMAND ${CMAKE_COMMAND} -E make_directory "${generated_file_path}"
|
||||
COMMAND ${CMAKE_COMMAND} ARGS
|
||||
-D verbose:BOOL=${verbose_output}
|
||||
-D build_configuration:STRING=${_hip_build_configuration}
|
||||
-D "generated_file:STRING=${generated_file}"
|
||||
-P "${custom_target_script}"
|
||||
WORKING_DIRECTORY "${hip_compile_output_dir}"
|
||||
COMMENT "${hip_build_comment_string}"
|
||||
)
|
||||
|
||||
# Make sure the build system knows the file is generated
|
||||
set_source_files_properties(${generated_file} PROPERTIES GENERATED TRUE)
|
||||
list(APPEND _hip_generated_files ${generated_file})
|
||||
endif()
|
||||
endforeach()
|
||||
|
||||
# Set the return parameter
|
||||
set(${_generated_files} ${_hip_generated_files})
|
||||
endmacro()
|
||||
|
||||
###############################################################################
|
||||
# HIP_ADD_EXECUTABLE
|
||||
###############################################################################
|
||||
macro(HIP_ADD_EXECUTABLE hip_target)
|
||||
# Separate the sources from the options
|
||||
HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN})
|
||||
HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options})
|
||||
set(HIP_CMAKE_CXX_LINK_EXECUTABLE ${CMAKE_CXX_LINK_EXECUTABLE})
|
||||
set(CMAKE_CXX_LINK_EXECUTABLE "${HIP_HIPCC_EXECUTABLE} <FLAGS> <CMAKE_CXX_LINK_FLAGS> <LINK_FLAGS> <OBJECTS> -o <TARGET>")
|
||||
add_executable(${hip_target} ${_cmake_options} ${_generated_files} ${_sources})
|
||||
set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE ${HIP_C_OR_CXX})
|
||||
#set(CMAKE_CXX_COMPILER ${ORIGINAL_CMAKE_CXX_COMPILER})
|
||||
endmacro()
|
||||
|
||||
# vim: ts=4:sw=4:expandtab:smartindent
|
||||
|
||||
@@ -0,0 +1,163 @@
|
||||
###############################################################################
|
||||
# Runs commands using HIPCC
|
||||
###############################################################################
|
||||
|
||||
###############################################################################
|
||||
# This file runs the hipcc commands to produce the desired output file
|
||||
# along with the dependency file needed by CMake to compute dependencies.
|
||||
#
|
||||
# Input variables:
|
||||
#
|
||||
# verbose:BOOL=<> OFF: Be as quiet as possible (default)
|
||||
# ON : Describe each step
|
||||
# build_configuration:STRING=<> Build configuration. Defaults to Debug.
|
||||
# generated_file:STRING=<> File to generate. Mandatory argument.
|
||||
|
||||
if(NOT build_configuration)
|
||||
set(build_configuration Debug)
|
||||
endif()
|
||||
if(NOT generated_file)
|
||||
message(FATAL_ERROR "You must specify generated_file on the command line")
|
||||
endif()
|
||||
|
||||
# Set these up as variables to make reading the generated file easier
|
||||
set(HIP_HIPCC_EXECUTABLE "@HIP_HIPCC_EXECUTABLE@") # path
|
||||
set(HIP_HOST_COMPILER "@CUDA_HOST_COMPILER@") # path
|
||||
set(HIP_PLATFORM "@HIP_PLATFORM@") #string
|
||||
set(CMAKE_COMMAND "@CMAKE_COMMAND@") # path
|
||||
set(HIP_run_make2cmake "@HIP_run_make2cmake@") # path
|
||||
|
||||
@HIP_HOST_FLAGS@
|
||||
@_HIP_HIPCC_FLAGS@
|
||||
@_HIP_HCC_FLAGS@
|
||||
@_HIP_NVCC_FLAGS@
|
||||
set(HIP_HIPCC_INCLUDE_ARGS "@HIP_HIPCC_INCLUDE_ARGS@") # list (needs to be in quotes to handle spaces properly)
|
||||
|
||||
set(cmake_dependency_file "@cmake_dependency_file@") # path
|
||||
set(source_file "@source_file@") # path
|
||||
set(host_flag "@host_flag@") # bool
|
||||
|
||||
# Determine compiler and compiler flags
|
||||
if(NOT host_flag)
|
||||
set(__CC ${HIP_HIPCC_EXECUTABLE})
|
||||
if(HIP_PLATFORM STREQUAL "hcc")
|
||||
set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_HCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_HCC_FLAGS_${build_configuration}})
|
||||
else()
|
||||
set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_NVCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_NVCC_FLAGS_${build_configuration}})
|
||||
endif()
|
||||
else()
|
||||
set(__CC ${HIP_HOST_COMPILER})
|
||||
set(__CC_FLAGS ${CMAKE_HOST_FLAGS} ${CMAKE_HOST_FLAGS_${build_configuration}})
|
||||
endif()
|
||||
set(__CC_INCLUDES ${HIP_HIPCC_INCLUDE_ARGS})
|
||||
|
||||
# hip_execute_process - Executes a command with optional command echo and status message.
|
||||
# status - Status message to print if verbose is true
|
||||
# command - COMMAND argument from the usual execute_process argument structure
|
||||
# ARGN - Remaining arguments are the command with arguments
|
||||
# HIP_result - Return value from running the command
|
||||
macro(hip_execute_process status command)
|
||||
set(_command ${command})
|
||||
if(NOT "x${_command}" STREQUAL "xCOMMAND")
|
||||
message(FATAL_ERROR "Malformed call to hip_execute_process. Missing COMMAND as second argument. (command = ${command})")
|
||||
endif()
|
||||
if(verbose)
|
||||
execute_process(COMMAND "${CMAKE_COMMAND}" -E echo -- ${status})
|
||||
# Build command string to print
|
||||
set(hip_execute_process_string)
|
||||
foreach(arg ${ARGN})
|
||||
# Escape quotes if any
|
||||
string(REPLACE "\"" "\\\"" arg ${arg})
|
||||
# Surround args with spaces with quotes
|
||||
if(arg MATCHES " ")
|
||||
list(APPEND hip_execute_process_string "\"${arg}\"")
|
||||
else()
|
||||
list(APPEND hip_execute_process_string ${arg})
|
||||
endif()
|
||||
endforeach()
|
||||
# Echo the command
|
||||
execute_process(COMMAND ${CMAKE_COMMAND} -E echo ${hip_execute_process_string})
|
||||
endif()
|
||||
# Run the command
|
||||
execute_process(COMMAND ${ARGN} RESULT_VARIABLE HIP_result)
|
||||
endmacro()
|
||||
|
||||
# Delete the target file
|
||||
hip_execute_process(
|
||||
"Removing ${generated_file}"
|
||||
COMMAND "${CMAKE_COMMAND}" -E remove "${generated_file}"
|
||||
)
|
||||
|
||||
# Generate the dependency file
|
||||
hip_execute_process(
|
||||
"Generating dependency file: ${cmake_dependency_file}.pre"
|
||||
COMMAND "${__CC}"
|
||||
-M
|
||||
"${source_file}"
|
||||
-o "${cmake_dependency_file}.pre"
|
||||
${__CC_FLAGS}
|
||||
${__CC_INCLUDES}
|
||||
)
|
||||
|
||||
if(HIP_result)
|
||||
message(FATAL_ERROR "Error generating ${generated_file}")
|
||||
endif()
|
||||
|
||||
# Generate the cmake readable dependency file to a temp file
|
||||
hip_execute_process(
|
||||
"Generating temporary cmake readable file: ${cmake_dependency_file}.tmp"
|
||||
COMMAND "${CMAKE_COMMAND}"
|
||||
-D "input_file:FILEPATH=${cmake_dependency_file}.pre"
|
||||
-D "output_file:FILEPATH=${cmake_dependency_file}.tmp"
|
||||
-D "verbose=${verbose}"
|
||||
-P "${HIP_run_make2cmake}"
|
||||
)
|
||||
|
||||
if(HIP_result)
|
||||
message(FATAL_ERROR "Error generating ${generated_file}")
|
||||
endif()
|
||||
|
||||
# Copy the file if it is different
|
||||
hip_execute_process(
|
||||
"Copy if different ${cmake_dependency_file}.tmp to ${cmake_dependency_file}"
|
||||
COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${cmake_dependency_file}.tmp" "${cmake_dependency_file}"
|
||||
)
|
||||
|
||||
if(HIP_result)
|
||||
message(FATAL_ERROR "Error generating ${generated_file}")
|
||||
endif()
|
||||
|
||||
# Delete the temporary file
|
||||
hip_execute_process(
|
||||
"Removing ${cmake_dependency_file}.tmp and ${cmake_dependency_file}.pre"
|
||||
COMMAND "${CMAKE_COMMAND}" -E remove "${cmake_dependency_file}.tmp" "${cmake_dependency_file}.pre"
|
||||
)
|
||||
|
||||
if(HIP_result)
|
||||
message(FATAL_ERROR "Error generating ${generated_file}")
|
||||
endif()
|
||||
|
||||
# Generate the output file
|
||||
hip_execute_process(
|
||||
"Generating ${generated_file}"
|
||||
COMMAND "${__CC}"
|
||||
-c
|
||||
"${source_file}"
|
||||
-o "${generated_file}"
|
||||
${__CC_FLAGS}
|
||||
${__CC_INCLUDES}
|
||||
)
|
||||
|
||||
if(HIP_result)
|
||||
# Make sure that we delete the output file
|
||||
hip_execute_process(
|
||||
"Removing ${generated_file}"
|
||||
COMMAND "${CMAKE_COMMAND}" -E remove "${generated_file}"
|
||||
)
|
||||
message(FATAL_ERROR "Error generating file ${generated_file}")
|
||||
else()
|
||||
if(verbose)
|
||||
message("Generated ${generated_file} successfully.")
|
||||
endif()
|
||||
endif()
|
||||
# vim: ts=4:sw=4:expandtab:smartindent
|
||||
@@ -0,0 +1,50 @@
|
||||
###############################################################################
|
||||
# Computes dependencies using HIPCC
|
||||
###############################################################################
|
||||
|
||||
###############################################################################
|
||||
# This file converts dependency files generated using hipcc to a format that
|
||||
# cmake can understand.
|
||||
|
||||
# Input variables:
|
||||
#
|
||||
# input_file:STRING=<> Dependency file to parse. Required argument
|
||||
# output_file:STRING=<> Output file to generate. Required argument
|
||||
|
||||
if(NOT input_file OR NOT output_file)
|
||||
message(FATAL_ERROR "You must specify input_file and output_file on the command line")
|
||||
endif()
|
||||
|
||||
file(READ ${input_file} depend_text)
|
||||
|
||||
if (NOT "${depend_text}" STREQUAL "")
|
||||
string(REPLACE " /" "\n/" depend_text ${depend_text})
|
||||
string(REGEX REPLACE "^.*:" "" depend_text ${depend_text})
|
||||
string(REGEX REPLACE "[ \\\\]*\n" ";" depend_text ${depend_text})
|
||||
|
||||
set(dependency_list "")
|
||||
|
||||
foreach(file ${depend_text})
|
||||
string(REGEX REPLACE "^ +" "" file ${file})
|
||||
if(NOT EXISTS "${file}")
|
||||
message(WARNING " Removing non-existent dependency file: ${file}")
|
||||
set(file "")
|
||||
endif()
|
||||
|
||||
if(NOT IS_DIRECTORY "${file}")
|
||||
get_filename_component(file_absolute "${file}" ABSOLUTE)
|
||||
list(APPEND dependency_list "${file_absolute}")
|
||||
endif()
|
||||
endforeach()
|
||||
endif()
|
||||
|
||||
# Remove the duplicate entries and sort them.
|
||||
list(REMOVE_DUPLICATES dependency_list)
|
||||
list(SORT dependency_list)
|
||||
|
||||
foreach(file ${dependency_list})
|
||||
set(hip_hipcc_depend "${hip_hipcc_depend} \"${file}\"\n")
|
||||
endforeach()
|
||||
|
||||
file(WRITE ${output_file} "# Generated by: FindHIP.cmake. Do not edit.\nSET(HIP_HIPCC_DEPEND\n ${hip_hipcc_depend})\n\n")
|
||||
# vim: ts=4:sw=4:expandtab:smartindent
|
||||
@@ -213,6 +213,9 @@ export HIP_ATP_MARKER=1
|
||||
```shell
|
||||
# Use profile to generate timeline view:
|
||||
/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
|
||||
|
||||
Or
|
||||
/opt/rocm/bin/rocm-profiler -e HIP_ATP_MARKER=1 -o <outputATPFileName> -A <applicationName> <applicationArguments>
|
||||
```
|
||||
|
||||
#### Using HIP_TRACE_API
|
||||
|
||||
@@ -669,20 +669,14 @@ The following C++ features are not supported:
|
||||
- Try/catch
|
||||
|
||||
## Kernel Compilation
|
||||
HIP now supports compiling C++/HIP kernels to binary. Eventhough HIP does not support fatbinary (yet), the user can specify the target for which the binary can be generated. The file format for binary is `.co` which means Code Object. The following command builds the binary using `hipcc`.
|
||||
hipcc now supports compiling C++/HIP kernels to binary code objects.
|
||||
The user can specify the target for which the binary can be generated. HIP/HCC does not yet support fat binaries so only a single target may be specified.
|
||||
The file format for binary is `.co` which means Code Object. The following command builds the code object using `hipcc`.
|
||||
|
||||
`hipcc --genisa --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]`
|
||||
`hipcc --genco --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]`
|
||||
```[TARGET GPU] = fiji/hawaii
|
||||
[INPUT FILE] = Name of the file containing kernels
|
||||
[OUTPUT FILE] = Name of the generated code object file```
|
||||
|
||||
Note that the kernel file should have `int main(){}` at the end it so that the binary is generated. This happens because HCC generates binaries at linking time rather than compilation
|
||||
|
||||
You need 3 things to run kernel in binary.
|
||||
1. Kernel Binary
|
||||
2. Name of kernel binary
|
||||
3. Name of the kernel
|
||||
|
||||
We already got first two of them. In order to get name of the kernel, try `objdump -x [OUTPUT FILE]`. OUTPUT FILE is file generated by hipcc during kernel compilation. The output from objdump has symbol to the kernel whose name is mangled with `grid_launch_parm`, `__functor`, `__cxxamp_trampoline`. An example of how it looks is `ZN12_GLOBAL__N_137_Z3Cpy16grid_launch_parmPfS0__functor19__cxxamp_trampolineEiiiiiiPKfPf` where `Cpy` is the name of the kernel written in C++.
|
||||
|
||||
|
||||
|
||||
@@ -12,7 +12,7 @@ There are two possible ways to transfer data from Host to Device (H2D) and Devic
|
||||
|
||||
#### On Large BAR Setup
|
||||
|
||||
There are two possible ways to transfer data from Host to Device (H2D)
|
||||
There are three possible ways to transfer data from Host to Device (H2D)
|
||||
* Using Staging Buffers
|
||||
* Using PinInPlace
|
||||
* Direct Memcpy
|
||||
@@ -24,12 +24,9 @@ There are two possible ways to transfer data from Host to Device (H2D)
|
||||
Some GPUs may not be able to directly access host memory, and in these cases we need to
|
||||
stage the copy through an optimized pinned staging buffer, to implement H2D and D2H copies.The copy is broken into buffer-sized chunks to limit the size of the buffer and also to provide better performance by overlapping the CPU copies with the DMA copies.
|
||||
|
||||
PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA
|
||||
engine.
|
||||
PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA engine.
|
||||
|
||||
By default staging buffers are used for unpinned memory transfers, however other ways can be used by enabling few environment variables (so no need to build the code again!!!)
|
||||
|
||||
Following environment variables can be used:
|
||||
By default staging buffers are used for unpinned memory transfers. Environment variables allow control over the unpinned copy algorithm and parameters:
|
||||
|
||||
- HIP_PININPLACE - This environment variable forces the use of PinInPlace logic for all unpinned memory copies
|
||||
|
||||
|
||||
@@ -32,7 +32,7 @@ THE SOFTWARE.
|
||||
// #define USE_MEMCPYTOSYMBOL
|
||||
//
|
||||
//Use the new HCC accelerator_view::copy instead of am_copy
|
||||
#define USE_AV_COPY 0
|
||||
#define USE_AV_COPY (__hcc_workweek__ >= 16351)
|
||||
|
||||
// Compile peer-to-peer support.
|
||||
// >= 2 : use HCC hc:accelerator::get_is_peer
|
||||
@@ -353,18 +353,36 @@ struct LockedBase {
|
||||
};
|
||||
|
||||
|
||||
class ihipModule_t{
|
||||
public:
|
||||
hsa_executable_t executable;
|
||||
hsa_code_object_t object;
|
||||
std::string fileName;
|
||||
void *ptr;
|
||||
size_t size;
|
||||
};
|
||||
|
||||
|
||||
class ihipFunction_t{
|
||||
public:
|
||||
hsa_executable_symbol_t kernel_symbol;
|
||||
uint64_t kernel;
|
||||
};
|
||||
|
||||
|
||||
template <typename MUTEX_TYPE>
|
||||
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
|
||||
{
|
||||
public:
|
||||
ihipStreamCriticalBase_t() :
|
||||
ihipStreamCriticalBase_t(hc::accelerator_view av) :
|
||||
_last_command_type(ihipCommandCopyH2H),
|
||||
_last_copy_signal(NULL),
|
||||
_signalCursor(0),
|
||||
_oldest_live_sig_id(1),
|
||||
_streamSigId(0),
|
||||
_kernelCnt(0),
|
||||
_signalCnt(0)
|
||||
_signalCnt(0),
|
||||
_av(av)
|
||||
{
|
||||
_signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1);
|
||||
};
|
||||
@@ -395,27 +413,14 @@ public:
|
||||
// 2 are required if a barrier packet is inserted.
|
||||
uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait().
|
||||
SIGSEQNUM _streamSigId; // Monotonically increasing unique signal id.
|
||||
|
||||
hc::accelerator_view _av;
|
||||
};
|
||||
|
||||
|
||||
typedef ihipStreamCriticalBase_t<StreamMutex> ihipStreamCritical_t;
|
||||
typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
|
||||
|
||||
class ihipModule_t{
|
||||
public:
|
||||
hsa_executable_t executable;
|
||||
hsa_code_object_t object;
|
||||
std::string fileName;
|
||||
void *ptr;
|
||||
size_t size;
|
||||
};
|
||||
|
||||
|
||||
class ihipFunction_t{
|
||||
public:
|
||||
hsa_executable_symbol_t kernel_symbol;
|
||||
uint64_t kernel;
|
||||
};
|
||||
|
||||
// Internal stream structure.
|
||||
class ihipStream_t {
|
||||
@@ -431,21 +436,31 @@ typedef uint64_t SeqNum_t ;
|
||||
|
||||
void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
|
||||
|
||||
int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType);
|
||||
|
||||
//---
|
||||
// Thread-safe accessors - these acquire / release mutex:
|
||||
bool lockopen_preKernelCommand();
|
||||
// Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
|
||||
LockedAccessor_StreamCrit_t lockopen_preKernelCommand();
|
||||
void lockclose_postKernelCommand(hc::completion_future &kernel_future);
|
||||
|
||||
int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType);
|
||||
|
||||
void locked_reclaimSignals(SIGSEQNUM sigNum);
|
||||
void locked_wait(bool assertQueueEmpty=false);
|
||||
SIGSEQNUM locked_lastCopySeqId() {LockedAccessor_StreamCrit_t crit(_criticalData); return lastCopySeqId(crit); };
|
||||
|
||||
hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); };
|
||||
|
||||
void locked_waitEvent(hipEvent_t event);
|
||||
void locked_recordEvent(hipEvent_t event);
|
||||
|
||||
//---
|
||||
|
||||
// Use this if we already have the stream critical data mutex:
|
||||
void wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty=false);
|
||||
|
||||
void launchModuleKernel(hsa_signal_t signal, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, uint32_t sharedMemBytes, void *kernarg, size_t kernSize, uint64_t kernel);
|
||||
void launchModuleKernel(hc::accelerator_view av, hsa_signal_t signal,
|
||||
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
|
||||
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
|
||||
uint32_t sharedMemBytes, void *kernarg, size_t kernSize, uint64_t kernel);
|
||||
|
||||
// Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function.
|
||||
SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sigId : 0; };
|
||||
@@ -462,7 +477,6 @@ public:
|
||||
//---
|
||||
//Public member vars - these are set at initialization and never change:
|
||||
SeqNum_t _id; // monotonic sequence ID
|
||||
hc::accelerator_view _av;
|
||||
unsigned _flags;
|
||||
|
||||
|
||||
@@ -672,7 +686,7 @@ extern const char *ihipErrorString(hipError_t);
|
||||
extern ihipCtx_t *ihipGetTlsDefaultCtx();
|
||||
extern void ihipSetTlsDefaultCtx(ihipCtx_t *ctx);
|
||||
extern hipError_t ihipSynchronize(void);
|
||||
extern hipError_t ihipCtxStackUpdate();
|
||||
extern void ihipCtxStackUpdate();
|
||||
|
||||
extern ihipDevice_t *ihipGetDevice(int);
|
||||
ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex);
|
||||
@@ -704,5 +718,12 @@ inline std::ostream & operator<<(std::ostream& os, const dim3& s)
|
||||
return os;
|
||||
}
|
||||
|
||||
// Stream printf functions:
|
||||
inline std::ostream& operator<<(std::ostream& os, const hipEvent_t& e)
|
||||
{
|
||||
os << "event:" << std::hex << static_cast<void*> (e);
|
||||
return os;
|
||||
}
|
||||
|
||||
|
||||
#endif
|
||||
|
||||
@@ -50,9 +50,8 @@ THE SOFTWARE.
|
||||
|
||||
#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.
|
||||
#define USE_GRID_LAUNCH_20 1
|
||||
#else
|
||||
#define USE_GRID_LAUNCH_20 0
|
||||
#error (HCC must support GRID_LAUNCH_20)
|
||||
#endif
|
||||
|
||||
#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01)
|
||||
@@ -532,7 +531,7 @@ __device__ void __threadfence_block(void);
|
||||
*
|
||||
* @warning __threadfence is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
|
||||
*/
|
||||
__device__ void __threadfence(void);
|
||||
__device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
|
||||
|
||||
/**
|
||||
* @brief threadfence_system makes writes to pinned system memory visible on host CPU.
|
||||
@@ -543,7 +542,7 @@ __device__ void __threadfence(void);
|
||||
*
|
||||
* @warning __threadfence_system is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
|
||||
*/
|
||||
__device__ void __threadfence_system(void);
|
||||
__device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
|
||||
|
||||
|
||||
// doxygen end Memory Fence
|
||||
@@ -633,7 +632,6 @@ extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
|
||||
#define KNRM "\x1B[0m"
|
||||
#define KGRN "\x1B[32m"
|
||||
|
||||
#if USE_GRID_LAUNCH_20
|
||||
#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
|
||||
do {\
|
||||
grid_launch_parm lp;\
|
||||
@@ -645,21 +643,6 @@ do {\
|
||||
_kernelName (lp, ##__VA_ARGS__);\
|
||||
ihipPostLaunchKernel(trueStream, lp);\
|
||||
} while(0)
|
||||
#else
|
||||
#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
|
||||
do {\
|
||||
grid_launch_parm lp;\
|
||||
lp.groupMemBytes = _groupMemBytes; \
|
||||
hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp)); \
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d stream=%p\n" KNRM, \
|
||||
#_kernelName, lp.gridDim.x, lp.gridDim.y, lp.gridDim.z, lp.groupDim.x, lp.groupDim.y, lp.groupDim.z, lp.groupMemBytes, (void*)(_stream));\
|
||||
}\
|
||||
_kernelName (lp, ##__VA_ARGS__);\
|
||||
ihipPostLaunchKernel(trueStream, lp);\
|
||||
} while(0)
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
#elif defined (__HCC_C__)
|
||||
|
||||
@@ -726,7 +726,7 @@ hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute
|
||||
hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags) ;
|
||||
|
||||
/**
|
||||
* @brief Get flags associated with host pointer
|
||||
* @brief Return flags associated with host pointer
|
||||
*
|
||||
* @param[out] flagsPtr Memory location to store flags
|
||||
* @param[in] hostPtr Host Pointer allocated through hipHostMalloc
|
||||
@@ -1186,13 +1186,12 @@ hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig );
|
||||
hipError_t hipCtxSynchronize ( void );
|
||||
|
||||
/**
|
||||
* @brief Get flags used for creating current/default context.
|
||||
* @brief Return flags used for creating default context.
|
||||
*
|
||||
* @param [out] flags
|
||||
*
|
||||
* @returns #hipSuccess.
|
||||
*/
|
||||
|
||||
hipError_t hipCtxGetFlags ( unsigned int* flags );
|
||||
|
||||
/**
|
||||
@@ -1235,6 +1234,53 @@ hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx);
|
||||
*/
|
||||
hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId);
|
||||
|
||||
/**
|
||||
* @brief Returns a handle to a compute device
|
||||
* @param [out] device
|
||||
* @param [in] ordinal
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInavlidDevice
|
||||
*/
|
||||
hipError_t hipDeviceGet(hipDevice_t *device, int ordinal);
|
||||
|
||||
/**
|
||||
* @brief Returns the compute capability of the device
|
||||
* @param [out] major
|
||||
* @param [out] minor
|
||||
* @param [in] device
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInavlidDevice
|
||||
*/
|
||||
hipError_t hipDeviceComputeCapability(int *major,int *minor,hipDevice_t device);
|
||||
|
||||
/**
|
||||
* @brief Returns an identifer string for the device.
|
||||
* @param [out] name
|
||||
* @param [in] len
|
||||
* @param [in] device
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInavlidDevice
|
||||
*/
|
||||
hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
|
||||
|
||||
/**
|
||||
* @brief Returns a PCI Bus Id string for the device.
|
||||
* @param [out] pciBusId
|
||||
* @param [in] len
|
||||
* @param [in] device
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInavlidDevice
|
||||
*/
|
||||
hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device);
|
||||
|
||||
/**
|
||||
* @brief Returns the total amount of memory on the device.
|
||||
* @param [out] bytes
|
||||
* @param [in] device
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInavlidDevice
|
||||
*/
|
||||
hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device);
|
||||
|
||||
/**
|
||||
* @brief Returns the approximate HIP driver version.
|
||||
@@ -1261,7 +1307,7 @@ hipError_t hipDriverGetVersion(int *driverVersion) ;
|
||||
hipError_t hipModuleLoad(hipModule_t *module, const char *fname);
|
||||
|
||||
/**
|
||||
* @brief Freeing the module
|
||||
* @brief Frees the module
|
||||
*
|
||||
* @param [in] module
|
||||
*
|
||||
@@ -1273,7 +1319,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname);
|
||||
hipError_t hipModuleUnload(hipModule_t module);
|
||||
|
||||
/**
|
||||
* @brief Function with kname will be extracted present in module
|
||||
* @brief Function with kname will be extracted if present in module
|
||||
*
|
||||
* @param [in] module
|
||||
* @param [in] kname
|
||||
@@ -1284,19 +1330,20 @@ hipError_t hipModuleUnload(hipModule_t module);
|
||||
hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, const char *kname);
|
||||
|
||||
/**
|
||||
* @brief returns device memory pointer and size of the kernel present in the module with symbol - name
|
||||
* @brief returns device memory pointer and size of the kernel present in the module with symbol @p name
|
||||
*
|
||||
* @param [in] moodule
|
||||
* @param [in] name
|
||||
* @param [out] dptr
|
||||
* @param [out[ bytes
|
||||
* @param [in] hmod
|
||||
* @param [in] name
|
||||
*
|
||||
* @returns hipSuccess, hipErrorInvalidValue, hipErrorNotInitialized
|
||||
*/
|
||||
hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);
|
||||
|
||||
|
||||
/**
|
||||
* @brief builds module from code object which resides in host memory. And image is pointer to that location.
|
||||
* @brief builds module from code object which resides in host memory. Image is pointer to that location.
|
||||
*
|
||||
* @param [in] image
|
||||
* @param [out] module
|
||||
@@ -1305,8 +1352,9 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, hipModule_t h
|
||||
*/
|
||||
hipError_t hipModuleLoadData(hipModule_t *module, const void *image);
|
||||
|
||||
|
||||
/**
|
||||
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kerneelparams or extra
|
||||
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra
|
||||
*
|
||||
* @param [in[ f
|
||||
* @param [in] gridDimX
|
||||
|
||||
@@ -63,7 +63,9 @@ inline std::string ToString(T v)
|
||||
template <>
|
||||
inline std::string ToString(hipEvent_t v)
|
||||
{
|
||||
return ToString(&v);
|
||||
std::ostringstream ss;
|
||||
ss << v;
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -484,6 +484,17 @@ inline static hipError_t hipStreamDestroy(hipStream_t stream)
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamWaitEvent(stream, event, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipStreamQuery(hipStream_t stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamQuery(stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
cudaError_t err = cudaDriverGetVersion(driverVersion);
|
||||
@@ -624,6 +635,26 @@ inline static hipError_t hipDeviceGet(hipDevice_t *device, int ordinal)
|
||||
return hipCUResultTohipError(cuDeviceGet(device, ordinal));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceComputeCapability(major,minor,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceGetName(name,len,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceTotalMem(bytes,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleLoad(hipModule_t *module, const char* fname)
|
||||
{
|
||||
return hipCUResultTohipError(cuModuleLoad(module, fname));
|
||||
|
||||
@@ -4,6 +4,7 @@ project(hip_base)
|
||||
install(DIRECTORY @hip_SOURCE_DIR@/bin DESTINATION . USE_SOURCE_PERMISSIONS)
|
||||
install(DIRECTORY @hip_SOURCE_DIR@/include DESTINATION . PATTERN "hip" EXCLUDE)
|
||||
install(FILES @PROJECT_BINARY_DIR@/.version DESTINATION bin)
|
||||
install(DIRECTORY @hip_SOURCE_DIR@/cmake DESTINATION .)
|
||||
|
||||
#############################
|
||||
# Packaging steps
|
||||
|
||||
@@ -1,8 +1,9 @@
|
||||
HCC_HOME?=/opt/rocm/hcc
|
||||
HCC = $(HCC_HOME)/bin/hcc
|
||||
|
||||
HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags`
|
||||
HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags`
|
||||
OPT=-O2
|
||||
HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` ${OPT}
|
||||
HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` ${OPT}
|
||||
|
||||
CPPAMP_CFLAGS= -std=c++amp -stdlib=libc++ -I$(HCC_HOME)/include
|
||||
CPPAMP_LDFLAGS= -std=c++amp -L$(HCC_HOME)/lib -Wl,--rpath=$(HCC_HOME)/lib -lc++ -lc++abi -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive
|
||||
@@ -19,7 +20,7 @@ $(error hcc_dialects requires hcc compiler and only runs on hcc platform)
|
||||
endif
|
||||
|
||||
|
||||
TARGETS=vadd_hc_arrayview vadd_hc_array vadd_amp_arrayview vadd_hip
|
||||
TARGETS=vadd_hc_arrayview vadd_hc_array vadd_hc_am vadd_amp_arrayview vadd_hip
|
||||
|
||||
all: $(TARGETS)
|
||||
|
||||
@@ -51,7 +52,7 @@ vadd_hc_array: vadd_hc_array.o
|
||||
vadd_hc_am.o: vadd_hc_am.cpp
|
||||
$(HCC) $(HCC_CFLAGS) -c $< -o $@
|
||||
vadd_hc_am: vadd_hc_am.o
|
||||
$(HCC) $(HCC_LDFLAGS) $< -o $@
|
||||
$(HCC) $(HCC_LDFLAGS) -lhc_am $< -o $@
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -32,23 +32,25 @@ int main(int argc, char *argv[])
|
||||
for (int i=0; i<sizeElements; i++) {
|
||||
A_h[i] = 1.618f * i;
|
||||
B_h[i] = 3.142f * i;
|
||||
C_h[i] = 0;
|
||||
}
|
||||
|
||||
av.copy(A_h, A_d); // C++ copy H2D
|
||||
av.copy(B_h, B_d); //C++ copy H2D
|
||||
av.copy(A_h, A_d, sizeBytes); // C++ copy H2D
|
||||
av.copy(B_h, B_d, sizeBytes); // C++ copy H2D
|
||||
|
||||
// Launch kernel onto AV.
|
||||
// Because the kernel PFE and the copies are submitted to same AV, they will execute in order
|
||||
// and we don't need additional synchronization to ensure the copies complete before the PFE begins.
|
||||
hc::completion_future cf=
|
||||
hc::parallel_for_each(av, hc::extent<1> (sizeElements),
|
||||
[&] (hc::index<1> idx) [[hc]] {
|
||||
[=] (hc::index<1> idx) [[hc]] {
|
||||
int i = idx[0];
|
||||
C_d[i] = A_d[i] + B_d[i];
|
||||
});
|
||||
|
||||
|
||||
// This copy is in same AV as the kernel and thus will wait for the kernel to finish before executing.
|
||||
av.copy(C_d, C_h); // C++ copy D2H
|
||||
av.copy(C_d, C_h, sizeBytes); // C++ copy D2H
|
||||
|
||||
|
||||
for (int i=0; i<sizeElements; i++) {
|
||||
|
||||
@@ -27,22 +27,25 @@ int main(int argc, char *argv[])
|
||||
hipMalloc(&B_d, sizeBytes);
|
||||
hipMalloc(&C_d, sizeBytes);
|
||||
|
||||
// Initialize host data
|
||||
// Initialize host memory
|
||||
for (int i=0; i<sizeElements; i++) {
|
||||
A_h[i] = 1.618f * i;
|
||||
B_h[i] = 3.142f * i;
|
||||
}
|
||||
|
||||
// H2D Copy
|
||||
hipMemcpy(A_d, A_h, sizeBytes, hipMemcpyHostToDevice);
|
||||
hipMemcpy(B_d, B_h, sizeBytes, hipMemcpyHostToDevice);
|
||||
|
||||
// Launch kernel onto default accelerator:
|
||||
// Launch kernel onto default accelerator
|
||||
int blockSize = 256; // pick arbitrary block size
|
||||
int blocks = (sizeElements+blockSize-1)/blockSize; // round up to launch enough blocks
|
||||
hipLaunchKernel(vadd_hip, dim3(blocks), dim3(blockSize), 0, 0, A_d, B_d, C_d, sizeElements);
|
||||
|
||||
// D2H Copy
|
||||
hipMemcpy(C_h, C_d, sizeBytes, hipMemcpyDeviceToHost);
|
||||
|
||||
// Verify
|
||||
for (int i=0; i<sizeElements; i++) {
|
||||
float ref= 1.618f * i + 3.142f * i;
|
||||
if (C_h[i] != ref) {
|
||||
|
||||
@@ -3,19 +3,31 @@ ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
OPT=
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
|
||||
ifeq (${HIP_PLATFORM}, hcc)
|
||||
GENCODEOBJECT_FLAGS=--target-isa-fiji
|
||||
endif
|
||||
|
||||
all: vcpy_isa.compile runKernel.hip.out
|
||||
|
||||
vcpy_isa.compile: vcpy_isa.cpp
|
||||
$(HIPCC) --gencodeobject $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co
|
||||
all: vcpy_kernel.compile runKernel.hip.out
|
||||
|
||||
runKernel.hip.out: runKernel.cpp
|
||||
$(HIPCC) runKernel.cpp -o runKernel.hip.out
|
||||
$(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out
|
||||
|
||||
ifeq (${HIP_PLATFORM}, hcc)
|
||||
|
||||
vcpy_kernel.compile: vcpy_kernel.cpp
|
||||
$(HIPCC) --genco --target-isa=fiji vcpy_kernel.cpp -o vcpy_kernel.co
|
||||
|
||||
clean:
|
||||
rm -f *.co *.out
|
||||
|
||||
endif
|
||||
|
||||
ifeq (${HIP_PLATFORM}, nvcc)
|
||||
|
||||
vcpy_kernel.compile: vcpy_kernel.cpp
|
||||
$(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx
|
||||
|
||||
clean:
|
||||
rm -f *.ptx *.out
|
||||
|
||||
endif
|
||||
|
||||
|
||||
@@ -27,12 +27,12 @@ THE SOFTWARE.
|
||||
#define SIZE LEN<<2
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define fileName "vcpy_isa.co"
|
||||
#define kernel_name "ZN12_GLOBAL__N_146_Z11hello_world16grid_launch_parmPfS0__functor19__cxxamp_trampolineEiiiiiiPKfPf"
|
||||
#define fileName "vcpy_kernel.co"
|
||||
#define kernel_name "hello_world"
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
#define fileName "vcpy_isa.ptx"
|
||||
#define fileName "vcpy_kernel.ptx"
|
||||
#define kernel_name "hello_world"
|
||||
#endif
|
||||
|
||||
@@ -66,8 +66,9 @@ int main(){
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
|
||||
uint32_t len = LEN;
|
||||
uint32_t one = 1;
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
uint32_t len = LEN;
|
||||
uint32_t one = 1;
|
||||
|
||||
std::vector<void*>argBuffer(5);
|
||||
uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
|
||||
@@ -79,7 +80,16 @@ int main(){
|
||||
memcpy(ptr32_t + 5, &one, sizeof(uint32_t));
|
||||
memcpy(&argBuffer[3], &Ad, sizeof(void*));
|
||||
memcpy(&argBuffer[4], &Bd, sizeof(void*));
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
uint32_t one = 1;
|
||||
std::vector<void*>argBuffer(3);
|
||||
uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
|
||||
memcpy(ptr32_t + 0, &one, sizeof(uint32_t));
|
||||
memcpy(&argBuffer[1], &Ad, sizeof(void*));
|
||||
memcpy(&argBuffer[2], &Bd, sizeof(void*));
|
||||
#endif
|
||||
|
||||
|
||||
size_t size = argBuffer.size()*sizeof(void*);
|
||||
|
||||
@@ -1,9 +0,0 @@
|
||||
#include<hip_runtime.h>
|
||||
|
||||
__global__ void hello_world(hipLaunchParm lp, float *a, float *b)
|
||||
{
|
||||
int tx = hipThreadIdx_x;
|
||||
b[tx] = a[tx];
|
||||
}
|
||||
|
||||
int main(){}
|
||||
@@ -1,6 +0,0 @@
|
||||
|
||||
extern "C" __global__ void hello_world(float *a, float *b)
|
||||
{
|
||||
int tx = threadIdx.x;
|
||||
b[tx] = a[tx];
|
||||
}
|
||||
@@ -1,38 +0,0 @@
|
||||
//
|
||||
// Generated by NVIDIA NVVM Compiler
|
||||
//
|
||||
// Compiler Build ID: CL-19856038
|
||||
// Cuda compilation tools, release 7.5, V7.5.17
|
||||
// Based on LLVM 3.4svn
|
||||
//
|
||||
|
||||
.version 4.3
|
||||
.target sm_20
|
||||
.address_size 64
|
||||
|
||||
// .globl hello_world
|
||||
|
||||
.visible .entry hello_world(
|
||||
.param .u64 hello_world_param_0,
|
||||
.param .u64 hello_world_param_1
|
||||
)
|
||||
{
|
||||
.reg .f32 %f<2>;
|
||||
.reg .b32 %r<2>;
|
||||
.reg .b64 %rd<8>;
|
||||
|
||||
|
||||
ld.param.u64 %rd1, [hello_world_param_0];
|
||||
ld.param.u64 %rd2, [hello_world_param_1];
|
||||
cvta.to.global.u64 %rd3, %rd2;
|
||||
cvta.to.global.u64 %rd4, %rd1;
|
||||
mov.u32 %r1, %tid.x;
|
||||
mul.wide.s32 %rd5, %r1, 4;
|
||||
add.s64 %rd6, %rd4, %rd5;
|
||||
ld.global.f32 %f1, [%rd6];
|
||||
add.s64 %rd7, %rd3, %rd5;
|
||||
st.global.f32 [%rd7], %f1;
|
||||
ret;
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,8 @@
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
|
||||
{
|
||||
int tx = hipThreadIdx_x;
|
||||
b[tx] = a[tx];
|
||||
}
|
||||
|
||||
@@ -27,11 +27,6 @@ THE SOFTWARE.
|
||||
using namespace hc::precise_math;
|
||||
#endif
|
||||
|
||||
#if __hcc_workweek__ > 16186
|
||||
#define USE_DYNAMIC_SHARED 1
|
||||
#else
|
||||
#define USE_DYNAMIC_SHARED 0
|
||||
#endif
|
||||
|
||||
#define HIP_SQRT_2 1.41421356237
|
||||
#define HIP_SQRT_PI 1.77245385091
|
||||
@@ -774,11 +769,6 @@ __device__ float __hip_ynf(int n, float x)
|
||||
}
|
||||
|
||||
|
||||
#if __hcc_workweek__ > 16186
|
||||
#define USE_DYNAMIC_SHARED 1
|
||||
#else
|
||||
#define USE_DYNAMIC_SHARED 0
|
||||
#endif
|
||||
|
||||
__device__ float acosf(float x)
|
||||
{
|
||||
@@ -1854,12 +1844,10 @@ __host__ __device__ int max(int arg1, int arg2)
|
||||
return (int)(hc::precise_math::fmax((float)arg1, (float)arg2));
|
||||
}
|
||||
|
||||
#if USE_DYNAMIC_SHARED
|
||||
__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr()
|
||||
{
|
||||
return hc::get_dynamic_group_segment_base_pointer();
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -29,22 +29,13 @@ THE SOFTWARE.
|
||||
// Stack of contexts
|
||||
thread_local std::stack<ihipCtx_t *> tls_ctxStack;
|
||||
|
||||
hipError_t ihipCtxStackUpdate()
|
||||
void ihipCtxStackUpdate()
|
||||
{
|
||||
//HIP_INIT_API();
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
if(tls_ctxStack.empty()) {
|
||||
tls_ctxStack.push(ihipGetTlsDefaultCtx());
|
||||
tls_ctxStack.push(ihipGetTlsDefaultCtx());
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidValue
|
||||
*/
|
||||
//---
|
||||
hipError_t hipInit(unsigned int flags)
|
||||
{
|
||||
HIP_INIT_API(flags);
|
||||
@@ -59,10 +50,6 @@ hipError_t hipInit(unsigned int flags)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init
|
||||
@@ -75,10 +62,6 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDeviceGet(hipDevice_t *device, int deviceId)
|
||||
{
|
||||
HIP_INIT_API(device, deviceId); // FIXME - review if we want to init
|
||||
@@ -93,11 +76,6 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId)
|
||||
return ihipLogStatus(e);
|
||||
};
|
||||
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
HIP_INIT_API(driverVersion);
|
||||
@@ -109,10 +87,6 @@ hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidValue
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxDestroy(hipCtx_t ctx)
|
||||
{
|
||||
HIP_INIT_API(ctx);
|
||||
@@ -135,10 +109,6 @@ hipError_t hipCtxDestroy(hipCtx_t ctx)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxPopCurrent(hipCtx_t* ctx)
|
||||
{
|
||||
HIP_INIT_API(ctx);
|
||||
@@ -159,10 +129,6 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidContext
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxPushCurrent(hipCtx_t ctx)
|
||||
{
|
||||
HIP_INIT_API(ctx);
|
||||
@@ -177,10 +143,6 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxGetCurrent(hipCtx_t* ctx)
|
||||
{
|
||||
HIP_INIT_API(ctx);
|
||||
@@ -194,10 +156,6 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxSetCurrent(hipCtx_t ctx)
|
||||
{
|
||||
HIP_INIT_API(ctx);
|
||||
@@ -212,10 +170,6 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidContext
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxGetDevice(hipDevice_t *device)
|
||||
{
|
||||
HIP_INIT_API(device);
|
||||
@@ -232,10 +186,6 @@ hipError_t hipCtxGetDevice(hipDevice_t *device)
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion)
|
||||
{
|
||||
HIP_INIT_API(apiVersion);
|
||||
@@ -247,10 +197,6 @@ hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion)
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig )
|
||||
{
|
||||
HIP_INIT_API(cacheConfig);
|
||||
@@ -260,10 +206,6 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig )
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig )
|
||||
{
|
||||
HIP_INIT_API(cacheConfig);
|
||||
@@ -273,10 +215,6 @@ hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig )
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config )
|
||||
{
|
||||
HIP_INIT_API(config);
|
||||
@@ -286,10 +224,6 @@ hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config )
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
|
||||
{
|
||||
HIP_INIT_API(pConfig);
|
||||
@@ -299,20 +233,12 @@ hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxSynchronize ( void )
|
||||
{
|
||||
HIP_INIT_API(1);
|
||||
return ihipSynchronize(); //TODP Shall check validity of ctx?
|
||||
}
|
||||
|
||||
/**
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
//---
|
||||
hipError_t hipCtxGetFlags ( unsigned int* flags )
|
||||
{
|
||||
HIP_INIT_API(flags);
|
||||
|
||||
@@ -160,7 +160,7 @@ hipError_t hipSetDevice(int deviceId)
|
||||
*/
|
||||
hipError_t hipDeviceSynchronize(void)
|
||||
{
|
||||
HIP_INIT_API(1);
|
||||
HIP_INIT_API();
|
||||
return ihipSynchronize();
|
||||
}
|
||||
|
||||
@@ -324,3 +324,40 @@ hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId)
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(major,minor, device);
|
||||
hipError_t e = hipSuccess;
|
||||
int deviceId= device->_deviceId;
|
||||
e = hipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId);
|
||||
e = hipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId);
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(name,len, device);
|
||||
hipError_t e = hipSuccess;
|
||||
int nameLen = strlen(device->_props.name);
|
||||
if(nameLen <= len)
|
||||
memcpy(name,device->_props.name,nameLen);
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(pciBusId,len, device);
|
||||
hipError_t e = hipSuccess;
|
||||
int deviceId= device->_deviceId;
|
||||
e = hipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId);
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(bytes, device);
|
||||
hipError_t e = hipSuccess;
|
||||
*bytes= device->_props.totalGlobalMem;
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
@@ -40,7 +40,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
|
||||
eh->_flags = flags;
|
||||
eh->_timestamp = 0;
|
||||
eh->_copySeqId = 0;
|
||||
*event = eh;
|
||||
*event = eh; // TODO - allocat the event directly, no copy needed.
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
@@ -72,9 +72,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
|
||||
{
|
||||
HIP_INIT_API(event, stream);
|
||||
|
||||
ihipEvent_t *eh = event;
|
||||
if (eh && eh->_state != hipEventStatusUnitialized) {
|
||||
eh->_stream = stream;
|
||||
if (event && event->_state != hipEventStatusUnitialized) {
|
||||
event->_stream = stream;
|
||||
|
||||
if (stream == NULL) {
|
||||
// If stream == NULL, wait on all queues.
|
||||
@@ -83,16 +82,16 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
|
||||
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
|
||||
ctx->locked_syncDefaultStream(true);
|
||||
|
||||
eh->_timestamp = hc::get_system_ticks();
|
||||
eh->_state = hipEventStatusRecorded;
|
||||
event->_timestamp = hc::get_system_ticks();
|
||||
event->_state = hipEventStatusRecorded;
|
||||
return ihipLogStatus(hipSuccess);
|
||||
} else {
|
||||
eh->_state = hipEventStatusRecording;
|
||||
event->_state = hipEventStatusRecording;
|
||||
// Clear timestamps
|
||||
eh->_timestamp = 0;
|
||||
eh->_marker = stream->_av.create_marker();
|
||||
|
||||
eh->_copySeqId = stream->locked_lastCopySeqId();
|
||||
event->_timestamp = 0;
|
||||
|
||||
// Record the event in the stream:
|
||||
stream->locked_recordEvent(event);
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
@@ -122,21 +121,19 @@ hipError_t hipEventSynchronize(hipEvent_t event)
|
||||
{
|
||||
HIP_INIT_API(event);
|
||||
|
||||
ihipEvent_t *eh = event;
|
||||
|
||||
if (eh) {
|
||||
if (eh->_state == hipEventStatusUnitialized) {
|
||||
if (event) {
|
||||
if (event->_state == hipEventStatusUnitialized) {
|
||||
return ihipLogStatus(hipErrorInvalidResourceHandle);
|
||||
} else if (eh->_state == hipEventStatusCreated ) {
|
||||
} else if (event->_state == hipEventStatusCreated ) {
|
||||
// Created but not actually recorded on any device:
|
||||
return ihipLogStatus(hipSuccess);
|
||||
} else if (eh->_stream == NULL) {
|
||||
} else if (event->_stream == NULL) {
|
||||
auto *ctx = ihipGetTlsDefaultCtx();
|
||||
ctx->locked_syncDefaultStream(true);
|
||||
return ihipLogStatus(hipSuccess);
|
||||
} else {
|
||||
eh->_marker.wait((eh->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
|
||||
eh->_stream->locked_reclaimSignals(eh->_copySeqId);
|
||||
event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
|
||||
event->_stream->locked_reclaimSignals(event->_copySeqId);
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
}
|
||||
@@ -196,12 +193,11 @@ hipError_t hipEventQuery(hipEvent_t event)
|
||||
{
|
||||
HIP_INIT_API(event);
|
||||
|
||||
ihipEvent_t *eh = event;
|
||||
|
||||
// TODO-stream - need to read state of signal here: The event may have become ready after recording..
|
||||
// TODO-HCC - use get_hsa_signal here.
|
||||
|
||||
if (eh->_state == hipEventStatusRecording) {
|
||||
if (event->_state == hipEventStatusRecording) {
|
||||
return ihipLogStatus(hipErrorNotReady);
|
||||
} else {
|
||||
return ihipLogStatus(hipSuccess);
|
||||
|
||||
@@ -137,7 +137,7 @@ ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex)
|
||||
};
|
||||
|
||||
|
||||
static thread_local ihipCtx_t *tls_defaultCtx = nullptr;
|
||||
static thread_local ihipCtx_t *tls_defaultCtx = nullptr;
|
||||
void ihipSetTlsDefaultCtx(ihipCtx_t *ctx)
|
||||
{
|
||||
tls_defaultCtx = ctx;
|
||||
@@ -162,7 +162,7 @@ hipError_t ihipSynchronize(void)
|
||||
{
|
||||
ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
return (hipSuccess);
|
||||
}
|
||||
|
||||
//=================================================================================================
|
||||
@@ -195,9 +195,9 @@ ihipSignal_t::~ihipSignal_t()
|
||||
//---
|
||||
ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags) :
|
||||
_id(0), // will be set by add function.
|
||||
_av(av),
|
||||
_flags(flags),
|
||||
_ctx(ctx)
|
||||
_ctx(ctx),
|
||||
_criticalData(av)
|
||||
{
|
||||
tprintf(DB_SYNC, " streamCreate: stream=%p\n", this);
|
||||
};
|
||||
@@ -246,7 +246,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
|
||||
{
|
||||
if (! assertQueueEmpty) {
|
||||
tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this);
|
||||
_av.wait();
|
||||
crit->_av.wait();
|
||||
}
|
||||
|
||||
if (crit->_last_copy_signal) {
|
||||
@@ -273,6 +273,28 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty)
|
||||
|
||||
};
|
||||
|
||||
#if USE_AV_COPY
|
||||
// Causes current stream to wait for specified event to complete:
|
||||
void ihipStream_t::locked_waitEvent(hipEvent_t event)
|
||||
{
|
||||
LockedAccessor_StreamCrit_t crit(_criticalData);
|
||||
|
||||
// TODO - check state of event here:
|
||||
crit->_av.create_blocking_marker(event->_marker);
|
||||
}
|
||||
#endif
|
||||
|
||||
// Create a marker in this stream.
|
||||
// Save state in the event so it can track the status of the event.
|
||||
void ihipStream_t::locked_recordEvent(hipEvent_t event)
|
||||
{
|
||||
// Lock the stream to prevent simultaneous access
|
||||
LockedAccessor_StreamCrit_t crit(_criticalData);
|
||||
|
||||
event->_marker = crit->_av.create_marker();
|
||||
event->_copySeqId = lastCopySeqId(crit);
|
||||
}
|
||||
|
||||
//=============================================================================
|
||||
|
||||
|
||||
@@ -388,11 +410,10 @@ int HIP_NUM_KERNELS_INFLIGHT = 128;
|
||||
//into the stream to mimic CUDA stream semantics. (some hardware uses separate
|
||||
//queues for data commands and kernel commands, and no implicit ordering is provided).
|
||||
//
|
||||
bool ihipStream_t::lockopen_preKernelCommand()
|
||||
LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
|
||||
{
|
||||
LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/);
|
||||
|
||||
bool addedSync = false;
|
||||
|
||||
if(crit->_kernelCnt > HIP_NUM_KERNELS_INFLIGHT){
|
||||
this->wait(crit);
|
||||
@@ -402,9 +423,8 @@ bool ihipStream_t::lockopen_preKernelCommand()
|
||||
// If switching command types, we need to add a barrier packet to synchronize things.
|
||||
if (crit->_last_command_type != ihipCommandKernel) {
|
||||
if (crit->_last_copy_signal) {
|
||||
addedSync = true;
|
||||
|
||||
hsa_queue_t * q = (hsa_queue_t*)_av.get_hsa_queue();
|
||||
hsa_queue_t * q = (hsa_queue_t*) (crit->_av.get_hsa_queue());
|
||||
if (HIP_DISABLE_HW_KERNEL_DEP == 0) {
|
||||
this->enqueueBarrier(q, crit->_last_copy_signal, NULL);
|
||||
tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n",
|
||||
@@ -422,7 +442,7 @@ bool ihipStream_t::lockopen_preKernelCommand()
|
||||
crit->_last_command_type = ihipCommandKernel;
|
||||
}
|
||||
|
||||
return addedSync;
|
||||
return crit;
|
||||
}
|
||||
|
||||
|
||||
@@ -433,6 +453,11 @@ void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFutu
|
||||
// We locked _criticalData in the lockopen_preKernelCommand() so OK to access here:
|
||||
_criticalData._last_kernel_future = kernelFuture;
|
||||
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
kernelFuture.wait();
|
||||
tprintf(DB_SYNC, " %s LAUNCH_BLOCKING for kernel completion\n", ToString(this).c_str());
|
||||
}
|
||||
|
||||
_criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
|
||||
};
|
||||
|
||||
@@ -457,7 +482,7 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
|
||||
needSync = 1;
|
||||
ihipSignal_t *depSignal = allocSignal(crit);
|
||||
hsa_signal_store_relaxed(depSignal->_hsaSignal,1);
|
||||
this->enqueueBarrier(static_cast<hsa_queue_t*>(_av.get_hsa_queue()), NULL, depSignal);
|
||||
this->enqueueBarrier(static_cast<hsa_queue_t*>(crit->_av.get_hsa_queue()), NULL, depSignal);
|
||||
*waitSignal = depSignal->_hsaSignal;
|
||||
} else if (crit->_last_copy_signal) {
|
||||
needSync = 1;
|
||||
@@ -487,7 +512,10 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
|
||||
}
|
||||
|
||||
|
||||
void ihipStream_t::launchModuleKernel(hsa_signal_t signal,
|
||||
// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av.
|
||||
void ihipStream_t::launchModuleKernel(
|
||||
hc::accelerator_view av,
|
||||
hsa_signal_t signal,
|
||||
uint32_t blockDimX,
|
||||
uint32_t blockDimY,
|
||||
uint32_t blockDimZ,
|
||||
@@ -500,11 +528,12 @@ void ihipStream_t::launchModuleKernel(hsa_signal_t signal,
|
||||
uint64_t kernel){
|
||||
hsa_status_t status;
|
||||
void *kern;
|
||||
hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(_av.get_hsa_kernarg_region());
|
||||
|
||||
hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(av.get_hsa_kernarg_region());
|
||||
status = hsa_amd_memory_pool_allocate(*pool, kernSize, 0, &kern);
|
||||
status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)_av.get_hsa_agent(), 0, kern);
|
||||
status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)av.get_hsa_agent(), 0, kern);
|
||||
memcpy(kern, kernarg, kernSize);
|
||||
hsa_queue_t *Queue = (hsa_queue_t*)_av.get_hsa_queue();
|
||||
hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue();
|
||||
const uint32_t queue_mask = Queue->size-1;
|
||||
uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue);
|
||||
hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]);
|
||||
@@ -1117,8 +1146,9 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c
|
||||
while (std::getline(ss, device_id, ',')) {
|
||||
if (atoi(device_id.c_str()) >= 0) {
|
||||
g_hip_visible_devices.push_back(atoi(device_id.c_str()));
|
||||
}else// Any device number after invalid number will not present
|
||||
} else { // Any device number after invalid number will not present
|
||||
break;
|
||||
}
|
||||
}
|
||||
// Print out the number of ids
|
||||
if (HIP_PRINT_ENV) {
|
||||
@@ -1312,11 +1342,11 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
|
||||
}
|
||||
}
|
||||
|
||||
void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
|
||||
void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
|
||||
{
|
||||
std::string streamString = ToString(stream);
|
||||
fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d %s\n" KNRM, \
|
||||
kernelName, lp->grid_dim.x, lp->grid_dim.y, lp->grid_dim.z, lp->group_dim.x, lp->group_dim.y, lp->group_dim.z,
|
||||
kernelName, lp->grid_dim.x, lp->grid_dim.y, lp->grid_dim.z, lp->group_dim.x, lp->group_dim.y, lp->group_dim.z,
|
||||
lp->dynamic_group_mem_bytes, streamString.c_str());\
|
||||
}
|
||||
|
||||
@@ -1327,7 +1357,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
|
||||
{
|
||||
HIP_INIT();
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
#if USE_GRID_LAUNCH_20
|
||||
lp->grid_dim.x = grid.x;
|
||||
lp->grid_dim.y = grid.y;
|
||||
lp->grid_dim.z = grid.z;
|
||||
@@ -1336,27 +1365,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
|
||||
lp->group_dim.z = block.z;
|
||||
lp->barrier_bit = barrier_bit_queue_default;
|
||||
lp->launch_fence = -1;
|
||||
#else
|
||||
lp->gridDim.x = grid.x;
|
||||
lp->gridDim.y = grid.y;
|
||||
lp->gridDim.z = grid.z;
|
||||
lp->groupDim.x = block.x;
|
||||
lp->groupDim.y = block.y;
|
||||
lp->groupDim.z = block.z;
|
||||
#endif
|
||||
stream->lockopen_preKernelCommand();
|
||||
// *av = &stream->_av;
|
||||
lp->av = &stream->_av;
|
||||
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
lp->av = &(crit->_av);
|
||||
lp->cf = new hc::completion_future;
|
||||
// lp->av = static_cast<void*>(av);
|
||||
// lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
|
||||
return (stream);
|
||||
}
|
||||
|
||||
|
||||
hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp)
|
||||
{
|
||||
HIP_INIT_API(stream, grid, block, lp);
|
||||
HIP_INIT();
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
#if USE_GRID_LAUNCH_20
|
||||
lp->grid_dim.x = grid;
|
||||
lp->grid_dim.y = 1;
|
||||
lp->grid_dim.z = 1;
|
||||
@@ -1365,28 +1385,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri
|
||||
lp->group_dim.z = block.z;
|
||||
lp->barrier_bit = barrier_bit_queue_default;
|
||||
lp->launch_fence = -1;
|
||||
#else
|
||||
lp->gridDim.x = grid;
|
||||
lp->gridDim.y = 1;
|
||||
lp->gridDim.z = 1;
|
||||
lp->groupDim.x = block.x;
|
||||
lp->groupDim.y = block.y;
|
||||
lp->groupDim.z = block.z;
|
||||
#endif
|
||||
stream->lockopen_preKernelCommand();
|
||||
// *av = &stream->_av;
|
||||
lp->av = &stream->_av;
|
||||
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
lp->av = &(crit->_av);
|
||||
lp->cf = new hc::completion_future;
|
||||
// lp->av = static_cast<void*>(av);
|
||||
// lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
|
||||
return (stream);
|
||||
}
|
||||
|
||||
|
||||
hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp)
|
||||
{
|
||||
HIP_INIT_API(stream, grid, block, lp);
|
||||
HIP_INIT();
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
#if USE_GRID_LAUNCH_20
|
||||
lp->grid_dim.x = grid.x;
|
||||
lp->grid_dim.y = grid.y;
|
||||
lp->grid_dim.z = grid.z;
|
||||
@@ -1395,28 +1405,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri
|
||||
lp->group_dim.z = 1;
|
||||
lp->barrier_bit = barrier_bit_queue_default;
|
||||
lp->launch_fence = -1;
|
||||
#else
|
||||
lp->gridDim.x = grid.x;
|
||||
lp->gridDim.y = grid.y;
|
||||
lp->gridDim.z = grid.z;
|
||||
lp->groupDim.x = block;
|
||||
lp->groupDim.y = 1;
|
||||
lp->groupDim.z = 1;
|
||||
#endif
|
||||
stream->lockopen_preKernelCommand();
|
||||
// *av = &stream->_av;
|
||||
lp->av = &stream->_av;
|
||||
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
lp->av = &(crit->_av);
|
||||
lp->cf = new hc::completion_future;
|
||||
// lp->av = static_cast<void*>(av);
|
||||
// lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
|
||||
return (stream);
|
||||
}
|
||||
|
||||
|
||||
hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp)
|
||||
{
|
||||
HIP_INIT_API(stream, grid, block, lp);
|
||||
HIP_INIT();
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
#if USE_GRID_LAUNCH_20
|
||||
lp->grid_dim.x = grid;
|
||||
lp->grid_dim.y = 1;
|
||||
lp->grid_dim.z = 1;
|
||||
@@ -1425,37 +1425,23 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
|
||||
lp->group_dim.z = 1;
|
||||
lp->barrier_bit = barrier_bit_queue_default;
|
||||
lp->launch_fence = -1;
|
||||
#else
|
||||
lp->gridDim.x = grid;
|
||||
lp->gridDim.y = 1;
|
||||
lp->gridDim.z = 1;
|
||||
lp->groupDim.x = block;
|
||||
lp->groupDim.y = 1;
|
||||
lp->groupDim.z = 1;
|
||||
#endif
|
||||
stream->lockopen_preKernelCommand();
|
||||
// *av = &stream->_av;
|
||||
lp->av = &stream->_av;
|
||||
lp->cf = new hc::completion_future;
|
||||
// lp->av = static_cast<void*>(av);
|
||||
// lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
|
||||
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
lp->av = &(crit->_av);
|
||||
lp->cf = new hc::completion_future; // TODO, is this necessary?
|
||||
return (stream);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Called after kernel finishes execution.
|
||||
//This releases the lock on the stream.
|
||||
void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
|
||||
{
|
||||
// stream->lockclose_postKernelCommand(cf);
|
||||
stream->lockclose_postKernelCommand(*lp.cf);
|
||||
if (HIP_LAUNCH_BLOCKING) {
|
||||
tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream);
|
||||
}
|
||||
stream->lockclose_postKernelCommand(*(lp.cf));
|
||||
}
|
||||
|
||||
|
||||
//
|
||||
//=================================================================================================
|
||||
// HIP API Implementation
|
||||
//
|
||||
@@ -1629,7 +1615,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src,dst,sizeBytes);
|
||||
crit->_av.copy(src,dst,sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
@@ -1677,7 +1663,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
|
||||
// TODO - remove, slow path.
|
||||
tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
|
||||
#if USE_AV_COPY
|
||||
_av.copy(src, dst, sizeBytes);
|
||||
crit->_av.copy(src, dst, sizeBytes);
|
||||
#else
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
#endif
|
||||
@@ -1879,7 +1865,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a
|
||||
stream = device->_defaultStream;
|
||||
}
|
||||
|
||||
*av = &(stream->_av);
|
||||
*av = stream->locked_getAv();
|
||||
|
||||
hipError_t err = hipSuccess;
|
||||
return ihipLogStatus(err);
|
||||
|
||||
@@ -119,7 +119,13 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
|
||||
HIP_INIT_API(ptr, sizeBytes);
|
||||
|
||||
hipError_t hip_status = hipSuccess;
|
||||
|
||||
// return NULL pointer when malloc size is 0
|
||||
if (sizeBytes == 0)
|
||||
{
|
||||
*ptr = NULL;
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
|
||||
if (ctx) {
|
||||
@@ -142,6 +148,8 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
}
|
||||
|
||||
//printf (" hipMalloc allocated %p\n", *ptr);
|
||||
|
||||
return ihipLogStatus(hip_status);
|
||||
}
|
||||
|
||||
@@ -672,9 +680,12 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
|
||||
|
||||
|
||||
|
||||
// TODO - make member function of stream?
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes)
|
||||
ihipMemsetKernel(hipStream_t stream,
|
||||
LockedAccessor_StreamCrit_t &crit,
|
||||
T * ptr, T val, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_computeUnits);
|
||||
const int threads_per_wg = 256;
|
||||
@@ -690,7 +701,7 @@ ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes)
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
crit->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
@@ -707,41 +718,6 @@ ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes)
|
||||
return cf;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
hc::completion_future
|
||||
ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes)
|
||||
{
|
||||
int wg = std::min((unsigned)8, stream->getDevice()->_computeUnits);
|
||||
const int threads_per_wg = 256;
|
||||
|
||||
int threads = wg * threads_per_wg;
|
||||
if (threads > sizeBytes) {
|
||||
threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;
|
||||
}
|
||||
|
||||
|
||||
hc::extent<1> ext(threads);
|
||||
auto ext_tile = ext.tile(threads_per_wg);
|
||||
|
||||
hc::completion_future cf =
|
||||
hc::parallel_for_each(
|
||||
stream->_av,
|
||||
ext_tile,
|
||||
[=] (hc::tiled_index<1> idx)
|
||||
__attribute__((hc))
|
||||
{
|
||||
int offset = amp_get_global_id(0);
|
||||
// TODO-HCC - change to hc_get_local_size()
|
||||
int stride = amp_get_local_size(0) * hc_get_num_groups(0) ;
|
||||
|
||||
for (int i=offset; i<sizeBytes; i+=stride) {
|
||||
c[i] = a[i];
|
||||
}
|
||||
});
|
||||
|
||||
return cf;
|
||||
}
|
||||
|
||||
|
||||
|
||||
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
|
||||
@@ -756,7 +732,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (stream) {
|
||||
stream->lockopen_preKernelCommand();
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
|
||||
hc::completion_future cf ;
|
||||
|
||||
@@ -765,7 +741,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
|
||||
try {
|
||||
value = value & 0xff;
|
||||
unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
|
||||
cf = ihipMemsetKernel<unsigned> (stream, static_cast<unsigned*> (dst), value32, sizeBytes/sizeof(unsigned));
|
||||
cf = ihipMemsetKernel<unsigned> (stream, crit, static_cast<unsigned*> (dst), value32, sizeBytes/sizeof(unsigned));
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
e = hipErrorInvalidValue;
|
||||
@@ -773,7 +749,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
|
||||
} else {
|
||||
// use a slow byte-per-workitem copy:
|
||||
try {
|
||||
cf = ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, sizeBytes);
|
||||
cf = ihipMemsetKernel<char> (stream, crit, static_cast<char*> (dst), value, sizeBytes);
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
e = hipErrorInvalidValue;
|
||||
@@ -808,7 +784,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (stream) {
|
||||
stream->lockopen_preKernelCommand();
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
|
||||
hc::completion_future cf ;
|
||||
|
||||
@@ -817,7 +793,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
|
||||
try {
|
||||
value = value & 0xff;
|
||||
unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
|
||||
cf = ihipMemsetKernel<unsigned> (stream, static_cast<unsigned*> (dst), value32, sizeBytes/sizeof(unsigned));
|
||||
cf = ihipMemsetKernel<unsigned> (stream, crit, static_cast<unsigned*> (dst), value32, sizeBytes/sizeof(unsigned));
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
e = hipErrorInvalidValue;
|
||||
@@ -825,7 +801,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
|
||||
} else {
|
||||
// use a slow byte-per-workitem copy:
|
||||
try {
|
||||
cf = ihipMemsetKernel<char> (stream, static_cast<char*> (dst), value, sizeBytes);
|
||||
cf = ihipMemsetKernel<char> (stream, crit, static_cast<char*> (dst), value, sizeBytes);
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
e = hipErrorInvalidValue;
|
||||
|
||||
@@ -104,7 +104,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
*module = new ihipModule_t;
|
||||
|
||||
if(module == NULL){
|
||||
return hipErrorInvalidValue;
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
@@ -117,7 +117,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
std::ifstream in(fname, std::ios::binary | std::ios::ate);
|
||||
|
||||
if(!in){
|
||||
return hipErrorFileNotFound;
|
||||
return ihipLogStatus(hipErrorFileNotFound);
|
||||
|
||||
}else{
|
||||
|
||||
@@ -130,12 +130,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
status = hsa_memory_allocate(sysRegion, size, (void**)&p);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorOutOfMemory;
|
||||
return ihipLogStatus(hipErrorOutOfMemory);
|
||||
}
|
||||
|
||||
char *ptr = (char*)p;
|
||||
if(!ptr){
|
||||
return hipErrorOutOfMemory;
|
||||
return ihipLogStatus(hipErrorOutOfMemory);
|
||||
}
|
||||
(*module)->ptr = p;
|
||||
(*module)->size = size;
|
||||
@@ -146,27 +146,33 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorSharedObjectInitFailed;
|
||||
return ihipLogStatus(hipErrorSharedObjectInitFailed);
|
||||
}
|
||||
|
||||
status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorNotInitialized;
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return ret;
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
hipError_t hipModuleUnload(hipModule_t hmod){
|
||||
hipError_t ret = hipSuccess;
|
||||
hsa_status_t status = hsa_executable_destroy(hmod->executable);
|
||||
if(status != HSA_STATUS_SUCCESS){ret = hipErrorInvalidValue; }
|
||||
if(status != HSA_STATUS_SUCCESS)
|
||||
{
|
||||
ret = hipErrorInvalidValue;
|
||||
}
|
||||
status = hsa_code_object_destroy(hmod->object);
|
||||
if(status != HSA_STATUS_SUCCESS){ret = hipErrorInvalidValue; }
|
||||
if(status != HSA_STATUS_SUCCESS)
|
||||
{
|
||||
ret = hipErrorInvalidValue;
|
||||
}
|
||||
delete hmod;
|
||||
return ret;
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){
|
||||
@@ -174,7 +180,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
|
||||
hipError_t ret = hipSuccess;
|
||||
|
||||
if(name == nullptr){
|
||||
return hipErrorInvalidValue;
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
if(ctx == nullptr){
|
||||
@@ -189,13 +195,13 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
|
||||
hsa_status_t status;
|
||||
status = hsa_executable_load_code_object(hmod->executable, gpuAgent, hmod->object, NULL);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorNotInitialized;
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
|
||||
status = hsa_executable_freeze(hmod->executable, NULL);
|
||||
status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(*func)->kernel_symbol);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorNotFound;
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
|
||||
status = hsa_executable_symbol_get_info((*func)->kernel_symbol,
|
||||
@@ -203,10 +209,10 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
|
||||
&(*func)->kernel);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorNotFound;
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
}
|
||||
return ret;
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod,
|
||||
@@ -215,6 +221,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod,
|
||||
return ihipModuleGetFunction(hfunc, hmod, name);
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
|
||||
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
|
||||
@@ -240,10 +247,10 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){
|
||||
kernSize = *(size_t*)(config[3]);
|
||||
}else{
|
||||
return hipErrorNotInitialized;
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
}else{
|
||||
return hipErrorInvalidValue;
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
/*
|
||||
Kernel argument preparation.
|
||||
@@ -262,7 +269,7 @@ Kernel argument preparation.
|
||||
/*
|
||||
Launch AQL packet
|
||||
*/
|
||||
hStream->launchModuleKernel(signal, blockDimX, blockDimY, blockDimZ,
|
||||
hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
|
||||
gridDimX, gridDimY, gridDimZ, sharedMemBytes, config[1], kernSize, f->kernel);
|
||||
|
||||
/*
|
||||
@@ -273,10 +280,10 @@ Kernel argument preparation.
|
||||
|
||||
|
||||
ihipPostLaunchKernel(hStream, lp);
|
||||
|
||||
|
||||
}
|
||||
|
||||
return ret;
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
|
||||
@@ -285,17 +292,17 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes,
|
||||
HIP_INIT_API(name);
|
||||
hipError_t ret = hipSuccess;
|
||||
if(dptr == NULL || bytes == NULL){
|
||||
return hipErrorInvalidValue;
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
if(name == NULL || hmod == NULL){
|
||||
return hipErrorNotInitialized;
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
else{
|
||||
hipFunction_t func;
|
||||
ihipModuleGetFunction(&func, hmod, name);
|
||||
*bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t);
|
||||
*dptr = reinterpret_cast<void*>(func->kernel);
|
||||
return ret;
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -303,7 +310,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image){
|
||||
HIP_INIT_API(image);
|
||||
hipError_t ret = hipSuccess;
|
||||
if(image == NULL || module == NULL){
|
||||
return hipErrorNotInitialized;
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}else{
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
*module = new ihipModule_t;
|
||||
@@ -318,12 +325,12 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image){
|
||||
status = hsa_memory_allocate(sysRegion, size, (void**)&p);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorOutOfMemory;
|
||||
return ihipLogStatus(hipErrorOutOfMemory);
|
||||
}
|
||||
|
||||
char *ptr = (char*)p;
|
||||
if(!ptr){
|
||||
return hipErrorOutOfMemory;
|
||||
return ihipLogStatus(hipErrorOutOfMemory);
|
||||
}
|
||||
(*module)->ptr = p;
|
||||
(*module)->size = size;
|
||||
@@ -333,15 +340,15 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image){
|
||||
status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorSharedObjectInitFailed;
|
||||
return ihipLogStatus(hipErrorSharedObjectInitFailed);
|
||||
}
|
||||
|
||||
status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return hipErrorNotInitialized;
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
}
|
||||
return ret;
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -66,7 +66,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
|
||||
}
|
||||
|
||||
//---
|
||||
hipError_t hipStreamCreate(hipStream_t *stream)
|
||||
hipError_t hipStreamCreate(hipStream_t *stream)
|
||||
{
|
||||
HIP_INIT_API(stream);
|
||||
|
||||
@@ -74,23 +74,42 @@ hipError_t hipStreamCreate(hipStream_t *stream)
|
||||
}
|
||||
|
||||
|
||||
#if USE_AV_COPY==0
|
||||
//---
|
||||
/**
|
||||
* @bug This function conservatively waits for all work in the specified stream to complete.
|
||||
*/
|
||||
#endif
|
||||
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
|
||||
{
|
||||
HIP_INIT_API(stream, event, flags);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
{
|
||||
// TODO-hcc Convert to use create_blocking_marker(...) functionality.
|
||||
// Currently we have a super-conservative version of this - block on host, and drain the queue.
|
||||
// This should create a barrier packet in the target queue.
|
||||
stream->locked_wait();
|
||||
e = hipSuccess;
|
||||
}
|
||||
if (event == nullptr) {
|
||||
e = hipErrorInvalidResourceHandle;
|
||||
|
||||
} else if (event->_state != hipEventStatusUnitialized) {
|
||||
|
||||
bool fastWait = false;
|
||||
|
||||
#if USE_AV_COPY
|
||||
if (stream != hipStreamNull) {
|
||||
stream->locked_waitEvent(event);
|
||||
|
||||
fastWait = true; // don't use the slow host-side synchronization.
|
||||
}
|
||||
// TODO - clean up if/else logic when USE_AV_COPY enabled.
|
||||
#endif
|
||||
|
||||
if (!fastWait) {
|
||||
// TODO-hcc Convert to use create_blocking_marker(...) functionality.
|
||||
// Currently we have a super-conservative version of this - block on host, and drain the queue.
|
||||
// This should create a barrier packet in the target queue.
|
||||
stream->locked_wait();
|
||||
e = hipSuccess;
|
||||
}
|
||||
} // else event not recorded, return immediately and don't create marker.
|
||||
|
||||
return ihipLogStatus(e);
|
||||
};
|
||||
|
||||
@@ -50,3 +50,12 @@ ctest
|
||||
ctest -R Memcpy
|
||||
```
|
||||
|
||||
|
||||
### If a test fails:
|
||||
|
||||
Extract the commandline from the testing log:
|
||||
$ grep -A3 -m2 hipMemcpy-size Testing/Temporary/LastTest.log
|
||||
36/47 Testing: hipMemcpy-size
|
||||
36/47 Test: hipMemcpy-size
|
||||
Command: "/home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory/hipMemcpy" "--tests" "0x6"
|
||||
Directory: /home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory
|
||||
|
||||
@@ -166,7 +166,8 @@ build_hip_executable (hipEventRecord hipEventRecord.cpp)
|
||||
build_hip_executable_libcpp (hipHcc hipHcc.cpp)
|
||||
#set_source_files_properties (hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++ )
|
||||
|
||||
build_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp)
|
||||
# __workweek fix.
|
||||
#build_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp)
|
||||
build_hip_executable (hipHostAlloc hipHostAlloc.cpp)
|
||||
build_hip_executable (hipHostGetFlags hipHostGetFlags.cpp)
|
||||
build_hip_executable (hipHostRegister hipHostRegister.cpp)
|
||||
|
||||
@@ -23,6 +23,8 @@ THE SOFTWARE.
|
||||
#include<fstream>
|
||||
#include<vector>
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN<<2
|
||||
|
||||
@@ -45,17 +47,17 @@ int main(){
|
||||
std::cout<<A[i] << " "<<B[i]<<std::endl;
|
||||
}
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
|
||||
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
|
||||
|
||||
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
|
||||
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
|
||||
HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
HIPCHECK(hipModuleLoad(&Module, fileName));
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
hipStream_t stream;
|
||||
hipStreamCreate(&stream);
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
void *args[2] = {&Ad, &Bd};
|
||||
|
||||
|
||||
@@ -73,9 +75,9 @@ int main(){
|
||||
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config);
|
||||
|
||||
hipStreamDestroy(stream);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
|
||||
HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
|
||||
|
||||
for(uint32_t i=0;i<LEN;i++){
|
||||
std::cout<<A[i]<<" - "<<B[i]<<std::endl;
|
||||
@@ -83,3 +85,4 @@ int main(){
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
@@ -28,11 +28,12 @@ THE SOFTWARE.
|
||||
#include <vector>
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
//#include "hcc_detail/AM.h"
|
||||
#include "hc_am.hpp"
|
||||
|
||||
#endif
|
||||
|
||||
#define USE_AV_COPY (__hcc_workweek__ >= 16351)
|
||||
|
||||
size_t Nbytes = 0;
|
||||
|
||||
//=================================================================================================
|
||||
@@ -409,11 +410,21 @@ void thread_noise_generator(int iters, size_t numBuffers, Dir addDir, Dir remove
|
||||
|
||||
if (addDir == Up) {
|
||||
for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) {
|
||||
#if USE_AV_COPY
|
||||
hc::AmPointerInfo info(p, p, bufferSize, acc, false, false);
|
||||
hc::am_memtracker_add(p, info);
|
||||
#else
|
||||
hc::am_memtracker_add(p, bufferSize, acc, false);
|
||||
#endif
|
||||
}
|
||||
} else if (addDir == Down) {
|
||||
for (char *p = basePtr+maxSize-bufferSize; p>=0; p-=bufferSize) {
|
||||
#if USE_AV_COPY
|
||||
hc::AmPointerInfo info(p, p, bufferSize, acc, false, false);
|
||||
hc::am_memtracker_add(p, info);
|
||||
#else
|
||||
hc::am_memtracker_add(p, bufferSize, acc, false);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -330,7 +330,7 @@ int main(int argc, char *argv[])
|
||||
parseMyArguments(argc, argv);
|
||||
|
||||
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
printf ("info: set device to %d tests=%x\n", p_gpuDevice, p_tests);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
if (p_tests & 0x01) {
|
||||
|
||||
@@ -21,7 +21,7 @@ THE SOFTWARE.
|
||||
#define SIZE 1024*1024*256
|
||||
|
||||
int main(){
|
||||
float *Ad, *B, *Bd, *Bm, *C, *Cd;
|
||||
float *Ad, *B, *Bd, *Bm, *C, *Cd, *ptr_0;
|
||||
B = (float*)malloc(SIZE);
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipHostMalloc((void**)&B, SIZE);
|
||||
@@ -31,12 +31,15 @@ int main(){
|
||||
|
||||
hipHostGetDevicePointer((void**)&Cd, C, 0/*flags*/);
|
||||
|
||||
HIPCHECK_API(hipMalloc((void**)&ptr_0,0), hipSuccess);
|
||||
|
||||
HIPCHECK_API(hipFree(Ad) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(Ad) , hipErrorInvalidValue);
|
||||
|
||||
HIPCHECK_API(hipFree(B) , hipErrorInvalidDevicePointer); // try to hipFree on malloced memory
|
||||
HIPCHECK_API(hipFree(Bd) , hipErrorInvalidDevicePointer);
|
||||
HIPCHECK_API(hipFree(Bm) , hipErrorInvalidDevicePointer);
|
||||
HIPCHECK_API(hipFree(ptr_0) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(Bd) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(Bm) , hipSuccess);
|
||||
|
||||
|
||||
@@ -6,8 +6,10 @@ include_directories( ${HIPTEST_SOURCE_DIR} )
|
||||
build_hip_executable (hipAPIStreamEnable hipAPIStreamEnable.cpp)
|
||||
build_hip_executable (hipAPIStreamDisable hipAPIStreamDisable.cpp)
|
||||
build_hip_executable (hipStreamL5 hipStreamL5.cpp)
|
||||
build_hip_executable (hipStreamWaitEvent hipStreamWaitEvent.cpp)
|
||||
|
||||
# TODO - seg fault
|
||||
#make_test(hipAPIStreamEnable " ")
|
||||
#make_test(hipAPIStreamDisable " ")
|
||||
make_test(hipStreamL5 " ")
|
||||
make_test(hipStreamWaitEvent " ")
|
||||
|
||||
@@ -0,0 +1,123 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// Test under-development. Calls async mem-copy API, experiment with functionality.
|
||||
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
#include <vector>
|
||||
unsigned p_streams = 6;
|
||||
|
||||
|
||||
//------
|
||||
// Structure for one stream;
|
||||
template <typename T>
|
||||
class Streamer {
|
||||
public:
|
||||
Streamer(size_t numElements);
|
||||
~Streamer();
|
||||
void runAsync();
|
||||
void waitComplete();
|
||||
|
||||
private:
|
||||
T *_A_h;
|
||||
T *_B_h;
|
||||
T *_C_h;
|
||||
|
||||
T *_A_d;
|
||||
T *_B_d;
|
||||
T *_C_d;
|
||||
|
||||
hipStream_t _stream;
|
||||
hipEvent_t _event;
|
||||
|
||||
size_t _numElements;
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
Streamer<T>::Streamer(size_t numElements) :
|
||||
_numElements(numElements)
|
||||
{
|
||||
HipTest::initArrays (&_A_d, &_B_d, &_C_d, &_A_h, &_B_h, &_C_h, numElements, true);
|
||||
|
||||
HIPCHECK(hipStreamCreate(&_stream));
|
||||
HIPCHECK(hipEventCreate(&_event));
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
void Streamer<T>::runAsync()
|
||||
{
|
||||
printf ("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0);
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, _stream, _A_d, _B_d, _C_d, _numElements);
|
||||
HIPCHECK(hipEventRecord(_event, _stream));
|
||||
HIPCHECK(hipStreamWaitEvent(_stream, _event, 0));
|
||||
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
//---
|
||||
//Parse arguments specific to this test.
|
||||
void parseMyArguments(int argc, char *argv[])
|
||||
{
|
||||
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
|
||||
|
||||
// parse args for this test:
|
||||
for (int i = 1; i < more_argc; i++) {
|
||||
const char *arg = argv[i];
|
||||
|
||||
if (!strcmp(arg, "--streams")) {
|
||||
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) {
|
||||
failed("Bad streams argument");
|
||||
}
|
||||
} else {
|
||||
failed("Bad argument '%s'", arg);
|
||||
}
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
|
||||
//---
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
parseMyArguments(argc, argv);
|
||||
|
||||
typedef Streamer<float> FloatStreamer;
|
||||
|
||||
std::vector<FloatStreamer *> streamers;
|
||||
|
||||
size_t numElements = N;
|
||||
|
||||
for (int i=0; i<p_streams; i++) {
|
||||
FloatStreamer * s = new FloatStreamer(numElements);
|
||||
streamers.push_back(s);
|
||||
}
|
||||
|
||||
for (int i=0; i<p_streams; i++) {
|
||||
streamers[i]->runAsync();
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -68,8 +68,9 @@ THE SOFTWARE.
|
||||
{\
|
||||
hipError_t localError = error; \
|
||||
if (localError != hipSuccess) { \
|
||||
printf("%serror: '%s'(%d) at %s:%d%s\n", \
|
||||
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
|
||||
KRED,hipGetErrorString(localError), localError,\
|
||||
#error,\
|
||||
__FILE__, __LINE__,KNRM); \
|
||||
failed("API returned error code.");\
|
||||
}\
|
||||
|
||||
@@ -65,6 +65,7 @@ syn keyword hipFunctionName expf __expf exp logf __logf log
|
||||
syn keyword hipType hipDeviceProp_t
|
||||
syn keyword hipType hipError_t
|
||||
syn keyword hipType hipStream_t
|
||||
syn keyword hipType hipEvent_t
|
||||
|
||||
" Runtime functions
|
||||
syn keyword hipFunctionName hipBindTexture hipBindTextureToArray
|
||||
|
||||
Ссылка в новой задаче
Block a user