AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'amd-develop' into amd-master

Browse Source 
Change-Id: I4e5fe28aea06abf2780e28bcf58d0742aad0c25d
This commit is contained in:
Maneesh Gupta
2016-10-07 22:20:58 +05:30
parent 4381557e26 99299a27cc
commit f5faeabb07
162 changed files with 451 additions and 780 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+82 -74
View File
@@ -16,8 +16,8 @@ endmacro()
#############################
# Determine HIP_BASE_VERSION
execute_process(COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/bin/hipconfig --version
OUTPUT_VARIABLE HIP_BASE_VERSION
OUTPUT_STRIP_TRAILING_WHITESPACE)
OUTPUT_VARIABLE HIP_BASE_VERSION
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REPLACE "." ";" VERSION_LIST ${HIP_BASE_VERSION})
list(GET VERSION_LIST 0 HIP_VERSION_MAJOR)
list(GET VERSION_LIST 1 HIP_VERSION_MINOR)
@@ -27,11 +27,11 @@ list(GET VERSION_LIST 1 HIP_VERSION_MINOR)
# use the commit date, instead of build date
# add xargs to remove strange trailing newline character
execute_process(COMMAND git show -s --format=@%ct
COMMAND xargs
COMMAND date -f - --utc +%y%U%w
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
OUTPUT_VARIABLE HIP_VERSION_PATCH
OUTPUT_STRIP_TRAILING_WHITESPACE)
COMMAND xargs
COMMAND date -f - --utc +%y%U%w
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
OUTPUT_VARIABLE HIP_VERSION_PATCH
OUTPUT_STRIP_TRAILING_WHITESPACE)
set(HIP_VERSION $HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH)
add_to_config(_versionInfo HIP_VERSION_MAJOR)
@@ -45,8 +45,8 @@ add_to_config(_versionInfo HIP_VERSION_PATCH)
if(NOT DEFINED HIP_PLATFORM)
if(NOT DEFINED ENV{HIP_PLATFORM})
execute_process(COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/bin/hipconfig --platform
OUTPUT_VARIABLE HIP_PLATFORM
OUTPUT_STRIP_TRAILING_WHITESPACE)
OUTPUT_VARIABLE HIP_PLATFORM
OUTPUT_STRIP_TRAILING_WHITESPACE)
else()
set(HIP_PLATFORM $ENV{HIP_PLATFORM} CACHE STRING "HIP Platform")
endif()
@@ -68,9 +68,9 @@ if(HIP_PLATFORM STREQUAL "hcc")
endif()
if(IS_ABSOLUTE ${HCC_HOME} AND EXISTS ${HCC_HOME} AND IS_DIRECTORY ${HCC_HOME})
execute_process(COMMAND ${HCC_HOME}/bin/hcc --version
COMMAND cut -d\ -f9
OUTPUT_VARIABLE HCC_VERSION
OUTPUT_STRIP_TRAILING_WHITESPACE)
COMMAND cut -d\ -f9
OUTPUT_VARIABLE HCC_VERSION
OUTPUT_STRIP_TRAILING_WHITESPACE)
message(STATUS "Looking for HCC in: " ${HCC_HOME} ". Found version: " ${HCC_VERSION})
else()
message(FATAL_ERROR "Don't know where to find HCC. Please specify abolute path using -DHCC_HOME")
@@ -100,9 +100,9 @@ endif()
# Determine HIP install path
if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT AND CMAKE_INSTALL_PREFIX MATCHES "/usr/local")
if(CMAKE_BUILD_TYPE MATCHES Debug)
set(CMAKE_INSTALL_PREFIX ${CMAKE_CURRENT_SOURCE_DIR} CACHE PATH "Installation path for HIP" FORCE)
set(CMAKE_INSTALL_PREFIX ${CMAKE_CURRENT_SOURCE_DIR} CACHE PATH "Installation path for HIP" FORCE)
elseif(CMAKE_BUILD_TYPE MATCHES Release)
set(CMAKE_INSTALL_PREFIX "/opt/rocm/hip" CACHE PATH "Installation path for HIP" FORCE)
set(CMAKE_INSTALL_PREFIX "/opt/rocm/hip" CACHE PATH "Installation path for HIP" FORCE)
else()
message(FATAL_ERROR "Invalid CMAKE_BUILD_TYPE specified. Valid values are Debug and Release")
endif()
@@ -123,12 +123,10 @@ if(NOT DEFINED HIP_LIB_TYPE)
endif()
add_to_config(_buildInfo HIP_LIB_TYPE)
# Check if we need to build clang hipify
if(NOT DEFINED BUILD_CLANG_HIPIFY)
if(NOT DEFINED ENV{BUILD_CLANG_HIPIFY})
set(BUILD_CLANG_HIPIFY 0)
else()
set(BUILD_CLANG_HIPIFY $ENV{BUILD_CLANG_HIPIFY})
# Check if we need to build hipify-clang
if(NOT DEFINED HIPIFY_CLANG_LLVM_DIR)
if(DEFINED ENV{HIPIFY_CLANG_LLVM_DIR})
set(HIPIFY_CLANG_LLVM_DIR $ENV{HIPIFY_CLANG_LLVM_DIR})
endif()
endif()
@@ -149,9 +147,7 @@ add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
add_custom_target(update_build_and_version_info ALL COMMAND make rebuild_cache)
# Build clang hipify if enabled
if(BUILD_CLANG_HIPIFY)
add_subdirectory(clang-hipify)
endif()
add_subdirectory(hipify-clang)
# Build hip_hcc if platform is hcc
if(HIP_PLATFORM STREQUAL "hcc")
@@ -166,7 +162,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}")
# Add remaining flags
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++")
set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include -stdlib=libc++")
# Set compiler and compiler flags
set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc")
@@ -175,17 +171,17 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
set(SOURCE_FILES src/device_util.cpp
src/hip_hcc.cpp
src/hip_context.cpp
src/hip_device.cpp
src/hip_error.cpp
src/hip_event.cpp
src/hip_ldg.cpp
src/hip_memory.cpp
src/hip_peer.cpp
src/hip_stream.cpp
src/hip_fp16.cpp
src/hip_module.cpp)
src/hip_hcc.cpp
src/hip_context.cpp
src/hip_device.cpp
src/hip_error.cpp
src/hip_event.cpp
src/hip_ldg.cpp
src/hip_memory.cpp
src/hip_peer.cpp
src/hip_stream.cpp
src/hip_fp16.cpp
src/hip_module.cpp)
if(${HIP_LIB_TYPE} EQUAL 0)
add_library(hip_hcc OBJECT ${SOURCE_FILES})
@@ -211,7 +207,7 @@ file(WRITE "${PROJECT_BINARY_DIR}/.version" ${_versionInfo})
# Build doxygen documentation
add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMAKE_CURRENT_SOURCE_DIR}/docs/doxygen-input/doxy.cfg
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs)
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs)
#############################
# Install steps
@@ -233,7 +229,7 @@ install(FILES ${PROJECT_BINARY_DIR}/.version DESTINATION bin)
# Install src, bin, include & cmake if necessary
execute_process(COMMAND test ${CMAKE_INSTALL_PREFIX} -ef ${CMAKE_CURRENT_SOURCE_DIR}
RESULT_VARIABLE INSTALL_SOURCE)
RESULT_VARIABLE INSTALL_SOURCE)
if(NOT ${INSTALL_SOURCE} EQUAL 0)
install(DIRECTORY src DESTINATION .)
install(DIRECTORY bin DESTINATION . USE_SOURCE_PERMISSIONS)
@@ -249,13 +245,24 @@ set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_base)
configure_file(packaging/hip_base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
configure_file(packaging/hip_base.postinst ${BUILD_DIR}/postinst @ONLY)
configure_file(packaging/hip_base.prerm ${BUILD_DIR}/prerm @ONLY)
add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
if(NOT BUILD_HIPIFY_CLANG)
add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
else()
add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS hipify-clang)
endif()
# Package: hip_hcc
set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_hcc)
@@ -263,46 +270,46 @@ configure_file(packaging/hip_hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
configure_file(packaging/hip_hcc.postinst ${BUILD_DIR}/postinst @ONLY)
configure_file(packaging/hip_hcc.prerm ${BUILD_DIR}/prerm @ONLY)
add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS hip_hcc)
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS hip_hcc)
# Package: hip_nvcc
set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_nvcc)
configure_file(packaging/hip_nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
# Package: hip_doc
set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_doc)
configure_file(packaging/hip_doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
# Package: hip_samples
set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_samples)
configure_file(packaging/hip_samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} .
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
COMMAND rm -rf *.deb *.rpm *.tar.gz
COMMAND make package
COMMAND cp *.deb ${PROJECT_BINARY_DIR}
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR})
# Package: all
if(POLICY CMP0037)
@@ -325,14 +332,15 @@ if(POLICY CMP0037)
cmake_policy(SET CMP0037 OLD)
endif()
add_custom_target(install_for_test COMMAND "${CMAKE_COMMAND}" --build . --target install
WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
execute_process(COMMAND getconf _NPROCESSORS_ONLN OUTPUT_VARIABLE DASH_JAY OUTPUT_STRIP_TRAILING_WHITESPACE)
add_custom_target(test COMMAND ${CMAKE_COMMAND} .
COMMAND make -j ${DASH_JAY}
COMMAND make test
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS install_for_test)
COMMAND make -j ${DASH_JAY}
COMMAND make test
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS install_for_test)
if(POLICY CMP0037)
cmake_policy(POP)
endif()
# vim: ts=4:sw=4:expandtab:smartindent
README.md
+1 -1
View File
@@ -33,7 +33,7 @@ HIP releases are typically of two types. The tag naming convention is different
- [HIP Porting Guide](docs/markdown/hip_porting_guide.md)
- [HIP Porting Driver Guide](docs/markdown/hip_porting_driver_api.md)
- [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](clang-hipify/README.md)
- [hipify-clang](hipify-clang/README.md)
- [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
- [Release Notes](RELEASE.md)
bin/hipcc
+2 -2
View File
@@ -91,7 +91,7 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPCC=$HCC;
$HIPCXXFLAGS = $HCCFLAGS;
$HIPCXXFLAGS .= " -I$HIP_PATH/include/hcc_detail/cuda";
$HIPCXXFLAGS .= " -I$HIP_PATH/include/hip/hcc_detail/cuda";
$HIPCXXFLAGS .= " -I$HSA_PATH/include";
$HIPCXXFLAGS .= " -Wno-deprecated-register";
$HIPLDFLAGS = "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
@@ -102,7 +102,7 @@ if ($HIP_PLATFORM eq "hcc") {
if ($HOST_OSNAME eq "ubuntu" and $HOST_OSVER eq "\"16.04\"") {
# No additional flags required
} else {
$HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=0 -Wl,--defsym=_binary_kernel_spir_start=0 -Wl,--defsym=_binary_kernel_cl_start=0 -Wl,--defsym=_binary_kernel_cl_end=0";
$HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=1 -Wl,--defsym=_binary_kernel_spir_start=1 -Wl,--defsym=_binary_kernel_cl_start=1 -Wl,--defsym=_binary_kernel_cl_end=1";
}
# Satisfy HCC dependencies
clang-hipify/CMakeLists.txt
-77
View File
@@ -1,77 +0,0 @@
cmake_minimum_required(VERSION 2.8.8)
project(hipify-clang)
# Find LLVM package
find_package(LLVM 3.8 QUIET PATHS ${LLVM_DIR} NO_DEFAULT_PATH)
if (NOT ${LLVM_FOUND})
message(FATAL_ERROR "Don't know where to find LLVM (v3.8) package. Please specify absolute path using -DLLVM_DIR")
endif()
list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
include(AddLLVM)
message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
include_directories(${LLVM_INCLUDE_DIRS})
link_directories(${LLVM_LIBRARY_DIRS})
add_definitions(${LLVM_DEFINITIONS})
add_llvm_executable(hipify-clang src/Cuda2Hip.cpp )
find_program(LIT_COMMAND lit)
# Link against LLVM and CLANG tools libraries
target_link_libraries(hipify-clang
clangASTMatchers
clangFrontend
clangTooling
clangParse
clangSerialization
clangSema
clangEdit
clangLex
clangAnalysis
clangDriver
clangAST
clangToolingCore
clangRewrite
clangBasic
LLVMProfileData
LLVMSupport
LLVMMCParser
LLVMMC
LLVMBitReader
LLVMOption
LLVMCore)
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"")
install(TARGETS hipify-clang
DESTINATION bin)
# tests
set(Python_ADDITIONAL_VERSIONS 2.7)
include(FindPythonInterp)
if( NOT PYTHONINTERP_FOUND )
message(FATAL_ERROR
"Unable to find Python interpreter, required for builds and testing\n\n"
"Please install Python or specify the PYTHON_EXECUTABLE CMake variable.")
endif()
set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
configure_file(
${CMAKE_SOURCE_DIR}/tests/clang-hipify/lit.site.cfg.in
${CMAKE_CURRENT_BINARY_DIR}/tests/clang-hipify/lit.site.cfg
@ONLY)
add_lit_testsuite(test-hipify "Running HIPify regression tests"
${CMAKE_SOURCE_DIR}/tests/clang-hipify
PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/clang-hipify/lit.site.cfg
DEPENDS hipify-clang lit
)
add_custom_target(test-clang-hipify)
add_dependencies(test-clang-hipify test-hipify)
set_target_properties(test-clang-hipify PROPERTIES FOLDER "Tests")
docs/Makefile
-5
View File
@@ -1,5 +0,0 @@
HIP_PATH ?= ..
DOXYGEN_SRC=$(wildcard $(HIP_PATH)/include/hcc_detail/*.h) $(wildcard $(HIP_PATH)/include/*.h) $(HIP_PATH)/src/hip_hcc.cpp $(wildcard doxygen-input/*)
doc: $(DOXYGEN_SRC)
HIP_PATH=$(HIP_PATH) doxygen doxygen-input/doxy.cfg
docs/doxygen-input/doxy.cfg
+2 -2
View File
@@ -760,8 +760,8 @@ WARN_LOGFILE =
INPUT = $(HIP_PATH)/docs/doxygen-input/mainpage.txt \
$(HIP_PATH)/docs/doxygen-input/sync.txt \
$(HIP_PATH)/include/ \
$(HIP_PATH)/include/hcc_detail/ \
$(HIP_PATH)/include/hip \
$(HIP_PATH)/include/hip/hcc_detail/ \
$(HIP_PATH)/src/
# This tag can be used to specify the character encoding of the source files
docs/markdown/hip_faq.md
+2 -2
View File
@@ -114,7 +114,7 @@ These offer pointer-based memory interfaces (as opposed to opaque buffers) and c
- [hcsparse](https://bitbucket.org/multicoreware/hcsparse)
- [hcrng](https://bitbucket.org/multicoreware/hcrng)
Additionally, some of the cublas routines are automatically converted to hipblas equivalents by the clang-hipify tool. These APIs use cublas or hcblas depending on the platform, and replace the need
Additionally, some of the cublas routines are automatically converted to hipblas equivalents by the hipify-clang tool. These APIs use cublas or hcblas depending on the platform, and replace the need
to use conditional compilation.
### How does HIP compare with OpenCL?
@@ -148,7 +148,7 @@ The tools also struggle with more complex CUDA applications, in particular those
### Does Hipify automatically convert all source code?
Typically, Hipify can automatically convert almost all run-time code, and the coordinate indexing device code (i.e. threadIdx.x -> hipThreadIdx_x).
Most device code needs no additional conversion, since HIP and CUDA have similar names for math and built-in functions.
The clang-hipify tool will automatically modify the kernel signature as needed (automating a step that used to be done manually)
The hipify-clang tool will automatically modify the kernel signature as needed (automating a step that used to be done manually)
Additional porting may be required to deal with architecture feature queries or with CUDA capabilities that HIP doesn't support.
In general, developers should always expect to perform some platform-specific tuning and optimization.
docs/markdown/hip_kernel_language.md
Regular → Executable
+2
View File
@@ -451,6 +451,8 @@ Following is the list of supported integer intrinsics. Note that intrinsics are
| unsigned int __ffsll(long long int x) <br><sub>Find the position of least signigicant bit set to 1 in a 64 bit signed integer.</sub> |
| unsigned int __popc ( unsigned int x ) <br><sub>Count the number of bits that are set to 1 in a 32 bit integer.</sub> |
| int __popcll ( unsigned long long int x )<br><sub>Count the number of bits that are set to 1 in a 64 bit integer.</sub> |
| int __mul24 ( int x, int y )<br><sub>Multiply two 24bit integers.</sub> |
| unsigned int __umul24 ( unsigned int x, unsigned int y )<br><sub>Multiply two 24bit unsigned integers.</sub> |
<sub><b id="f3"><sup>[1]</sup></b>
The hcc implementation of __ffs() and __ffsll() contains code to add a constant +1 to produce the ffs result format.
For the cases where this overhead is not acceptable and programmer is willing to specialize for the platform,
docs/markdown/hip_porting_guide.md
+48 -25
View File
@@ -411,29 +411,55 @@ For example:
Device Code:
```
// Cuda Device Code
__constant__ float Array[1024];
__global__ void Inc(float *Out){
Int tx = hipThreadIdx_x;
Out[tx] = Array[tx] + 1;
#include<hip_runtime_api.h>
#include<hip_runtime.h>
#include<iostream>
#ifdef __HIP_PLATFORM_HCC__
__global__ void Inc(hipLaunchParm lp, float *Ad, float *Out)
#endif
#ifdef __HIP_PLATFORM_NVCC__
__constant__ float Ad[1024];
__global__ void Inc(hipLaunchParm lp, float *Out)
#endif
{
int tx = hipThreadIdx_x;
Out[tx] = Ad[tx] + 1.0f;
}
// HIP Device Code
__global__ void Inc(hipLaunchParm lp, float *Array, float *Out){
Int tx = hipThreadIdx_x;
Out[tx] = Array[tx] + 1;
int main()
{
float *A, *Ad;
float *Out, *Outd;
A = new float[1024];
Out = new float[1024];
for(uint32_t i=0;i<1024;i++)
{
A[i] = 1.0f*i;
Out[i] = 0.0f;
}
hipMalloc((void**)&Ad, 1024*sizeof(float));
hipMalloc((void**)&Outd, 1024*sizeof(float));
hipMemcpy(Outd, Out, 1024*sizeof(float), hipMemcpyHostToDevice);
#ifdef __HIP_PLATFORM_HCC__
assert(hipSuccess == hipMemcpy(Ad, A, 1024*sizeof(float), hipMemcpyHostToDevice));
hipLaunchKernel(Inc, dim3(1,1,1), dim3(1024,1,1), 0, 0, Ad, Outd);
#endif
#ifdef __HIP_PLATFORM_NVCC__
assert(hipSuccess == hipMemcpyToSymbol(Ad, A, 1024*sizeof(float)));
hipLaunchKernel(Inc, dim3(1,1,1), dim3(1024,1,1), 0, 0, Outd);
#endif
hipMemcpy(Out, Outd, 1024*sizeof(float), hipMemcpyDeviceToHost);
std::cout<<Out[10]<<" "<<A[10]<<std::endl;
assert(Out[10] - A[10] == 1.0f);
}
```
Host Code:
```
// CUDA Host Code
cudaMemcpyToSymbol(Array, hostArray, sizeofArray);
// HIP Host Code
hipMemcpy(Array, hostArray, sizeofArray);
```
## threadfence_system
Threadfence_system makes all device memory writes, all writes to mapped host memory, and all writes to peer memory visible to CPU and other GPU devices.
Some implementations can provide this behavior by flushing the GPU L2 cache.
@@ -544,13 +570,10 @@ HIP_PRINT_ENV = 1 : Print HIP environment variables.
HIP_LAUNCH_BLOCKING = 0 : Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING.
HIP_DB = 0 : Print various debug info. Bitmask, see hip_hcc.cpp for more information.
HIP_TRACE_API = 0 : Trace each HIP API call. Print function name and return code to stderr as program executes.
HIP_STAGING_SIZE = 64 : Size of each staging buffer (in KB)
HIP_STAGING_BUFFERS = 2 : Number of staging buffers to use in each direction. 0=use hsa_memory_copy.
HIP_PININPLACE = 0 : For unpinned transfers, pin the memory in-place in chunks before doing the copy. Under development.
HIP_STREAM_SIGNALS = 2 : Number of signals to allocate when new stream is created (signal pool will grow on demand)
HIP_TRACE_API_COLOR = green : Color to use for HIP_API. None/Red/Green/Yellow/Blue/Magenta/Cyan/White
HIP_ATP_MARKER = 0 : Add HIP function begin/end to ATP file generated with CodeXL
HIP_VISIBLE_DEVICES = 0 : Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence
HIP_DISABLE_HW_KERNEL_DEP = 1 : Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)
HIP_DISABLE_HW_COPY_DEP = 1 : Disable HW dependencies before copy commands - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)
HIP_NUM_KERNELS_INFLIGHT = 128 : Number of kernels per stream
```
hipify-clang/CMakeLists.txt
+82
View File
@@ -0,0 +1,82 @@
cmake_minimum_required(VERSION 2.8.8)
project(hipify-clang)
set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE)
# Find LLVM package
find_package(LLVM 3.8 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH)
if (NOT ${LLVM_FOUND})
message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM (v3.8) package using -DHIPIFY_CLANG_LLVM_DIR")
else()
list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
include(AddLLVM)
message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
include_directories(${LLVM_INCLUDE_DIRS})
link_directories(${LLVM_LIBRARY_DIRS})
add_definitions(${LLVM_DEFINITIONS})
add_llvm_executable(hipify-clang src/Cuda2Hip.cpp)
find_program(LIT_COMMAND lit)
# Link against LLVM and CLANG tools libraries
target_link_libraries(hipify-clang
clangASTMatchers
clangFrontend
clangTooling
clangParse
clangSerialization
clangSema
clangEdit
clangLex
clangAnalysis
clangDriver
clangAST
clangToolingCore
clangRewrite
clangBasic
LLVMProfileData
LLVMSupport
LLVMMCParser
LLVMMC
LLVMBitReader
LLVMOption
LLVMCore)
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"")
install(TARGETS hipify-clang DESTINATION bin)
# tests
set(Python_ADDITIONAL_VERSIONS 2.7)
include(FindPythonInterp)
if(NOT PYTHONINTERP_FOUND)
message(FATAL_ERROR
"Unable to find Python interpreter, required for builds and testing\n\n"
"Please install Python or specify the PYTHON_EXECUTABLE CMake variable.")
endif()
set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
configure_file(
${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in
${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg
@ONLY)
add_lit_testsuite(test-hipify "Running HIPify regression tests"
${CMAKE_SOURCE_DIR}/tests/hipify-clang
PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg
DEPENDS hipify-clang lit
)
add_custom_target(test-hipify-clang)
add_dependencies(test-hipify-clang test-hipify)
set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests")
set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE)
endif()
# vim: ts=4:sw=4:expandtab:smartindent
clang-hipify/README.md → hipify-clang/README.md
+2 -2
View File
@@ -25,12 +25,12 @@ tar xvfJ clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz
- Enable build of hipify-clang and specify path to LLVM.
Note LLVM_DIR must be a full absolute path to the location extracted above. Here's an example assuming we extract the clang 3.8 package into ~/HIP/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/
Note HIPIFY_CLANG_LLVM_DIR must be a full absolute path to the location extracted above. Here's an example assuming we extract the clang 3.8 package into ~/HIP/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/
```shell
cd HIP
mkdir build
cd build
cmake -DBUILD_CLANG_HIPIFY=1 -DLLVM_DIR=~/HIP/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/ -DCMAKE_BUILD_TYPE=Release ..
cmake -DHIPIFY_CLANG_LLVM_DIR=~/HIP/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/ -DCMAKE_BUILD_TYPE=Release ..
make
make install
```
clang-hipify/src/Cuda2Hip.cpp → hipify-clang/src/Cuda2Hip.cpp
+5 -5
View File
@@ -107,9 +107,9 @@ struct cuda2hipMap {
cuda2hipRename["__CUDACC__"] = {"__HIPCC__", CONV_DEF, API_RUNTIME};
// CUDA includes
cuda2hipRename["cuda.h"] = {"hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER};
cuda2hipRename["cuda_runtime.h"] = {"hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME};
cuda2hipRename["cuda_runtime_api.h"] = {"hip_runtime_api.h", CONV_INCLUDE, API_RUNTIME};
cuda2hipRename["cuda.h"] = {"hip/hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER};
cuda2hipRename["cuda_runtime.h"] = {"hip/hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME};
cuda2hipRename["cuda_runtime_api.h"] = {"hip/hip_runtime_api.h", CONV_INCLUDE, API_RUNTIME};
// HIP includes
// TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h
@@ -1939,7 +1939,7 @@ public:
} while (false);
if (PP->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 &&
countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace->size() > 0) {
StringRef repName = "#include <hip_runtime.h>\n";
StringRef repName = "#include <hip/hip_runtime.h>\n";
SourceManager *SM = Result.SourceManager;
Replacement Rep(*SM, SM->getLocForStartOfFile(SM->getMainFileID()), 0, repName);
Replace->insert(Rep);
@@ -1961,7 +1961,7 @@ private:
void HipifyPPCallbacks::handleEndSource() {
if (Match->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 &&
countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace->size() > 0) {
StringRef repName = "#include <hip_runtime.h>\n";
StringRef repName = "#include <hip/hip_runtime.h>\n";
Replacement Rep(*_sm, _sm->getLocForStartOfFile(_sm->getMainFileID()), 0, repName);
Replace->insert(Rep);
countReps[CONV_INCLUDE_CUDA_MAIN_H]++;
include/hip
-1
View File
@@ -1 +0,0 @@
../include
include/hcc.h → include/hip/hcc.h
+1 -1
View File
@@ -2,7 +2,7 @@
#define HCC_H
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
#include <hcc_detail/hcc_acc.h>
#include "hip/hcc_detail/hcc_acc.h"
#endif
#endif
include/hcc_detail/cuda/cuda.h → include/hip/hcc_detail/cuda/cuda.h
View File
include/hcc_detail/cuda/math_functions.h → include/hip/hcc_detail/cuda/math_functions.h
View File
include/hcc_detail/hcc_acc.h → include/hip/hcc_detail/hcc_acc.h
View File
include/hcc_detail/hipComplex.h → include/hip/hcc_detail/hipComplex.h
View File
include/hcc_detail/hip_blas.h → include/hip/hcc_detail/hip_blas.h
View File
include/hcc_detail/hip_fp16.h → include/hip/hcc_detail/hip_fp16.h
+1 -1
View File
@@ -20,7 +20,7 @@ THE SOFTWARE.
#ifndef HIP_FP16_H
#define HIP_FP16_H
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
typedef struct{
unsigned x: 16;
include/hcc_detail/hip_hcc.h → include/hip/hcc_detail/hip_hcc.h
+1 -1
View File
@@ -109,7 +109,7 @@ extern const char *API_COLOR_END;
#endif
#define DB_SHOW_TID 1
#define DB_SHOW_TID 0
#if DB_SHOW_TID
#define COMPUTE_TID_STR \
include/hcc_detail/hip_ldg.h → include/hip/hcc_detail/hip_ldg.h
+2 -4
View File
@@ -21,11 +21,9 @@ THE SOFTWARE.
#define HIP_LDG_H
#if __HCC__
#include"hip_vector_types.h"
#include"host_defines.h"
#if __hcc_workweek__ >= 16164
#include"hip/hip_vector_types.h"
#include"hip/hcc_detail/host_defines.h"
#include "hip/hip_vector_types.h"
#include "hip/hcc_detail/host_defines.h"
__device__ char __ldg(const char* );
include/hcc_detail/hip_runtime.h → include/hip/hcc_detail/hip_runtime.h
Regular → Executable
+3
View File
@@ -430,6 +430,9 @@ __device__ unsigned int atomicInc(unsigned int* address,
__device__ unsigned int atomicDec(unsigned int* address,
unsigned int val);
//__mul24 __umul24
__device__ int __mul24(int arg1, int arg2);
__device__ unsigned int __umul24(unsigned int arg1, unsigned int arg2);
// integer intrinsic function __poc __clz __ffs __brev
__device__ unsigned int __popc( unsigned int input);
include/hcc_detail/hip_runtime_api.h → include/hip/hcc_detail/hip_runtime_api.h
View File
include/hcc_detail/hip_texture.h → include/hip/hcc_detail/hip_texture.h
View File
include/hcc_detail/hip_util.h → include/hip/hcc_detail/hip_util.h
View File
include/hcc_detail/hip_vector_types.h → include/hip/hcc_detail/hip_vector_types.h
View File
include/hcc_detail/host_defines.h → include/hip/hcc_detail/host_defines.h
View File
include/hcc_detail/trace_helper.h → include/hip/hcc_detail/trace_helper.h
View File
include/hipComplex.h → include/hip/hipComplex.h
View File
include/hip_common.h → include/hip/hip_common.h
View File
include/hip_fp16.h → include/hip/hip_fp16.h
View File
include/hip_runtime.h → include/hip/hip_runtime.h
View File
include/hip_runtime_api.h → include/hip/hip_runtime_api.h
View File
include/hip_vector_types.h → include/hip/hip_vector_types.h
View File
include/hipblas.h → include/hip/hipblas.h
View File
include/nvcc_detail/hipComplex.h → include/hip/nvcc_detail/hipComplex.h
View File
include/nvcc_detail/hip_blas.h → include/hip/nvcc_detail/hip_blas.h
View File
include/nvcc_detail/hip_runtime.h → include/hip/nvcc_detail/hip_runtime.h
View File
include/nvcc_detail/hip_runtime_api.h → include/hip/nvcc_detail/hip_runtime_api.h
+3 -3
View File
@@ -329,8 +329,8 @@ inline static hipError_t hipMemset(void* devPtr,int value, size_t count) {
return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
}
inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count) {
return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count));
inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hipStream_t stream = 0) {
return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream));
}
inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int device)
@@ -338,7 +338,7 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int dev
cudaDeviceProp cdprop;
cudaError_t cerror;
cerror = cudaGetDeviceProperties(&cdprop,device);
strcpy(p_prop->name,cdprop.name);
strncpy(p_prop->name,cdprop.name, 256);
p_prop->totalGlobalMem = cdprop.totalGlobalMem ;
p_prop->sharedMemPerBlock = cdprop.sharedMemPerBlock;
p_prop->regsPerBlock = cdprop.regsPerBlock;
packaging/convert_md_to_html.sh
+6 -6
View File
@@ -21,24 +21,24 @@ trap cleanup EXIT
export GRIPURL=$hip_srcdir
export GRIPHOME=$workdir
echo "CACHE_DIRECTORY = '$html_destdir/asset'" > $workdir/settings.py
mkdir -p $html_destdir $html_destdir/clang-hipify $html_destdir/docs/markdown
mkdir -p $html_destdir $html_destdir/hipify-clang $html_destdir/docs/markdown
# convert all md files to html
pushd $hip_srcdir
for f in *.md clang-hipify/*.md docs/markdown/*.md; do grip --export --no-inline $f $html_destdir/${f%.*}.html; done
for f in *.md hipify-clang/*.md docs/markdown/*.md; do grip --export --no-inline $f $html_destdir/${f%.*}.html; done
popd
# convert absolute links to relative links
pushd $html_destdir
for f in *.html; do sed -i "s?$GRIPURL/??g" $f; done
for f in clang-hipify/*.html; do sed -i "s?$GRIPURL/?../?g" $f; done
for f in hipify-clang/*.html; do sed -i "s?$GRIPURL/?../?g" $f; done
for f in docs/markdown/*.html; do sed -i "s?$GRIPURL/?../../?g" $f; done
popd
# update document titles
pushd $html_destdir
for f in *.html; do sed -i "s?.md - Grip??g" $f; done
for f in clang-hipify/*.html; do sed -i "s?.md - Grip??g" $f; done
for f in hipify-clang/*.html; do sed -i "s?.md - Grip??g" $f; done
for f in docs/markdown/*.html; do sed -i "s?.md - Grip??g" $f; done
popd
@@ -46,8 +46,8 @@ popd
pushd $html_destdir
for f in *.html; do sed -i "s?.md\"?.html\"?g" $f; done
for f in *.html; do sed -i "s?.md#?.html#?g" $f; done
for f in clang-hipify/*.html; do sed -i "s?.md\"?.html\"?g" $f; done
for f in clang-hipify/*.html; do sed -i "s?.md#?.html#?g" $f; done
for f in hipify-clang/*.html; do sed -i "s?.md\"?.html\"?g" $f; done
for f in hipify-clang/*.html; do sed -i "s?.md#?.html#?g" $f; done
for f in docs/markdown/*.html; do sed -i "s?.md\"?.html\"?g" $f; done
for f in docs/markdown/*.html; do sed -i "s?.md#?.html#?g" $f; done
popd
packaging/hip_base.postinst
+1 -4
View File
@@ -22,12 +22,9 @@ done
popd >/dev/null
# Soft-link to headers
HIPINCDIR=$HIPDIR/include
HIPINCDIR=$HIPDIR/include/hip
ROCMINCDIR=$ROCMDIR/include
mkdir -p $ROCMINCDIR
pushd $ROCMINCDIR
ln -s $HIPINCDIR hip
popd
pushd $HIPINCDIR
ln -s $HIPINCDIR hip
popd
packaging/hip_base.prerm
-3
View File
@@ -28,8 +28,5 @@ ROCMINCDIR=$ROCMDIR/include
pushd $ROCMINCDIR
rm hip
popd
pushd $HIPINCDIR
rm hip
popd
rmdir --ignore-fail-on-non-empty $ROCMINCDIR
packaging/hip_base.txt
+4 -1
View File
@@ -2,7 +2,10 @@ cmake_minimum_required(VERSION 2.8.3)
project(hip_base)
install(DIRECTORY @hip_SOURCE_DIR@/bin DESTINATION . USE_SOURCE_PERMISSIONS)
install(DIRECTORY @hip_SOURCE_DIR@/include DESTINATION . PATTERN "hip" EXCLUDE)
if(@BUILD_HIPIFY_CLANG@)
install(PROGRAMS @PROJECT_BINARY_DIR@/hipify-clang/hipify-clang DESTINATION bin)
endif()
install(DIRECTORY @hip_SOURCE_DIR@/include DESTINATION .)
install(FILES @PROJECT_BINARY_DIR@/.version DESTINATION bin)
install(DIRECTORY @hip_SOURCE_DIR@/cmake DESTINATION .)
packaging/hip_hcc.txt
+10 -3
View File
@@ -1,7 +1,14 @@
cmake_minimum_required(VERSION 2.8.3)
project(hip_hcc)
install(DIRECTORY @PROJECT_BINARY_DIR@/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
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)
endif()
install(FILES @PROJECT_BINARY_DIR@/.buildInfo DESTINATION lib)
#############################
# Packaging steps
@@ -20,12 +27,12 @@ set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR}
set(CPACK_GENERATOR "TGZ;DEB;RPM")
set(CPACK_BINARY_DEB "ON")
set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm")
set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), hcc_lc (= @HCC_VERSION@)")
set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), hcc_lc (= @HCC_VERSION@), rocm-profiler")
set(CPACK_BINARY_RPM "ON")
set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst")
set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm")
set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")
set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, hcc_lc = @HCC_VERSION@")
set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, hcc_lc = @HCC_VERSION@, rocm-profiler")
set(CPACK_SOURCE_GENERATOR "TGZ")
include(CPack)
samples/0_Intro/bit_extract/bit_extract.cpp
+1 -1
View File
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
#include <stdio.h>
#include <iostream>
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#ifdef __HIP_PLATFORM_HCC__
#include <hc.hpp>
#endif
samples/0_Intro/hcc_dialects/vadd_hip.cpp
+1 -1
View File
@@ -1,4 +1,4 @@
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
__global__ void vadd_hip(hipLaunchParm lp, const float *a, const float *b, float *c, int N)
{
samples/0_Intro/module_api/runKernel.cpp
+2 -2
View File
@@ -17,8 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<hip_runtime.h>
#include<hip_runtime_api.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include<iostream>
#include<fstream>
#include<vector>
samples/0_Intro/square/square.hipref.cpp
+1 -1
View File
@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <stdio.h>
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#define CHECK(cmd) \
{\
samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp
+1 -1
View File
@@ -2,7 +2,7 @@
#include <iostream>
#include <sstream>
#include <iomanip>
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#include "ResultDatabase.h"
samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+1 -1
View File
@@ -17,7 +17,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include"hip_runtime.h"
#include "hip/hip_runtime.h"
#include<iostream>
#include<time.h>
#include"ResultDatabase.h"
samples/1_Utils/hipInfo/hipInfo.cpp
+1 -1
View File
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
#include <iostream>
#include <iomanip>
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define KNRM "\x1B[0m"
#define KRED "\x1B[31m"
samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp
+1 -1
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include<iostream>
// hip header file
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define WIDTH 1024
samples/2_Cookbook/0_MatrixTranspose/Readme.md
+1 -1
View File
@@ -95,6 +95,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
- [hipify-clang](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/hipify-clang/README.md)
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
samples/2_Cookbook/1_hipEvent/Readme.md
+2 -2
View File
@@ -17,7 +17,7 @@ Programmers familiar with CUDA, OpenCL will be able to quickly learn and start c
We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to get the performance score for memory transfer and kernel execution time.
## hipEnvent_t
## hipEvent_t
We'll learn how to use the event management functionality of HIP runtime api. In the same sourcecode, we used for MatrixTranspose we will declare the following events as follows:
@@ -69,6 +69,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
- [hipify-clang](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/hipify-clang/README.md)
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
samples/2_Cookbook/1_hipEvent/hipEvent.cpp
+1 -1
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include<iostream>
// hip header file
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define HEIGHT 1024
samples/2_Cookbook/2_HIP_ATP_MARKER/MatrixTranspose.cpp
+1 -1
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include<iostream>
// hip header file
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define HEIGHT 1024
samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md
+1 -1
View File
@@ -46,6 +46,6 @@ Note this trace mode uses colors. "less -r" can handle raw control characters an
- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP)
- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md)
- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL)
- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md)
- [hipify-clang](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/hipify-clang/README.md)
- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md)
- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md)
samples/2_Cookbook/3_shared_memory/sharedMemory.cpp
+1 -1
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include<iostream>
// hip header file
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define WIDTH 1024
samples/2_Cookbook/4_shfl/shfl.cpp
+1 -1
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include<iostream>
// hip header file
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define WIDTH 4
samples/2_Cookbook/5_2dshfl/2dshfl.cpp
+1 -1
View File
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include<iostream>
// hip header file
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define WIDTH 4
samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp
+1 -1
View File
@@ -5,7 +5,7 @@
#include <cmath>
// header file for the GPU API
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#include <hipblas.h>
#define N (1024 * 500)
src/device_util.cpp
Regular → Executable
+16 -4
View File
@@ -17,11 +17,12 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include"hip_runtime.h"
#include<hc.hpp>
#include<grid_launch.h>
#include <hc.hpp>
#include <grid_launch.h>
#include <hc_math.hpp>
#include "hip/hip_runtime.h"
// TODO: Choose whether default is precise math or fast math based on compilation flag.
#ifdef __HCC_ACCELERATOR__
using namespace hc::precise_math;
@@ -1671,6 +1672,17 @@ __device__ unsigned int atomicDec(unsigned int* address,
return hc::__atomic_wrapdec(address,val);
}
//__mul24 __umul24
__device__ int __mul24(int arg1,
int arg2)
{
return hc::__mul24(arg1, arg2);
}
__device__ unsigned int __umul24(unsigned int arg1,
unsigned int arg2)
{
return hc::__mul24(arg1, arg2);
}
__device__ unsigned int test__popc(unsigned int input)
{
src/hip_context.cpp
+3 -3
View File
@@ -22,9 +22,9 @@ THE SOFTWARE.
#include <stack>
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
// Stack of contexts
thread_local std::stack<ihipCtx_t *> tls_ctxStack;
src/hip_device.cpp
+3 -3
View File
@@ -17,9 +17,9 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
//-------------------------------------------------------------------------------------------------
//Devices
src/hip_error.cpp
+3 -3
View File
@@ -17,9 +17,9 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
//-------------------------------------------------------------------------------------------------
//-------------------------------------------------------------------------------------------------
src/hip_event.cpp
+3 -3
View File
@@ -17,9 +17,9 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
//-------------------------------------------------------------------------------------------------
//-------------------------------------------------------------------------------------------------
src/hip_hcc.cpp
+21 -14
View File
@@ -36,15 +36,15 @@ THE SOFTWARE.
#include <vector>
#include <algorithm>
#include <atomic>
#include <hc.hpp>
#include <hc_am.hpp>
#include "hsa/hsa_ext_amd.h"
#include "libhsakmt/hsakmt.h"
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hsa_ext_amd.h"
#include "hsakmt.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
@@ -53,10 +53,6 @@ THE SOFTWARE.
//=================================================================================================
const int release = 1;
#define MEMCPY_D2H_STAGING_VS_PININPLACE_COPY_THRESHOLD 4194304
#define MEMCPY_H2D_DIRECT_VS_STAGING_COPY_THRESHOLD 65336
#define MEMCPY_H2D_STAGING_VS_PININPLACE_COPY_THRESHOLD 1048576
const char *API_COLOR = KGRN;
const char *API_COLOR_END = KNRM;
@@ -69,6 +65,9 @@ int HIP_ATP_MARKER= 0;
int HIP_DB= 0;
int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */
int HIP_NUM_KERNELS_INFLIGHT = 128;
int HIP_BLOCKING_SYNC = 0;
//#define DISABLE_COPY_EXT 1
std::once_flag hip_initialized;
@@ -177,7 +176,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
{
if (! assertQueueEmpty) {
tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this);
crit->_av.wait();
crit->_av.wait(HIP_BLOCKING_SYNC ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
}
crit->_kernelCnt = 0;
@@ -257,8 +256,8 @@ void ihipStream_t::lockclose_postKernelCommand(hc::accelerator_view *av)
{
if (HIP_LAUNCH_BLOCKING) {
// TODO - fix this so it goes through proper stream::wait() call.
av->wait(); // direct wait OK since we know the stream is locked.
// 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());
}
@@ -1008,7 +1007,9 @@ void ihipInit()
READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" );
READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Number of kernels per stream ");
READ_ENV_I(release, HIP_BLOCKING_SYNC, 0, "Use blocking synchronization for stream waits. This may increase latency but is friendlier to other processes. If 0, spin-wait.");
READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced.");
// 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) {
@@ -1451,7 +1452,13 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
{
LockedAccessor_StreamCrit_t crit (_criticalData);
#if DISABLE_COPY_EXT
#warning ("Disabled copy_ext path, P2P host staging copies will not work")
// Note - peer-to-peer copies which require host staging will not work in this path.
crit->_av.copy(src, dst, sizeBytes);
#else
crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine);
#endif
}
}
src/hip_ldg.cpp
+2 -1
View File
@@ -18,7 +18,8 @@ THE SOFTWARE.
*/
#include <hc.hpp>
#include"hcc_detail/hip_ldg.h"
#include "hip/hcc_detail/hip_ldg.h"
__device__ char __ldg(const char* ptr)
{
src/hip_memory.cpp
+7 -5
View File
@@ -17,12 +17,14 @@
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include <hsa.h>
#include <hc_am.hpp>
#include <hsa_ext_amd.h>
#include "hsa/hsa.h"
#include "hsa/hsa_ext_amd.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
//-------------------------------------------------------------------------------------------------
//-------------------------------------------------------------------------------------------------
// Memory
src/hip_module.cpp
+8 -6
View File
@@ -17,17 +17,19 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hsa/hsa.h"
#include "hsa/hsa_ext_amd.h"
#include "hsa/amd_hsa_kernel_code.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include <fstream>
#include <stdio.h>
#include <stdlib.h>
#include <elf.h>
#include "hsa/hsa.h"
#include "hsa/hsa_ext_amd.h"
#include "hsa/amd_hsa_kernel_code.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
//TODO Use Pool APIs from HCC to get memory regions.
namespace hipdrv {
src/hip_peer.cpp
+3 -3
View File
@@ -19,9 +19,9 @@ THE SOFTWARE.
#include <hc_am.hpp>
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
// Peer access functions.
src/hip_stream.cpp
+4 -4
View File
@@ -17,9 +17,9 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_hcc.h"
#include "hip/hcc_detail/trace_helper.h"
//-------------------------------------------------------------------------------------------------
@@ -52,7 +52,7 @@ hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
e = hipErrorInvalidDevice;
}
return ihipLogStatus(e);
return e;
}
tests/clang-hipify/axpy.cu → tests/hipify-clang/axpy.cu
View File
tests/clang-hipify/lit.cfg → tests/hipify-clang/lit.cfg
View File
tests/clang-hipify/lit.site.cfg.in → tests/hipify-clang/lit.site.cfg.in
View File
tests/src/CMakeLists.txt
-243
View File
@@ -1,243 +0,0 @@
cmake_minimum_required(VERSION 2.6)
# remove CMAKE_CXX_COMPILER entry from cache since it will be pointing to hipcc
unset(CMAKE_CXX_COMPILER CACHE)
# remove HIP_PATH entry from cache since we might be running tests with a different configuration
unset(HIP_PATH CACHE)
project(HIP_Unit_Tests)
include(CTest)
set(HIPTEST_SOURCE_DIR ${PROJECT_SOURCE_DIR})
string(ASCII 27 Esc)
set(ColorReset "${Esc}[m")
set(Red "${Esc}[31m")
set(Magenta "${Esc}[35m")
# Enable multi-gpu tests
if(NOT DEFINED HIP_MULTI_GPU)
set(HIP_MULTI_GPU 0 CACHE BOOL "Run tests requiring more than one GPU")
endif()
# We are going to use HIP source...
get_filename_component(HIP_SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/../.. ABSOLUTE)
# ...so we first need to determine the options to cascade to HIP build.
if(DEFINED HIP_PLATFORM)
set(ENV{HIP_PLATFORM} ${HIP_PLATFORM})
endif()
if(DEFINED HCC_HOME)
get_filename_component(HCC_HOME ${HCC_HOME} ABSOLUTE)
set(ENV{HCC_HOME} ${HCC_HOME})
set(ENV{HIP_DEVELOPER} 1)
endif()
if(DEFINED HIP_LIB_TYPE)
set(ENV{HIP_LIB_TYPE} ${HIP_LIB_TYPE})
endif()
# Purge previous HIP installation...
execute_process(
COMMAND "${CMAKE_COMMAND}" -E remove_directory hip
OUTPUT_QUIET
ERROR_QUIET
)
execute_process(
COMMAND "${CMAKE_COMMAND}" -E make_directory hip
OUTPUT_QUIET
ERROR_QUIET
)
message(STATUS "Configuring HIP")
# ...and now build HIP locally.
execute_process(
COMMAND "${CMAKE_COMMAND}" -DCMAKE_INSTALL_PREFIX=${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ${HIP_SRC_PATH}
WORKING_DIRECTORY hip
RESULT_VARIABLE hip_build_result
OUTPUT_VARIABLE hip_build_log
OUTPUT_STRIP_TRAILING_WHITESPACE
ERROR_QUIET
)
message("${Magenta}${hip_build_log}${ColorReset}")
if(hip_build_result)
message(FATAL_ERROR "Error configuring HIP")
else()
message(STATUS "Configuring HIP - done")
message(STATUS "Building HIP")
endif()
execute_process(
COMMAND "${CMAKE_COMMAND}" --build . --target install
WORKING_DIRECTORY hip
RESULT_VARIABLE hip_build_result
OUTPUT_VARIABLE hip_build_log
OUTPUT_STRIP_TRAILING_WHITESPACE
ERROR_QUIET
)
# Show HIP build errors if any.
if(hip_build_result)
message("${Red}${hip_build_log}${ColorReset}")
message(FATAL_ERROR "Error building HIP")
else()
# Building HIP is successful. Point HIP_PATH to this location.
message(STATUS "Building HIP - done")
get_filename_component(HIP_PATH ${CMAKE_CURRENT_BINARY_DIR}/hip/localbuild ABSOLUTE)
endif()
# Add a target to rebuild HIP if HIP source changes.
add_custom_target(
hip ALL
COMMAND "${CMAKE_COMMAND}" --build . --target install
WORKING_DIRECTORY hip
)
# Determine HIP_PLATFORM
execute_process(COMMAND ${HIP_PATH}/bin/hipconfig --platform OUTPUT_VARIABLE HIP_PLATFORM)
if(${HIP_PLATFORM} STREQUAL "hcc")
MESSAGE("HIP_PLATFORM=hcc")
elseif(${HIP_PLATFORM} STREQUAL "nvcc")
MESSAGE("HIP_PLATFORM=nvcc")
#Need C++11 for threads in some of the tests.
add_definitions(-std=c++11)
# NVCC does not not support -rdynamic option
set(CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS)
set(CMAKE_SHARED_LIBRARY_LINK_C_FLAGS)
else()
MESSAGE(FATAL_ERROR "UNKNOWN HIP_PLATFORM=" ${HIP_PLATFORM})
endif()
set(HIPCC ${HIP_PATH}/bin/hipcc)
set(CMAKE_CXX_COMPILER ${HIPCC} CACHE FILEPATH "CXX Compiler" FORCE)
add_library(test_common OBJECT test_common.cpp)
# usage : build_hip_executable(exe_name CPP_FILES)
macro(build_hip_executable exe cpp)
add_executable(${exe} ${cpp} ${ARGN} $<TARGET_OBJECTS:test_common>)
add_dependencies(${exe} hip)
endmacro()
# Make a hip executable, using libc++
macro(build_hip_executable_libcpp exe cpp)
build_hip_executable( ${exe} ${cpp} ${ARGN})
if(${HIP_PLATFORM} STREQUAL "hcc")
set_source_files_properties(${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --stdlib=libc++)
endif()
endmacro()
function(make_named_test exe testname)
add_test(NAME ${testname}
COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
)
set_tests_properties(${testname}
PROPERTIES PASS_REGULAR_EXPRESSION "PASSED"
)
endfunction()
macro(make_test exe)
string(REPLACE " " "" smush_args ${ARGN})
set(testname ${PROJECT_NAME}/${exe}${smush_args}.tst)
make_named_test(${exe} ${testname} ${ARGN})
endmacro()
macro(make_hipify_test sourceFile)
#string(REPLACE " " "" smush_args ${ARGN})
set(testname ${sourceFile}${smush_args}.tst)
add_test(NAME ${testname}
COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN}
)
endmacro()
macro(make_test_matches exe match_string)
string(REPLACE " " "" smush_args ${ARGN})
set(testname ${exe}${smush_args}.tst)
add_test(NAME ${testname}
COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
)
set_tests_properties(${testname}
PROPERTIES PASS_REGULAR_EXPRESSION ${match_string}
)
endmacro()
macro(build_hip_executable_sm35 exe cpp)
build_hip_executable( ${exe} ${cpp} ${ARGN})
if(${HIP_PLATFORM} STREQUAL "nvcc")
set_source_files_properties(${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --gpu-architecture=sm_35)
endif()
endmacro()
build_hip_executable(hipGetDeviceAttribute hipGetDeviceAttribute.cpp)
build_hip_executable(hipEnvVar hipEnvVar.cpp)
build_hip_executable(hipEnvVarDriver hipEnvVarDriver.cpp)
build_hip_executable(hipEventRecord hipEventRecord.cpp)
build_hip_executable_libcpp(hipHcc hipHcc.cpp)
#set_source_files_properties(hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++)
# __workweek fix.
#build_hip_executable_libcpp(hipPointerAttrib hipPointerAttrib.cpp)
build_hip_executable(hipHostMalloc hipHostMalloc.cpp)
build_hip_executable(hipHostGetFlags hipHostGetFlags.cpp)
build_hip_executable(hipHostRegister hipHostRegister.cpp)
build_hip_executable(hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
build_hip_executable(hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp)
build_hip_executable(hipFuncGetDevice hipFuncGetDevice.cpp)
build_hip_executable(hipFuncSetDevice hipFuncSetDevice.cpp)
build_hip_executable(hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
build_hip_executable(hipPeerToPeer_simple hipPeerToPeer_simple.cpp)
build_hip_executable(hipTestMemcpyPin hipTestMemcpyPin.cpp)
build_hip_executable(hipDynamicShared hipDynamicShared.cpp)
build_hip_executable(hipLaunchParm hipLaunchParm.cpp)
if(${HIP_PLATFORM} STREQUAL "hcc")
build_hip_executable(hipArray hipArray.cpp)
endif()
make_test(hipEventRecord --iterations 10)
make_test(hipEnvVarDriver " ")
make_test(hipLaunchParm " ")
#TODO -reenable
#make_test(hipPointerAttrib " ")
make_test(hipHostMalloc " ")
# BS- comment out since test appears broken - asks for device pointer but pointer was never allocated.
#make_test(hipHostGetFlags " ")
make_test(hipHcc " ")
make_test(hipHostRegister " ")
make_test(hipRandomMemcpyAsync " ")
make_test(hipFuncSetDeviceFlags " ")
make_test(hipFuncGetDevice " ")
make_test(hipFuncDeviceSynchronize " ")
make_test(hipTestMemcpyPin " ")
if(${HIP_MULTI_GPU})
make_test(hipPeerToPeer_simple " ") # use current device for copy, this fails.
make_test(hipPeerToPeer_simple --memcpyWithPeer)
make_test(hipPeerToPeer_simple --mirrorPeers) # mirror mapping: test to ensure mirror doesn't destroy orig mapping.
endif()
if(${HIP_PLATFORM} STREQUAL "hcc")
make_test(hipArray " ")
make_test(hipFuncSetDevice " ")
make_test(hipDynamicShared " ")
endif()
make_hipify_test(specialFunc.cu)
# Add subdirs here:
add_subdirectory(context)
add_subdirectory(deviceLib)
add_subdirectory(runtimeApi)
add_subdirectory(kernel)
# vim: ts=4:sw=4:expandtab:smartindent
tests/src/Functional/Negative/Device/hipDeviceGetAttribute.cpp
+2 -2
View File
@@ -1,5 +1,5 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include"hipDeviceUtil.h"
int main()
tests/src/Functional/Negative/Device/hipDeviceUtil.h
+1 -1
View File
@@ -1,7 +1,7 @@
#ifndef HIPDEVICEUTIL_H
#define HIPDEVICEUTIL_H
#include<hip/hip_runtime_api.h>
#include "hip/hip_runtime_api.h"
#include<iostream>
#define HIP_CHECK(status, func) \
tests/src/Functional/Negative/Device/hipGetDevice.cpp
+2 -2
View File
@@ -1,5 +1,5 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include"hipDeviceUtil.h"
int main()
tests/src/Functional/Negative/Device/hipGetDeviceCount.cpp
+2 -2
View File
@@ -1,5 +1,5 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include"hipDeviceUtil.h"
int main()
tests/src/Functional/Negative/Device/hipGetDeviceProperties.cpp
+2 -2
View File
@@ -1,5 +1,5 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include"hipDeviceUtil.h"
int main()
tests/src/Functional/Negative/Device/hipSetDevice.cpp
+2 -2
View File
@@ -1,5 +1,5 @@
#include<hip/hip_runtime.h>
#include<hip/hip_runtime_api.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include"hipDeviceUtil.h"
int main()
tests/src/context/CMakeLists.txt
-9
View File
@@ -1,9 +0,0 @@
cmake_minimum_required (VERSION 2.6)
# Functions for kernel attributes (grid_launch, __launch_bounds__, etc)
project (kernel)
include_directories( ${HIPTEST_SOURCE_DIR} )
build_hip_executable_libcpp (hipCtx_simple hipCtx_simple.cpp)
make_test(hipCtx_simple " " )
tests/src/context/hipCtx_simple.cpp
+1 -1
View File
@@ -26,7 +26,7 @@ THE SOFTWARE.
* HIT_END
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
int main(int argc, char *argv[])
tests/src/deviceLib/CMakeLists.txt
-46
View File
@@ -1,46 +0,0 @@
cmake_minimum_required (VERSION 2.6)
project (deviceLib)
include_directories( ${HIPTEST_SOURCE_DIR} )
build_hip_executable (hip_ballot hip_ballot.cpp)
make_test(hip_ballot " " )
build_hip_executable (hip_anyall hip_anyall.cpp)
make_test(hip_anyall " " )
build_hip_executable (hip_popc hip_popc.cpp)
make_test(hip_popc " " )
build_hip_executable (hip_clz hip_clz.cpp)
make_test(hip_clz " " )
build_hip_executable (hip_brev hip_brev.cpp)
make_test(hip_brev " " )
build_hip_executable (hip_ffs hip_ffs.cpp)
make_test(hip_ffs " " )
build_hip_executable_sm35 (hip_test_ldg hip_test_ldg.cpp)
make_test(hip_test_ldg " " )
build_hip_executable (hipSimpleAtomicsTest hipSimpleAtomicsTest.cpp)
make_test(hipSimpleAtomicsTest " ")
build_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp)
make_test(hipMathFunctionsHost " ")
build_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp)
make_test(hipMathFunctionsDevice " ")
build_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp)
make_test(hipIntrinsics " ")
build_hip_executable (hipTestDevice hipTestDevice.cpp)
make_test(hipTestDevice " ")
build_hip_executable (hipTestDeviceDouble hipTestDeviceDouble.cpp)
make_test(hipTestDeviceDouble " ")
tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
tests/src/deviceLib/hipIntegerIntrinsics.cpp
+3 -3
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
@@ -35,7 +35,7 @@ __device__ void integer_intrinsics()
__ffs((int)10);
__ffsll((long long)10);
//__hadd((int)1, (int)3);
//__mul24((int)1, (int)2);
__mul24((int)1, (int)2);
//__mul64hi((long long)1, (long long)2);
//__mulhi((int)1, (int)2);
__popc((unsigned int)4);
@@ -45,7 +45,7 @@ __device__ void integer_intrinsics()
//__rhadd((int)1, (int)2);
//__sad((int)1, (int)2, 0);
//__uhadd((unsigned int)1, (unsigned int)3);
//__umul24((unsigned int)1, (unsigned int)2);
__umul24((unsigned int)1, (unsigned int)2);
//__umul64hi((unsigned long long)1, (unsigned long long)2);
//__umulhi((unsigned int)1, (unsigned int)2);
//__urhadd((unsigned int)1, (unsigned int)2);
tests/src/deviceLib/hipMathFunctions.cpp
+1 -1
View File
@@ -30,7 +30,7 @@ THE SOFTWARE.
* HIT_END
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
int main(int argc, char *argv[])
tests/src/deviceLib/hipSimpleAtomicsTest.cpp
+1 -1
View File
@@ -31,7 +31,7 @@ THE SOFTWARE.
// Includes HIP Runtime
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#include <test_common.h>
#define EXIT_WAIVED 2
tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
+1 -1
View File
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "test_common.h"
#pragma GCC diagnostic ignored "-Wall"
tests/src/deviceLib/hipTestDevice.cpp
+2 -2
View File
@@ -24,8 +24,8 @@ THE SOFTWARE.
*/
#include"test_common.h"
#include"hip_runtime.h"
#include"hip_runtime_api.h"
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#define N 512
#define SIZE N*sizeof(float)
tests/src/deviceLib/hipTestDeviceDouble.cpp
+2 -2
View File
@@ -24,8 +24,8 @@ THE SOFTWARE.
*/
#include"test_common.h"
#include"hip_runtime.h"
#include"hip_runtime_api.h"
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#define N 512
#define SIZE N*sizeof(double)
tests/src/deviceLib/hip_anyall.cpp
+1 -1
View File
@@ -29,7 +29,7 @@ THE SOFTWARE.
#include <stdio.h>
#include <iostream>
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#define HIP_ASSERT(x) (assert((x)==hipSuccess))
__global__ void
tests/src/deviceLib/hip_ballot.cpp
+1 -1
View File
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include <iostream>
#include <hip_runtime.h>
#include "hip/hip_runtime.h"
#define HIP_ASSERT(x) (assert((x)==hipSuccess))
__global__ void
tests/src/deviceLib/hip_brev.cpp
+1 -1
View File
@@ -31,7 +31,7 @@ THE SOFTWARE.
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#define HIP_ASSERT(x) (assert((x)==hipSuccess))

Some files were not shown because too many files have changed in this diff Show More