AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

Merge branch 'amd-develop' into amd-master

Parcourir la source 
Change-Id: Iac707fbc6311ce80a919f18d3c793db1175f8064
Cette révision appartient à :
Maneesh Gupta
2016-12-12 15:06:29 +05:30
Parent 7fa65ee250 428a1bc79f
révision 21a04603d1
22 fichiers modifiés avec 616 ajouts et 185 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
CMakeLists.txt
+10 -30
Voir le fichier
@@ -115,16 +115,6 @@ else()
message(FATAL_ERROR "Don't know where to install HIP. Please specify absolute path using -DCMAKE_INSTALL_PREFIX")
endif()
# Set if we need to build shared or static library
if(NOT DEFINED HIP_LIB_TYPE)
if(NOT DEFINED ENV{HIP_LIB_TYPE})
set(HIP_LIB_TYPE 1)
else()
set(HIP_LIB_TYPE $ENV{HIP_LIB_TYPE})
endif()
endif()
add_to_config(_buildInfo HIP_LIB_TYPE)
# Check if we need to build hipify-clang
if(NOT DEFINED HIPIFY_CLANG_LLVM_DIR)
if(DEFINED ENV{HIPIFY_CLANG_LLVM_DIR})
@@ -173,7 +163,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
set(SOURCE_FILES_SHARED
set(SOURCE_FILES_RUNTIME
src/hip_hcc.cpp
src/hip_context.cpp
src/hip_device.cpp
@@ -184,27 +174,23 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/hip_stream.cpp
src/hip_module.cpp)
set(SOURCE_FILES_STATIC
set(SOURCE_FILES_DEVICE
src/device_util.cpp
src/hip_ldg.cpp
src/hip_fp16.cpp)
if(${HIP_LIB_TYPE} EQUAL 0)
add_library(hip_hcc OBJECT ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
elseif(${HIP_LIB_TYPE} EQUAL 1)
add_library(hip_hcc STATIC ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
else()
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
add_library(hip_hcc SHARED ${SOURCE_FILES_SHARED})
add_library(hip_device STATIC ${SOURCE_FILES_STATIC})
add_dependencies(hip_device hip_hcc)
endif()
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
add_dependencies(hip_hcc_static hip_hcc)
add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})
add_dependencies(hip_device hip_hcc)
# Generate hcc_version.txt
add_custom_target(query_hcc_version COMMAND ${HCC_HOME}/bin/hcc --version > ${PROJECT_BINARY_DIR}/hcc_version.tmp)
add_custom_target(check_hcc_version COMMAND ${CMAKE_COMMAND} -E copy_if_different ${PROJECT_BINARY_DIR}/hcc_version.tmp ${PROJECT_BINARY_DIR}/hcc_version.txt DEPENDS query_hcc_version)
set_source_files_properties(${PROJECT_BINARY_DIR}/hcc_version.txt PROPERTIES GENERATED TRUE)
set_source_files_properties(${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
set_source_files_properties(${SOURCE_FILES_RUNTIME} ${SOURCE_FILES_DEVICE} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
add_dependencies(hip_hcc check_hcc_version update_build_and_version_info)
# Generate .hipInfo
@@ -223,13 +209,7 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA
#############################
# Install hip_hcc if platform is hcc
if(HIP_PLATFORM STREQUAL "hcc")
if(${HIP_LIB_TYPE} EQUAL 0)
install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
elseif(${HIP_LIB_TYPE} EQUAL 1)
install(TARGETS hip_hcc DESTINATION lib)
else()
install(TARGETS hip_hcc hip_device DESTINATION lib)
endif()
install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib)
install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_ir.ll DESTINATION lib)
# Install .hipInfo
bin/hipcc
+20 -12
Voir le fichier
@@ -76,6 +76,7 @@ if ($HIP_PLATFORM eq "hcc") {
$HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc";
$HCC_VERSION=`${HCC_HOME}/bin/hcc --version | cut -d" " -f9 | tr -d "\n"`;
$HCC_VERSION_MAJOR=`${HCC_HOME}/bin/hcc --version | cut -d" " -f9 | cut -d"." -f1 | tr -d "\n"`;
$ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
@@ -92,14 +93,14 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPCXXFLAGS = $HCCFLAGS;
#### GCC system includes workaround ####
$WA = $ENV{'HIP_HCC_SYS_INCLUDES_WA'} // 0;
if (${WA}) {
$HCC_WA_FLAGS = " ";
if ($HCC_VERSION_MAJOR eq 1) {
my $GCC_CUR_VER = `gcc -dumpversion`;
my $GPP_CUR_VER = `g++ -dumpversion`;
$GCC_CUR_VER =~ s/\R//g;
$GPP_CUR_VER =~ s/\R//g;
if (${GCC_CUR_VER} eq ${GPP_CUR_VER}) {
$HIPCXXFLAGS .= "-I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER}";
$HCC_WA_FLAGS .= " -I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER} ";
}
}
@@ -192,7 +193,8 @@ my $needHipHcc = ($HIP_PLATFORM eq 'hcc'); # set if we need to link hip_hcc
my $printHipVersion = 0; # print HIP version
my $runCmd = 1;
my $buildDeps = 0;
my $linkType = 0;
my $setLinkType = 0;
my @options = ();
my @inputs = ();
@@ -260,6 +262,7 @@ foreach $arg (@ARGV)
if(($arg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
{
$HIPCXXFLAGS .= " -stdlib=libstdc++";
$HIPCXXFLAGS .= $HCC_WA_FLAGS;
$setStdLib = 1;
}
if($arg eq '--version') {
@@ -273,11 +276,20 @@ foreach $arg (@ARGV)
$compileOnly = 1;
$buildDeps = 1;
}
if($arg eq '-use_fast_math') {
print "In fast Math";
$HIPCXXFLAGS .= " -DHIP_FAST_MATH ";
}
if(($arg eq '-use-staticlib') and ($setLinkType eq 0))
{
$linkType = 0;
$setLinkType = 1;
}
if(($arg eq '-use-sharedlib') and ($setLinkType eq 0))
{
$linkType = 1;
$setLinkType = 1;
}
if ($arg =~ m/^-/) {
# options start with -
@@ -322,16 +334,12 @@ if ($buildDeps and $HIP_PLATFORM eq 'nvcc') {
if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
{
$HIPCXXFLAGS .= " -stdlib=libstdc++";
$HIPCXXFLAGS .= $HCC_WA_FLAGS;
}
if ($needHipHcc) {
$HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 1;
# TODO - remove the old sea-of-objects solution:
if ($HIP_LIB_TYPE eq 0) {
substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o $HIP_PATH/lib/hip_context.cpp.o $HIP_PATH/lib/hip_module.cpp.o ";
} elsif ($HIP_LIB_TYPE eq 1) {
substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ;
if ($linkType eq 0) {
substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc_static -lhip_device " ;
} else {
substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc -lhip_device ";
}
docs/markdown/hip_profiling.md
+11 -2
Voir le fichier
@@ -325,9 +325,18 @@ Some key information from the trace above.
Chicken bits are environment variables which cause the HIP, HCC, or HSA driver to disable some feature or optimization.
These are not intended for production but can be useful diagnose synchronization problems in the application (or driver).
Some of the most useful chicken bits are described here:
Some of the most useful chicken bits are described here. These bits are supported on the ROCm path:
HIP provides 3 environment variables in the HIP_*_BLOCKING family. These introduce additional synchronization and can be useful to isolate synchronization problems. Specifically, if the code works with this flag set, then it indicates the kernels are executing correctly, and any failures likely are causes by improper or missing synchronization. These flags will have performance impact and are not intended for production use.
- HIP_LAUNCH_BLOCKING=1 : Waits on the host after each kernel launch. Equivalent to setting CUDA_LAUNCH_BLOCKING.
- HIP_LAUNCH_BLOCKING_KERNELS: A comma-separated list of kernel names. The HIP runtime will wait on the host after one of the named kernels executes. This provides a more targeted version of HIP_LAUNCH_BLOCKING and may be useful to isolate exactly which kernel needs further analysis if HIP_LAUNCH_BLOCKING=1 improves functionality. There is no indication if kernel names are spelled incorrectly. One mechanism to verify that the blocking is working is to run with HIP_DB=api+sync and search for debug messages with "LAUNCH_BLOCKING".
- HIP_API_BLOCKING : Forces hipMemcpyAsync and hipMemsetAsync to be host-synchronous, meaning they will wait for the requested operation to complete before returning to the caller.
These options cause HCC to serialize. Useful if you have libraries or code which is calling HCC kernels directly rather than using HIP.
- HCC_SERIALZIE_KERNELS : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize.
- HCC_SERIALIZE_COPY : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize.0
- HIP_LAUNCH_BLOCKING=1 : On ROCm, this flag waits on the host after each kernel launches and after each memory copy command. On CUDA, the waits are only enforced after each kernel launch. This is useful to isolate synchronization problems. Specifically, if the code works with this flag set, then it indicates the kernels and memory management code are correct, and any failures likely are causes by improper or missing synchronization.
- HSA_ENABLE_SDMA=0 : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine. This flag is useful to isolate issues with the hardware copy engines.
- HSA_ENABLE_INTERRUPT=0 : Causes completion signals to be detected with memory-based polling rather than interrupts. Can be useful to diagnose interrupt storm issues in the driver.
- HSA_DISABLE_CACHE=1 : Disables the GPU L2 data cache.
hipify-clang/src/Cuda2Hip.cpp
+38 -18
Voir le fichier
@@ -2004,6 +2004,11 @@ static cl::opt<bool> Inplace("inplace",
cl::value_desc("inplace"),
cl::cat(ToolTemplateCategory));
static cl::opt<bool> NoBackup("no-backup",
cl::desc("Don't create a backup file for the hipified source"),
cl::value_desc("no-backup"),
cl::cat(ToolTemplateCategory));
static cl::opt<bool> NoOutput("no-output",
cl::desc("Don't write any translated output to stdout"),
cl::value_desc("no-output"),
@@ -2107,29 +2112,41 @@ int main(int argc, const char **argv) {
if (N) {
NoOutput = PrintStats = true;
}
if (NoOutput) {
if (Inplace) {
llvm::errs() << "Conflict: both -no-output and -inplace options are specified.\n";
return 1;
}
if (!dst.empty()) {
llvm::errs() << "Conflict: both -no-output and -o options are specified.\n";
return 1;
}
}
if (dst.empty()) {
dst = fileSources[0];
if (!Inplace) {
size_t pos = dst.rfind(".cu");
if (pos != std::string::npos) {
dst = dst.substr(0, pos) + ".hip.cu";
size_t pos = dst.rfind(".");
if (pos != std::string::npos && pos+1 < dst.size()) {
dst = dst.substr(0, pos) + ".hip." + dst.substr(pos+1, dst.size()-pos-1);
} else {
llvm::errs() << "Input .cu file was not specified.\n";
return 1;
dst += ".hip.cu";
}
}
} else {
if (Inplace) {
llvm::errs() << "Conflict: both -o and -inplace options are specified.";
llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
return 1;
}
dst += ".cu";
dst += ".hip";
}
// backup source file since tooling may change "inplace"
if (!NoBackup || !Inplace) {
std::ifstream source(fileSources[0], std::ios::binary);
std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
dest << source.rdbuf();
source.close();
dest.close();
}
// copy source file since tooling makes changes "inplace"
std::ifstream source(fileSources[0], std::ios::binary);
std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
dest << source.rdbuf();
source.close();
dest.close();
RefactoringTool Tool(OptionsParser.getCompilations(), dst);
ast_matchers::MatchFinder Finder;
@@ -2168,15 +2185,18 @@ int main(int argc, const char **argv) {
if (!Tool.applyAllReplacements(Rewrite)) {
DEBUG(dbgs() << "Skipped some replacements.\n");
}
Result = Rewrite.overwriteChangedFiles();
if (!Inplace) {
size_t pos = dst.rfind(".cu");
if (!NoOutput) {
Result = Rewrite.overwriteChangedFiles();
}
if (!Inplace && !NoOutput) {
size_t pos = dst.rfind(".");
if (pos != std::string::npos) {
rename(dst.c_str(), dst.substr(0, pos).c_str());
}
}
if (NoOutput) {
remove(dst.c_str());
}
if (PrintStats) {
printStats(fileSources[0], PPCallbacks, Callback);
}
include/hip/hcc_detail/hip_runtime.h
+4 -3
Voir le fichier
@@ -771,7 +771,8 @@ extern "C" __device__ void __threadfence(void);
*
* @warning __threadfence_system is a stub and map to no-op.
*/
__device__ void __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
//__device__ void __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
__device__ void __threadfence_system(void) ;
__device__ unsigned __hip_ds_bpermute(int index, unsigned src);
__device__ float __hip_ds_bpermutef(int index, float src);
@@ -828,7 +829,7 @@ extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
extern void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp);
// Due to multiple overloaded versions of ihipPreLaunchKernel, the numBlocks3D and blockDim3D can be either size_t or dim3 types
@@ -838,7 +839,7 @@ do {\
lp.dynamic_group_mem_bytes = _groupMemBytes; \
hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp, #_kernelName)); \
_kernelName (lp, ##__VA_ARGS__);\
ihipPostLaunchKernel(trueStream, lp);\
ihipPostLaunchKernel(#_kernelName, trueStream, lp);\
} while(0)
include/hip/hcc_detail/hip_runtime_api.h
+131 -2
Voir le fichier
@@ -57,6 +57,19 @@ typedef struct ihipDevice_t *hipDevice_t;
typedef struct ihipStream_t *hipStream_t;
//TODO: IPC implementation
#define hipIpcMemLazyEnablePeerAccess 0
typedef struct ihipIpcMemHandle_t *hipIpcMemHandle_t;
//TODO: IPC event handle currently unsupported
struct ihipIpcEventHandle_t;
typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
//END TODO
typedef struct ihipModule_t *hipModule_t;
typedef struct ihipFunction_t *hipFunction_t;
@@ -1281,6 +1294,18 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags);
*/
hipError_t hipDeviceDisablePeerAccess (int peerDeviceId);
/**
* @brief Get information on memory allocations.
*
* @param [out] pbase - BAse pointer address
* @param [out] psize - Size of allocation
* @param [in] dptr- Device Pointer
*
* @returns #hipSuccess, #hipErrorInvalidDevicePointer
*
* @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
*/
hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr );
#ifndef USE_PEER_NON_UNIFIED
#define USE_PEER_NON_UNIFIED 1
@@ -1590,14 +1615,15 @@ hipError_t hipDeviceComputeCapability(int *major,int *minor,hipDevice_t device);
hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
/**
* @brief Returns a PCI Bus Id string for the device.
* @brief Returns a PCI Bus Id string for the device, overloaded to take int device ID.
* @param [out] pciBusId
* @param [in] len
* @param [in] device
*
* @returns #hipSuccess, #hipErrorInavlidDevice
*/
hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device);
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,int device);
/**
* @brief Returns a handle to a compute device.
@@ -1779,13 +1805,116 @@ hipError_t hipProfilerStop();
* @}
*/
//TODO: implement IPC apis
/**
* @brief Gets an interprocess memory handle for an existing device memory
* allocation
*
* Takes a pointer to the base of an existing device memory allocation created
* with hipMalloc and exports it for use in another process. This is a
* lightweight operation and may be called multiple times on an allocation
* without adverse effects.
*
* If a region of memory is freed with hipFree and a subsequent call
* to hipMalloc returns memory with the same device address,
* hipIpcGetMemHandle will return a unique handle for the
* new memory.
*
* @param handle - Pointer to user allocated hipIpcMemHandle to return
* the handle in.
* @param devPtr - Base pointer to previously allocated device memory
*
* @returns
* hipSuccess,
* hipErrorInvalidResourceHandle,
* hipErrorMemoryAllocation,
* hipErrorMapBufferObjectFailed,
*
*/
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *devPtr);
/**
* @brief Opens an interprocess memory handle exported from another process
* and returns a device pointer usable in the local process.
*
* Maps memory exported from another process with hipIpcGetMemHandle into
* the current device address space. For contexts on different devices
* hipIpcOpenMemHandle can attempt to enable peer access between the
* devices as if the user called hipDeviceEnablePeerAccess. This behavior is
* controlled by the hipIpcMemLazyEnablePeerAccess flag.
* hipDeviceCanAccessPeer can determine if a mapping is possible.
*
* Contexts that may open hipIpcMemHandles are restricted in the following way.
* hipIpcMemHandles from each device in a given process may only be opened
* by one context per device per other process.
*
* Memory returned from hipIpcOpenMemHandle must be freed with
* hipIpcCloseMemHandle.
*
* Calling hipFree on an exported memory region before calling
* hipIpcCloseMemHandle in the importing context will result in undefined
* behavior.
*
* @param devPtr - Returned device pointer
* @param handle - hipIpcMemHandle to open
* @param flags - Flags for this operation. Must be specified as hipIpcMemLazyEnablePeerAccess
*
* @returns
* hipSuccess,
* hipErrorMapBufferObjectFailed,
* hipErrorInvalidResourceHandle,
* hipErrorTooManyPeers
*
* @note No guarantees are made about the address returned in @p *devPtr.
* In particular, multiple processes may not receive the same address for the same @p handle.
*
*/
hipError_t hipIpcOpenMemHandle(void **devPtr,
hipIpcMemHandle_t handle, unsigned int flags);
/**
* @brief Close memory mapped with hipIpcOpenMemHandle
*
* Unmaps memory returnd by hipIpcOpenMemHandle. The original allocation
* in the exporting process as well as imported mappings in other processes
* will be unaffected.
*
* Any resources used to enable peer access will be freed if this is the
* last mapping using them.
*
* @param devPtr - Device pointer returned by hipIpcOpenMemHandle
*
* @returns
* hipSuccess,
* hipErrorMapBufferObjectFailed,
* hipErrorInvalidResourceHandle,
*
*/
hipError_t hipIpcCloseMemHandle(void *devPtr);
// hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
// hipError_t hipIpcCloseMemHandle(void *devPtr);
// // hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
// hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
#ifdef __cplusplus
} /* extern "c" */
#endif
#ifdef __cplusplus
/**
* @brief Returns a PCI Bus Id string for the device.
* @param [out] pciBusId
* @param [in] len
* @param [hipDevice_t] device
*
* @returns #hipSuccess, #hipErrorInavlidDevice
*/
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
#endif
/**
*-------------------------------------------------------------------------------------------------
include/hip/hcc_detail/hip_vector_types.h
+4
Voir le fichier
@@ -343,6 +343,10 @@ __HIP_DEVICE__ double2 make_double2(double, double );
__HIP_DEVICE__ double3 make_double3(double, double, double );
__HIP_DEVICE__ double4 make_double4(double, double, double, double );
extern __HIP_DEVICE__ double __longlong_as_double(long long int x);
extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
/*
///---
// Inline functions for creating vector types from basic types
include/hip/hip_runtime_api.h
+1
Voir le fichier
@@ -210,6 +210,7 @@ typedef enum hipError_t {
hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically not seen in production systems.
hipErrorHostMemoryAlreadyRegistered = 1061, ///< Produced when trying to lock a page-locked memory.
hipErrorHostMemoryNotRegistered = 1062, ///< Produced when trying to unlock a non-page-locked memory.
hipErrorMapBufferObjectFailed = 1071, ///< Produced when the IPC memory attach failed from ROCr.
hipErrorTbd ///< Marker that more error codes are needed.
} hipError_t;
include/hip/nvcc_detail/hip_runtime_api.h
+12 -2
Voir le fichier
@@ -639,6 +639,11 @@ inline static hipError_t hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned in
return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags));
}
inline static hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
{
return hipCUResultTohipError(cuMemGetAddressRange( pbase , psize , dptr));
}
inline static hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count )
{
return hipCUDAErrorTohipError(cudaMemcpyPeer(dst, dstDevice, src, srcDevice, count));
@@ -765,9 +770,14 @@ 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)
inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,int device)
{
return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
return hipCUDAErrorTohipError(cudaDeviceGetPCIBusId(pciBusId,len,device));
}
inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,hipDevice_t device)
{
return hipCUResultTohipError(cuDeviceGetPCIBusId(pciBusId,len,device));
}
inline static hipError_t hipDeviceGetByPCIBusId(int* device, const int *pciBusId)
packaging/hip_hcc.txt
+3 -8
Voir le fichier
@@ -1,14 +1,9 @@
cmake_minimum_required(VERSION 2.8.3)
project(hip_hcc)
if(@HIP_LIB_TYPE@ EQUAL 0)
install(DIRECTORY @PROJECT_BINARY_DIR@/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
elseif(@HIP_LIB_TYPE@ EQUAL 1)
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.a DESTINATION lib)
else()
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
endif()
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib)
src/device_util.cpp
+16 -14
Voir le fichier
@@ -14,8 +14,7 @@ all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
@@ -23,6 +22,7 @@ THE SOFTWARE.
#include <hc.hpp>
#include <grid_launch.h>
#include <hc_math.hpp>
#include "device_util.h"
#include "hip/hip_runtime.h"
@@ -33,17 +33,6 @@ THE SOFTWARE.
This is the best place to put them because the device
global variables need to be initialized at the start.
*/
#define NUM_PAGES_PER_THREAD 16
#define SIZE_OF_PAGE 64
#define NUM_THREADS_PER_CU 64
#define NUM_CUS_PER_GPU 64
#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
#define SIZE_OF_HEAP SIZE_MALLOC
size_t g_malloc_heap_size = SIZE_OF_HEAP;
__attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
__attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
@@ -2597,7 +2586,18 @@ __HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
return d4;
}
__device__ void __threadfence_system(void){
__HIP_DEVICE__ double __longlong_as_double(long long int x)
{
return static_cast<double>(x);
}
__HIP_DEVICE__ long long __double_as_longlong(double x)
{
return static_cast<long long>(x);
}
__HIP_DEVICE__ void __threadfence_system(void){
// no-op
}
@@ -3390,3 +3390,5 @@ __host__ double norm4d(double a, double b, double c, double d)
{
return std::sqrt(a*a + b*b + c*c + d*d);
}
src/device_util.h
+38
Voir le fichier
@@ -0,0 +1,38 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifndef DEVICE_UTIL_H
#define DEVICE_UTIL_H
/*
Heap size computation for malloc and free device functions.
*/
#define NUM_PAGES_PER_THREAD 16
#define SIZE_OF_PAGE 64
#define NUM_THREADS_PER_CU 64
#define NUM_CUS_PER_GPU 64
#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
#define SIZE_OF_HEAP SIZE_MALLOC
#endif
src/hip_device.cpp
+25 -7
Voir le fichier
@@ -23,6 +23,7 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "hip_hcc.h"
#include "trace_helper.h"
#include "device_util.h"
//-------------------------------------------------------------------------------------------------
//Devices
@@ -97,8 +98,6 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig)
return ihipLogStatus(hipSuccess);
}
extern "C" size_t g_malloc_heap_size;
hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit_t limit)
{
HIP_INIT_API(pValue, limit);
@@ -106,7 +105,7 @@ hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit_t limit)
return ihipLogStatus(hipErrorInvalidValue);
}
if(limit == hipLimitMallocHeapSize) {
*pValue = g_malloc_heap_size;
*pValue = (size_t)SIZE_OF_HEAP;
return ihipLogStatus(hipSuccess);
}else{
return ihipLogStatus(hipErrorUnsupportedLimit);
@@ -306,7 +305,7 @@ hipError_t hipSetDeviceFlags( unsigned int flags)
}
}
unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax;
unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax;
if (flags & (~supportedFlags)) {
e = hipErrorInvalidValue;
@@ -338,15 +337,34 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
return ihipLogStatus(e);
}
hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
#ifdef __cplusplus
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
{
HIP_INIT_API(pciBusId,len, device);
HIP_INIT_API(pciBusId, len, device);
hipError_t e = hipSuccess;
int deviceId= device->_deviceId;
e = ihipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId);
int tempPciBusId = 0;
e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, deviceId);
if( e == hipSuccess) {
std::string tempPciStr = std::to_string(tempPciBusId);
memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
}
return ihipLogStatus(e);
}
#endif
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
{
HIP_INIT_API(pciBusId, len, device);
hipError_t e = hipSuccess;
int tempPciBusId = 0;
e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, device);
if( e == hipSuccess) {
std::string tempPciStr = std::to_string(tempPciBusId);
memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
}
return ihipLogStatus(e);
}
hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
{
HIP_INIT_API(bytes, device);
src/hip_hcc.cpp
+62 -37
Voir le fichier
@@ -61,6 +61,9 @@ const char *API_COLOR = KGRN;
const char *API_COLOR_END = KNRM;
int HIP_LAUNCH_BLOCKING = 0;
std::string HIP_LAUNCH_BLOCKING_KERNELS;
std::vector<std::string> g_hipLaunchBlockingKernels;
int HIP_API_BLOCKING = 0;
int HIP_PRINT_ENV = 0;
int HIP_TRACE_API= 0;
@@ -78,15 +81,17 @@ int HIP_WAIT_MODE = 0;
int HIP_FORCE_P2P_HOST = 0;
int HIP_DENY_PEER_ACCESS = 0;
// Force async copies to actually use the synchronous copy interface.
// Force async copies to actually use the synchronous copy interface.
int HIP_FORCE_SYNC_COPY = 0;
int HIP_COHERENT_HOST_ALLOC = 0;
#define HIP_USE_PRODUCT_NAME 0
#define HIP_USE_PRODUCT_NAME 1
//#define DISABLE_COPY_EXT 1
@@ -105,7 +110,7 @@ unsigned g_numLogicalThreads;
std::atomic<int> g_lastShortTid(1);
// Indexed by short-tid:
//
//
std::vector<ProfTrigger> g_dbStartTriggers;
std::vector<ProfTrigger> g_dbStopTriggers;
@@ -133,12 +138,12 @@ void recordApiTrace(std::string *fullStr, const std::string &apiStr)
if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) {
printf ("info: resume profiling at %lu\n", apiSeqNum);
RESUME_PROFILING;
RESUME_PROFILING;
g_dbStartTriggers.pop_back();
};
if ((tid < g_dbStopTriggers.size()) && (apiSeqNum >= g_dbStopTriggers[tid].nextTrigger())) {
printf ("info: stop profiling at %lu\n", apiSeqNum);
STOP_PROFILING;
STOP_PROFILING;
g_dbStopTriggers.pop_back();
};
@@ -211,8 +216,8 @@ hipError_t ihipSynchronize(void)
//=================================================================================================
ShortTid::ShortTid() :
_apiSeqNum(0)
{
_shortTid = g_lastShortTid.fetch_add(1);
{
_shortTid = g_lastShortTid.fetch_add(1);
if (COMPILE_HIP_DB && HIP_TRACE_API) {
std::stringstream tid_ss;
@@ -282,7 +287,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
} else if (HIP_WAIT_MODE == 2) {
waitMode = hc::hcWaitModeActive;
}
crit->_av.wait(waitMode);
}
@@ -358,13 +363,24 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
//---
// Must be called after kernel finishes, this releases the lock on the stream so other commands can submit.
void ihipStream_t::lockclose_postKernelCommand(hc::accelerator_view *av)
void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::accelerator_view *av)
{
bool blockThisKernel = false;
if (HIP_LAUNCH_BLOCKING) {
if (!g_hipLaunchBlockingKernels.empty()) {
std::string kernelNameString(kernelName);
for (auto o=g_hipLaunchBlockingKernels.begin(); o!=g_hipLaunchBlockingKernels.end(); o++) {
if ((*o == kernelNameString)) {
//printf ("force blocking for kernel %s\n", o->c_str());
blockThisKernel = true;
}
}
}
if (HIP_LAUNCH_BLOCKING || blockThisKernel) {
// TODO - fix this so it goes through proper stream::wait() call.// direct wait OK since we know the stream is locked.
av->wait(hc::hcWaitModeActive);
tprintf(DB_SYNC, " %s LAUNCH_BLOCKING for kernel completion\n", ToString(this).c_str());
tprintf(DB_SYNC, "%s LAUNCH_BLOCKING for kernel '%s' completion\n", ToString(this).c_str(), kernelName);
}
_criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
@@ -385,17 +401,11 @@ void ihipStream_t::launchModuleKernel(
uint32_t gridDimY,
uint32_t gridDimZ,
uint32_t groupSegmentSize,
uint32_t privateSegmentSize,
uint32_t privateSegmentSize,
void *kernarg,
size_t kernSize,
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());
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);
memcpy(kern, kernarg, kernSize);
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);
@@ -410,7 +420,7 @@ void ihipStream_t::launchModuleKernel(
dispatch_packet->grid_size_z = blockDimZ * gridDimZ;
dispatch_packet->group_segment_size = groupSegmentSize;
dispatch_packet->private_segment_size = privateSegmentSize;
dispatch_packet->kernarg_address = kern;
dispatch_packet->kernarg_address = kernarg;
dispatch_packet->kernel_object = kernel;
uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
(1 << HSA_PACKET_HEADER_BARRIER) |
@@ -455,7 +465,7 @@ bool ihipCtxCriticalBase_t<CtxMutex>::addPeerWatcher(const ihipCtx_t *thisCtx, i
auto match = std::find(_peers.begin(), _peers.end(), peerWatcher);
if (match == std::end(_peers)) {
// Not already a peer, let's update the list:
tprintf(DB_COPY, "addPeerWatcher. Allocations on %s now visible to peerWatcher %s.\n",
tprintf(DB_COPY, "addPeerWatcher. Allocations on %s now visible to peerWatcher %s.\n",
thisCtx->toString().c_str(), peerWatcher->toString().c_str());
_peers.push_back(peerWatcher);
recomputePeerAgents();
@@ -473,7 +483,7 @@ bool ihipCtxCriticalBase_t<CtxMutex>::removePeerWatcher(const ihipCtx_t *thisCtx
auto match = std::find(_peers.begin(), _peers.end(), peerWatcher);
if (match != std::end(_peers)) {
// Found a valid peer, let's remove it.
tprintf(DB_COPY, "removePeerWatcher. Allocations on %s no longer visible to former peerWatcher %s.\n",
tprintf(DB_COPY, "removePeerWatcher. Allocations on %s no longer visible to former peerWatcher %s.\n",
thisCtx->toString().c_str(), peerWatcher->toString().c_str());
_peers.remove(peerWatcher);
recomputePeerAgents();
@@ -813,7 +823,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)
prop->arch.hasFloatAtomicAdd = 0;
prop->arch.hasGlobalInt64Atomics = 1;
prop->arch.hasSharedInt64Atomics = 1;
prop->arch.hasDoubles = 1;
prop->arch.hasDoubles = 1;
prop->arch.hasWarpVote = 1;
prop->arch.hasWarpBallot = 1;
prop->arch.hasWarpShuffle = 1;
@@ -1181,7 +1191,7 @@ std::string HIP_DB_callback(void *var_ptr, const char *envVarString)
tokenize(e, '+', &tokens);
for (auto t=tokens.begin(); t!= tokens.end(); t++) {
for (int i=0; i<DB_MAX_FLAG; i++) {
if (!strcmp(t->c_str(), dbName[i]._shortName)) {
if (!strcmp(t->c_str(), dbName[i]._shortName)) {
*var_ptr_int |= (1<<i);
} // TODO - else throw error?
}
@@ -1243,7 +1253,15 @@ void ihipInit()
//-- READ HIP_PRINT_ENV env first, since it has impact on later env var reading
// TODO: In HIP/hcc, this variable blocks after both kernel commmands and data transfer. Maybe should be bit-mask for each command type?
READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." );
READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP kernel launches 'host-synchronous', so they block until any kernel launches. Alias: CUDA_LAUNCH_BLOCKING." );
READ_ENV_S(release, HIP_LAUNCH_BLOCKING_KERNELS, 0, "Comma-separated list of kernel names to make host-synchronous, so they block until completed.");
if (!HIP_LAUNCH_BLOCKING_KERNELS.empty()) {
tokenize(HIP_LAUNCH_BLOCKING_KERNELS, ',', &g_hipLaunchBlockingKernels);
}
READ_ENV_I(release, HIP_API_BLOCKING, 0, "Make HIP APIs 'host-synchronous', so they block until completed. Impacts hipMemcpyAsync, hipMemsetAsync." );
READ_ENV_C(release, HIP_DB, 0, "Print debug info. Bitmask (HIP_DB=0xff) or flags separated by '+' (HIP_DB=api+sync+mem+copy)", HIP_DB_callback);
if ((HIP_DB & (1<<DB_API)) && (HIP_TRACE_API == 0)) {
// Set HIP_TRACE_API default before we read it, so it is printed correctly.
@@ -1252,20 +1270,26 @@ void ihipInit()
READ_ENV_I(release, HIP_TRACE_API, 0, "Trace each HIP API call. Print function name and return code to stderr as program executes.");
READ_ENV_S(release, HIP_TRACE_API_COLOR, 0, "Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White");
READ_ENV_I(release, HIP_PROFILE_API, 0, "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, 0x2=full API name including args.");
READ_ENV_S(release, HIP_DB_START_API, 0, "Comma-separated list of tid.api_seq_num for when to start debug and profiling.");
READ_ENV_S(release, HIP_DB_STOP_API, 0, "Comma-separated list of tid.api_seq_num for when to stop debug and profiling.");
READ_ENV_C(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the sequence are visible to HIP applications and they are enumerated in the order of sequence.", HIP_VISIBLE_DEVICES_callback );
READ_ENV_I(release, HIP_WAIT_MODE, 0, "Force synchronization mode. 1= force yield, 2=force spin, 0=defaults specified in application");
READ_ENV_I(release, HIP_FORCE_P2P_HOST, 0, "Force use of host/staging copy for peer-to-peer copies.1=always use copies, 2=always return false for hipDeviceCanAccessPeer");
READ_ENV_I(release, HIP_FORCE_SYNC_COPY, 0, "Force all copies (even hipMemcpyAsync) to use sync copies");
// TODO - review, can we remove this?
READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced.");
READ_ENV_I(release, HIP_COHERENT_HOST_ALLOC, 0, "If set, all host memory will be allocated as fine-grained system memory. This allows threadfence_system to work but prevents host memory from being cached on GPU which may have performance impact.");
// Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled.
if (HIP_DB && !COMPILE_HIP_DB) {
fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)\n", HIP_DB);
@@ -1398,6 +1422,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
{
if (HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {
std::stringstream os_pre;
std::stringstream os;
@@ -1515,11 +1540,11 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
//---
//Called after kernel finishes execution.
//This releases the lock on the stream.
void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp)
{
tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
stream->lockclose_postKernelCommand(lp.av);
stream->lockclose_postKernelCommand(kernelName, lp.av);
MARKER_END();
}
@@ -1716,10 +1741,10 @@ unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDev
// hipMemKind must be "resolved" to a specific direction - cannot be default.
void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
const hc::AmPointerInfo *dstPtrInfo,
const hc::AmPointerInfo *srcPtrInfo,
hc::hcCommandKind *hcCopyDir,
void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
const hc::AmPointerInfo *dstPtrInfo,
const hc::AmPointerInfo *srcPtrInfo,
hc::hcCommandKind *hcCopyDir,
ihipCtx_t **copyDevice,
bool *forceUnpinnedCopy)
{
@@ -1749,11 +1774,11 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
if (HIP_FORCE_P2P_HOST & 0x1) {
*forceUnpinnedCopy = true;
tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n",
tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n",
(*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
} else {
tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst.\n",
tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst.\n",
(*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
}
} else {
@@ -1789,7 +1814,7 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
{
LockedAccessor_StreamCrit_t crit (_criticalData);
tprintf (DB_COPY, "copySync copyDev:%d dst=%p (phys_dev:%d, isDevMem:%d) src=%p(phys_dev:%d, isDevMem:%d) sz=%zu dir=%s forceUnpinnedCopy=%d\n",
copyDevice ? copyDevice->getDeviceNum():-1,
copyDevice ? copyDevice->getDeviceNum():-1,
dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
@@ -1846,7 +1871,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
bool forceUnpinnedCopy;
resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &copyDevice, &forceUnpinnedCopy);
tprintf (DB_COPY, "copyASync copyDev:%d dst=%p (phys_dev:%d, isDevMem:%d) src=%p(phys_dev:%d, isDevMem:%d) sz=%zu dir=%s forceUnpinnedCopy=%d\n",
copyDevice ? copyDevice->getDeviceNum():-1,
copyDevice ? copyDevice->getDeviceNum():-1,
dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
@@ -1883,8 +1908,8 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
};
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
if (HIP_API_BLOCKING) {
tprintf(DB_SYNC, "%s LAUNCH_BLOCKING for completion of hipMemcpyAsync(sz=%zu)\n", ToString(this).c_str(), sizeBytes);
this->wait(crit);
}
src/hip_hcc.h
+41 -18
Voir le fichier
@@ -25,6 +25,7 @@ THE SOFTWARE.
#include <hc.hpp>
#include <hsa/hsa.h>
#include "hsa/hsa_ext_amd.h"
#include "hip_util.h"
@@ -32,7 +33,7 @@ THE SOFTWARE.
#error("This version of HIP requires a newer version of HCC.");
#endif
#define USE_DISPATCH_HSA_KERNEL 0
#define USE_DISPATCH_HSA_KERNEL 1
//
@@ -45,6 +46,7 @@ extern const int release;
// TODO - this blocks both kernels and memory ops. Perhaps should have separate env var for kernels?
extern int HIP_LAUNCH_BLOCKING;
extern int HIP_API_BLOCKING;
extern int HIP_PRINT_ENV;
extern int HIP_PROFILE_API;
@@ -56,6 +58,8 @@ extern int HIP_STREAM_SIGNALS; /* number of signals to allocate at stream creat
extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */
extern int HIP_FORCE_P2P_HOST;
extern int HIP_COHERENT_HOST_ALLOC;
//---
// Chicken bits for disabling functionality to work around potential issues:
@@ -156,11 +160,6 @@ extern const char *API_COLOR_END;
#endif
// Compile code that force hipHostMalloc only allocates finegrained system memory.
#ifndef HIP_COHERENT_HOST_ALLOC
#define HIP_COHERENT_HOST_ALLOC 0
#endif
// Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.
@@ -369,17 +368,18 @@ struct LockedBase {
MUTEX_TYPE _mutex;
};
class ihipModule_t{
/**
* HIP IPC Handle Size
*/
#define HIP_IPC_HANDLE_SIZE 64
class ihipIpcMemHandle_t
{
public:
hsa_executable_t executable;
hsa_code_object_t object;
std::string fileName;
void *ptr;
size_t size;
hsa_amd_ipc_memory_t ipc_handle; ///< ipc memory handle on ROCr
char reserved[HIP_IPC_HANDLE_SIZE];
size_t psize;
};
class ihipFunction_t{
public:
ihipFunction_t(const char *name) {
@@ -401,6 +401,29 @@ public:
uint64_t _kernel;
};
class ihipModule_t {
public:
hsa_executable_t executable;
hsa_code_object_t object;
std::string fileName;
void *ptr;
size_t size;
ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {}
~ihipModule_t() {
for (int i = 0; i < hipFunctionTable.size(); ++i) {
ihipFunction_t *func = hipFunctionTable[i];
delete func;
}
hipFunctionTable.clear();
}
void registerFunction(ihipFunction_t* func) {
hipFunctionTable.push_back(func);
}
private:
std::vector<ihipFunction_t*> hipFunctionTable;
};
template <typename MUTEX_TYPE>
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
@@ -455,7 +478,7 @@ public:
//---
// 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::accelerator_view *av);
void lockclose_postKernelCommand(const char *kernelName, hc::accelerator_view *av);
void locked_wait(bool assertQueueEmpty=false);
@@ -497,9 +520,9 @@ private:
// The unsigned return is hipMemcpyKind
unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem);
void resolveHcMemcpyDirection(unsigned hipMemKind,
const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo,
hc::hcCommandKind *hcCopyDir,
void resolveHcMemcpyDirection(unsigned hipMemKind,
const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo,
hc::hcCommandKind *hcCopyDir,
ihipCtx_t **copyDevice,
bool *forceUnpinnedCopy);
src/hip_ir.ll
+22
Voir le fichier
@@ -12,4 +12,26 @@ define void @__threadfence_block() #1 {
ret void
}
define linkonce_odr spir_func i32 @__rocm_dp4a(i32 %in1, i32 %in2, i32 %in3) {
%val1 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0","=v,v,v"(i32 %in1, i32 %in2)
%ret1 = add i32 %val1, %in3
%val2 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:BYTE_1","=v,v,v"(i32 %in1, i32 %in2)
%ret2 = add i32 %ret1, %val2
%val3 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_2 src1_sel:BYTE_2","=v,v,v"(i32 %in1, i32 %in2)
%ret3 = add i32 %val3, %ret2
%val4 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:BYTE_3","=v,v,v"(i32 %in1, i32 %in2)
%ret4 = add i32 %val4, %ret3
ret i32 %ret4
}
define linkonce_odr spir_func i32 @__rocm_hfma(i32 %in1, i32 %in2, i32 %in3) {
tail call void asm "v_mac_f16 $0, $1, $2","v,v,v"(i32 %in1, i32 %in2, i32 %in3)
ret i32 %in3
}
define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
%val = tail call i32 asm "v_add_f16 $0, $1, $2","=v,v,v"(i32 %in1, i32 %in2)
ret i32 %val
}
attributes #1 = { alwaysinline nounwind }
src/hip_memory.cpp
+85 -24
Voir le fichier
@@ -28,6 +28,7 @@ THE SOFTWARE.
#include "hip_hcc.h"
#include "trace_helper.h"
#include "hip/hcc_detail/hip_texture.h"
#include <hc_am.hpp>
//-------------------------------------------------------------------------------------------------
//-------------------------------------------------------------------------------------------------
@@ -156,20 +157,6 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
return ihipLogStatus(hip_status);
}
void ihipReadSingleEnv(int *var_ptr, const char *var_name1, const char *description)
{
char * env = getenv(var_name1);
// Default is set when variable is initialized (at top of this file), so only override if we find
// an environment variable.
if (env) {
long int v = strtol(env, NULL, 0);
*var_ptr = (int) (v);
}
if (HIP_PRINT_ENV) {
printf ("%-30s = %2d : %s\n", var_name1, *var_ptr, description);
}
}
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
{
@@ -192,16 +179,12 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
const unsigned supportedFlags = hipHostMallocPortable | hipHostMallocMapped | hipHostMallocWriteCombined;
// Read from environment variable of HIP_COHERENT_HOST_ALLOC
int coherent_alloc=0;
ihipReadSingleEnv(&coherent_alloc, "HIP_COHERENT_HOST_ALLOC", "Flag to force allocate finegrained system memory");
if (flags & ~supportedFlags) {
hip_status = hipErrorInvalidValue;
}
else {
auto device = ctx->getWriteableDevice();
if(coherent_alloc){
if(HIP_COHERENT_HOST_ALLOC){
// Force to allocate finedgrained system memory
*ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
if(sizeBytes < 1 && (*ptr == NULL)){
@@ -852,13 +835,13 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
}
}
stream->lockclose_postKernelCommand(&crit->_av);
stream->lockclose_postKernelCommand("hipMemsetAsync", &crit->_av);
if (HIP_LAUNCH_BLOCKING) {
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset [stream:%p].\n", __func__, (void*)stream);
if (HIP_API_BLOCKING) {
tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str());
cf.wait();
tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
//tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
}
} else {
e = hipErrorInvalidValue;
@@ -905,7 +888,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes )
// TODO - is hipMemset supposed to be async?
cf.wait();
stream->lockclose_postKernelCommand(&crit->_av);
stream->lockclose_postKernelCommand("hipMemset", &crit->_av);
if (HIP_LAUNCH_BLOCKING) {
@@ -1037,3 +1020,81 @@ hipError_t hipFreeArray(hipArray* array)
return ihipLogStatus(hipStatus);
}
hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
{
HIP_INIT_API ( pbase , psize , dptr );
hipError_t hipStatus = hipSuccess;
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 );
am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , dptr );
if (status == AM_SUCCESS) {
*pbase = amPointerInfo._devicePointer;
*psize = amPointerInfo._sizeBytes;
}
else
hipStatus = hipErrorInvalidDevicePointer;
return hipStatus;
}
//TODO: IPC implementaiton:
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
HIP_INIT_API ( handle, devPtr);
hipError_t hipStatus = hipSuccess;
// Get the size of allocated pointer
size_t psize;
hc::accelerator acc;
hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 );
am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , devPtr );
if (status == AM_SUCCESS) {
psize = (size_t)amPointerInfo._sizeBytes;
}
else
hipStatus = hipErrorInvalidResourceHandle;
// Save the size of the pointer to hipIpcMemHandle
(*handle)->psize = psize;
// Create HSA ipc memory
hsa_status_t hsa_status =
hsa_amd_ipc_memory_create(devPtr, psize, &(*handle)->ipc_handle);
if(hsa_status!= HSA_STATUS_SUCCESS)
hipStatus = hipErrorMemoryAllocation;
return hipStatus;
}
hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
// HIP_INIT_API ( devPtr, handle.handle , flags);
hipError_t hipStatus = hipSuccess;
// Get the current device agent.
hc::accelerator acc;
hsa_agent_t *agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if(!agent)
return hipErrorInvalidResourceHandle;
//Attach ipc memory
hsa_status_t hsa_status =
hsa_amd_ipc_memory_attach(&handle->ipc_handle, handle->psize, 1, agent, devPtr);
if(hsa_status != HSA_STATUS_SUCCESS)
hipStatus = hipErrorMapBufferObjectFailed;
return hipStatus;
}
hipError_t hipIpcCloseMemHandle(void *devPtr){
HIP_INIT_API ( devPtr );
hipError_t hipStatus = hipSuccess;
hsa_status_t hsa_status =
hsa_amd_ipc_memory_detach(devPtr);
if(hsa_status != HSA_STATUS_SUCCESS)
return hipErrorInvalidResourceHandle;
return hipStatus;
}
// hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
// return hipSuccess;
// }
src/hip_module.cpp
+56 -5
Voir le fichier
@@ -126,7 +126,6 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
} else {
*module = new ihipModule_t;
size_t size = std::string::size_type(in.tellg());
void *p = NULL;
hsa_agent_t agent = currentDevice->_hsaAgent;
@@ -164,7 +163,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
return ihipLogStatus(ret);
}
hipError_t hipModuleUnload(hipModule_t hmod){
hipError_t hipModuleUnload(hipModule_t hmod)
{
// TODO - improve this synchronization so it is thread-safe.
// Currently we want for all inflight activity to complete, but don't prevent another
// thread from launching new kernels before we finish this operation.
ihipSynchronize();
hipError_t ret = hipSuccess;
hsa_status_t status = hsa_executable_destroy(hmod->executable);
if(status != HSA_STATUS_SUCCESS)
@@ -172,6 +176,11 @@ hipError_t hipModuleUnload(hipModule_t hmod){
ret = hipErrorInvalidValue;
}
status = hsa_code_object_destroy(hmod->object);
if(status != HSA_STATUS_SUCCESS)
{
ret = hipErrorInvalidValue;
}
status = hsa_memory_free(hmod->ptr);
if(status != HSA_STATUS_SUCCESS)
{
ret = hipErrorInvalidValue;
@@ -193,6 +202,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
}else{
*func = new ihipFunction_t(name);
hmod->registerFunction(*func);
int deviceId = ctx->getDevice()->_deviceId;
ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
@@ -268,11 +278,17 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
&groupSegmentSize);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
uint32_t privateSegmentSize;
status = hsa_executable_symbol_get_info(f->_kernelSymbol,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
&privateSegmentSize);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
privateSegmentSize += sharedMemBytes;
@@ -300,7 +316,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
aql.group_segment_size = groupSegmentSize;
aql.private_segment_size = privateSegmentSize;
aql.kernel_object = f->_kernel;
aql.setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
(1 << HSA_PACKET_HEADER_BARRIER) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
@@ -315,13 +331,32 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
hsa_signal_t signal;
status = hsa_signal_create(1, 0, NULL, &signal);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
/*
Allocate kernarg
*/
void *kern = nullptr;
hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(lp.av->get_hsa_kernarg_region());
status = hsa_amd_memory_pool_allocate(*pool, kernArgSize, 0, &kern);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)lp.av->get_hsa_agent(), 0, kern);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
memcpy(kern, config[1], kernArgSize);
/*
Launch AQL packet
*/
hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, config[1], kernArgSize, f->_kernel);
gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel);
/*
@@ -330,10 +365,26 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
/*
Destroy kernarg
*/
status = hsa_amd_memory_pool_free(kern);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
/*
Destroy the signal
*/
status = hsa_signal_destroy(signal);
if(status != HSA_STATUS_SUCCESS){
return ihipLogStatus(hipErrorNotFound);
}
#endif // USE_DISPATCH_HSA_KERNEL
ihipPostLaunchKernel(hStream, lp);
ihipPostLaunchKernel(f->_kernelName, hStream, lp);
}
tests/src/context/hipCtx_simple.cpp
+1 -1
Voir le fichier
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp HCC_OPTIONS --stdlib=libc++
* BUILD: %t %s ../test_common.cpp HCC_OPTIONS -stdlib=libc++
* RUN: %t
* HIT_END
*/
tests/src/hipDrvGetPCIBusId.cpp
+34
Voir le fichier
@@ -0,0 +1,34 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include<iostream>
#include<fstream>
#include<vector>
int main(){
hipInit(0);
hipDevice_t device;
hipDeviceGet(&device, 0);
char pciBusId[100];
hipDeviceGetPCIBusId(pciBusId,100,device);
printf("PCI Bus ID= %s\n",pciBusId);
return 0;
}
tests/src/hipHcc.cpp
+1 -1
Voir le fichier
@@ -22,7 +22,7 @@ THE SOFTWARE.
// Test the HCC-specific API extensions for HIP:
/* HIT_START
* BUILD: %t %s HCC_OPTIONS --stdlib=libc++
* BUILD: %t %s HCC_OPTIONS -stdlib=libc++
* RUN: %t
* HIT_END
*/
tests/src/kernel/hipLanguageExtensions.cpp
+1 -1
Voir le fichier
@@ -22,7 +22,7 @@ THE SOFTWARE.
// Collection of code to make sure that various features in the hip kernel language compile.
/* HIT_START
* BUILD: %t %s ../test_common.cpp HCC_OPTIONS --stdlib=libc++
* BUILD: %t %s ../test_common.cpp HCC_OPTIONS -stdlib=libc++
* RUN: %t
* HIT_END
*/