AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

Merge 'master' into 'amd-master'

Explorar el Código 
Change-Id: Ia44cb648c8ae16ff424bcf64fb9a9338c7678c0c


[ROCm/hip commit: 425254aa13]
Este commit está contenido en:
Jenkins
2019-11-07 05:22:19 -05:00
padre 4173159cd1 4c66b0a5a9
commit 62a0771088
Se han modificado 44 ficheros con 1958 adiciones y 1147 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
projects/hip/CMakeLists.txt
+8 -2
Ver fichero
@@ -163,11 +163,14 @@ if(NOT CMAKE_BUILD_TYPE)
endif()
# Determine HIP install path
if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT AND CMAKE_INSTALL_PREFIX MATCHES "/usr/local")
if (UNIX)
set(HIP_DEFAULT_INSTALL_PREFIX "/opt/rocm/hip")
endif()
if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
if(CMAKE_BUILD_TYPE MATCHES Debug)
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 ${HIP_DEFAULT_INSTALL_PREFIX} CACHE PATH "Installation path for HIP" FORCE)
else()
message(FATAL_ERROR "Invalid CMAKE_BUILD_TYPE specified. Valid values are Debug and Release")
endif()
@@ -184,7 +187,10 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
set(COMPILE_HIP_ATP_MARKER 0)
else()
set(COMPILE_HIP_ATP_MARKER $ENV{COMPILE_HIP_ATP_MARKER})
message(WARNING "HIP Markers are deprecated, please use roctracer/rocTX marker APIs.")
endif()
else()
message(WARNING "HIP Markers are deprecated, please use roctracer/rocTX marker APIs.")
endif()
add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
projects/hip/hipify-clang/README.md
+79 -77
Ver fichero
@@ -146,9 +146,9 @@ To run it:
* Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:
- Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38`
- Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.5.32`
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38`
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.5.32`
5. Ensure [`CUB`](https://github.com/NVlabs/cub) of the version corresponding to CUDA's version is installed.
@@ -190,9 +190,9 @@ To run it:
On Linux the following configurations are tested:
Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.4.38
Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.5.32
Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.4.38
Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.5.32
Minimum build system requirements for the above configurations:
@@ -207,7 +207,7 @@ cmake
-DCMAKE_INSTALL_PREFIX=../dist \
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38 \
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.5.32 \
-DCUDA_CUB_ROOT_DIR=/srv/git/CUB \
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \
..
@@ -264,74 +264,76 @@ Linux 5.2.0 - Platform OS
64 - hipify-clang binary bitness
64 - python 2.7.12 binary bitness
========================================
-- Testing: 65 tests, 12 threads --
PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 65)
PASS: hipify :: unit_tests/device/math_functions.cu (2 of 65)
PASS: hipify :: unit_tests/device/atomics.cu (3 of 65)
PASS: hipify :: unit_tests/device/device_symbols.cu (4 of 65)
PASS: hipify :: unit_tests/headers/headers_test_02.cu (5 of 65)
PASS: hipify :: unit_tests/headers/headers_test_03.cu (6 of 65)
PASS: hipify :: unit_tests/headers/headers_test_01.cu (7 of 65)
PASS: hipify :: unit_tests/headers/headers_test_04.cu (8 of 65)
PASS: hipify :: unit_tests/headers/headers_test_05.cu (9 of 65)
PASS: hipify :: unit_tests/headers/headers_test_07.cu (10 of 65)
PASS: hipify :: unit_tests/headers/headers_test_06.cu (11 of 65)
PASS: hipify :: unit_tests/headers/headers_test_11.cu (12 of 65)
PASS: hipify :: unit_tests/headers/headers_test_08.cu (13 of 65)
PASS: hipify :: unit_tests/headers/headers_test_10.cu (14 of 65)
PASS: hipify :: unit_tests/headers/headers_test_09.cu (15 of 65)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (16 of 65)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (17 of 65)
PASS: hipify :: unit_tests/libraries/CUB/cub_01.cu (18 of 65)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (19 of 65)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (20 of 65)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (21 of 65)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (22 of 65)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (23 of 65)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (24 of 65)
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (25 of 65)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (26 of 65)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (27 of 65)
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (28 of 65)
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (29 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (30 of 65)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (31 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (32 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (33 of 65)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (34 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (35 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (36 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (37 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (38 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (39 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (40 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (41 of 65)
PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (42 of 65)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (43 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (44 of 65)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (45 of 65)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (46 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (47 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (48 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (49 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (50 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (52 of 65)
PASS: hipify :: unit_tests/samples/MallocManaged.cpp (53 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (54 of 65)
PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (55 of 65)
PASS: hipify :: unit_tests/samples/allocators.cu (56 of 65)
PASS: hipify :: unit_tests/samples/coalescing.cu (57 of 65)
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (58 of 65)
PASS: hipify :: unit_tests/samples/axpy.cu (59 of 65)
PASS: hipify :: unit_tests/samples/cudaRegister.cu (60 of 65)
PASS: hipify :: unit_tests/samples/intro.cu (61 of 65)
PASS: hipify :: unit_tests/samples/square.cu (62 of 65)
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (63 of 65)
PASS: hipify :: unit_tests/samples/vec_add.cu (64 of 65)
PASS: hipify :: unit_tests/libraries/CUB/cub_02.cu (18 of 65)
Testing Time: 3.01s
Expected Passes : 65
-- Testing: 67 tests, 12 threads --
PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 67)
PASS: hipify :: unit_tests/device/math_functions.cu (2 of 67)
PASS: hipify :: unit_tests/device/atomics.cu (3 of 67)
PASS: hipify :: unit_tests/device/device_symbols.cu (4 of 67)
PASS: hipify :: unit_tests/headers/headers_test_01.cu (5 of 67)
PASS: hipify :: unit_tests/headers/headers_test_02.cu (6 of 67)
PASS: hipify :: unit_tests/headers/headers_test_03.cu (7 of 67)
PASS: hipify :: unit_tests/headers/headers_test_05.cu (8 of 67)
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 67)
PASS: hipify :: unit_tests/headers/headers_test_06.cu (10 of 67)
PASS: hipify :: unit_tests/headers/headers_test_07.cu (11 of 67)
PASS: hipify :: unit_tests/headers/headers_test_10.cu (12 of 67)
PASS: hipify :: unit_tests/headers/headers_test_11.cu (13 of 67)
PASS: hipify :: unit_tests/headers/headers_test_08.cu (14 of 67)
PASS: hipify :: unit_tests/kernel_launch/kernel_launch_01.cu (15 of 67)
PASS: hipify :: unit_tests/headers/headers_test_09.cu (16 of 67)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (17 of 67)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (18 of 67)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (19 of 67)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (20 of 67)
PASS: hipify :: unit_tests/libraries/CUB/cub_03.cu (21 of 67)
PASS: hipify :: unit_tests/libraries/CUB/cub_01.cu (22 of 67)
PASS: hipify :: unit_tests/libraries/CUB/cub_02.cu (23 of 67)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (24 of 67)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (25 of 67)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (26 of 67)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (27 of 67)
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (28 of 67)
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (29 of 67)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (30 of 67)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (31 of 67)
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (32 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (33 of 67)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (34 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (35 of 67)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (36 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (37 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (38 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (39 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (40 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (41 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (42 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (43 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (44 of 67)
PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (45 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (46 of 67)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (47 of 67)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (48 of 67)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (49 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (50 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (51 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (52 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (53 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (54 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (55 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (56 of 67)
PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (57 of 67)
PASS: hipify :: unit_tests/samples/MallocManaged.cpp (58 of 67)
PASS: hipify :: unit_tests/samples/allocators.cu (59 of 67)
PASS: hipify :: unit_tests/samples/coalescing.cu (60 of 67)
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (61 of 67)
PASS: hipify :: unit_tests/samples/axpy.cu (62 of 67)
PASS: hipify :: unit_tests/samples/intro.cu (63 of 67)
PASS: hipify :: unit_tests/samples/cudaRegister.cu (64 of 67)
PASS: hipify :: unit_tests/samples/square.cu (65 of 67)
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (66 of 67)
PASS: hipify :: unit_tests/samples/vec_add.cu (67 of 67)
Testing Time: 3.07s
Expected Passes : 67
[100%] Built target test-hipify
```
### <a name="windows"></a >Windows
@@ -340,13 +342,13 @@ On Windows 10 the following configurations are tested:
LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn 5.1.10 - 7.1.4.18
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn 7.0.5.15 - 7.6.4.38
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn 7.0.5.15 - 7.6.5.32
LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn 7.0.5.15 - 7.6.4.38
LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn 7.0.5.15 - 7.6.5.32
Build system requirements for the latest configuration LLVM 9.0.0/CUDA 10.1 Update 2:
Python 3.6.0 - 3.8.0, cmake 3.5.1 - 3.15.5, Visual Studio 2017 (15.5.2) - 2019 (16.3.5).
Python 3.6.0 - 3.8.0, cmake 3.5.1 - 3.16.0, Visual Studio 2017 (15.5.2) - 2019 (16.3.8).
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`:
@@ -360,7 +362,7 @@ cmake
-DCMAKE_PREFIX_PATH=f:/LLVM/9.0.0/dist \
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38 \
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.5.32 \
-DCUDA_CUB_ROOT_DIR=f:/GIT/cub \
-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \
-Thost=x64
projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp
+12 -12
Ver fichero
@@ -117,7 +117,7 @@ namespace perl {
{sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam}
};
void generateHeader(unique_ptr<ostream>& streamPtr) {
void generateHeader(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << "#!/usr/bin/perl -w" << endl_2;
*streamPtr.get() << sCopyright << endl;
*streamPtr.get() << sImportant << endl_2;
@@ -148,7 +148,7 @@ namespace perl {
*streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << endl_2;
}
void generateStatFunctions(unique_ptr<ostream>& streamPtr) {
void generateStatFunctions(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "totalStats" << " {" << endl;
*streamPtr.get() << tab << my << "%count = %{ shift() };" << endl;
*streamPtr.get() << tab << my << "$total = 0;" << endl;
@@ -175,11 +175,11 @@ namespace perl {
}
}
void generateSimpleSubstitutions(unique_ptr<ostream>& streamPtr) {
void generateSimpleSubstitutions(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "simpleSubstitutions" << " {" << endl;
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) {
for (auto &ma : CUDA_INCLUDE_MAP) {
if (Statistics::isUnsupported(ma.second)) continue;
if (i == ma.second.type) {
string sCUDA = ma.first.str();
@@ -190,7 +190,7 @@ namespace perl {
}
}
} else {
for (auto& ma : CUDA_RENAMES_MAP()) {
for (auto &ma : CUDA_RENAMES_MAP()) {
if (Statistics::isUnsupported(ma.second)) continue;
if (i == ma.second.type) {
*streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << endl;
@@ -201,7 +201,7 @@ namespace perl {
*streamPtr.get() << "}" << endl;
}
void generateExternShared(unique_ptr<ostream>& streamPtr) {
void generateExternShared(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << endl;
*streamPtr.get() << sub << "transformExternShared" << " {" << endl;
*streamPtr.get() << tab << no_warns << endl;
@@ -210,7 +210,7 @@ namespace perl {
*streamPtr.get() << tab << "$ft{'extern_shared'} += $k;" << endl << "}" << endl;
}
void generateKernelLaunch(unique_ptr<ostream>& streamPtr) {
void generateKernelLaunch(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << "# CUDA Kernel Launch Syntax" << endl << sub << "transformKernelLaunch" << " {" << endl;
*streamPtr.get() << tab << no_warns << endl;
*streamPtr.get() << tab << my_k << endl_2;
@@ -251,13 +251,13 @@ namespace perl {
*streamPtr.get() << tab_2 << "$Tkernels{$1}++;" << endl_tab << "}" << endl << "}" << endl;
}
void generateCubNamespace(unique_ptr<ostream>& streamPtr) {
void generateCubNamespace(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "transformCubNamespace" << " {" << endl_tab << my_k << endl;
*streamPtr.get() << tab << "$k += s/using\\s*namespace\\s*cub/using namespace hipcub/g;" << endl;
*streamPtr.get() << tab << "$k += s/\\bcub::\\b/hipcub::/g;" << endl << tab << return_k << "}" << endl;
}
void generateHostFunctions(unique_ptr<ostream>& streamPtr) {
void generateHostFunctions(unique_ptr<ostream> &streamPtr) {
*streamPtr.get() << endl << sub << "transformHostFunctions" << " {" << endl_tab << my_k << endl;
set<string> &funcSet = DeviceSymbolFunctions0;
const string s0 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
@@ -271,7 +271,7 @@ namespace perl {
default: funcSet = DeviceSymbolFunctions0;
}
unsigned int count = 0;
for (auto& f : funcSet) {
for (auto &f : funcSet) {
const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f);
if (found != CUDA_RUNTIME_FUNCTION_MAP.end()) {
*streamPtr.get() << (count ? ",\n" : "") << tab_2 << "\"" << found->second.hipName.str() << "\"";
@@ -291,12 +291,12 @@ namespace perl {
*streamPtr.get() << tab << return_k << "}" << endl;
}
void generateDeviceFunctions(unique_ptr<ostream>& streamPtr) {
void generateDeviceFunctions(unique_ptr<ostream> &streamPtr) {
unsigned int countUnsupported = 0;
unsigned int countSupported = 0;
stringstream sSupported;
stringstream sUnsupported;
for (auto& ma : CUDA_DEVICE_FUNC_MAP) {
for (auto &ma : CUDA_DEVICE_FUNC_MAP) {
bool isUnsupported = Statistics::isUnsupported(ma.second);
(isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << tab_2 << "\"" << ma.first.str() << "\"";
if (isUnsupported) countUnsupported++;
projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp
+3 -3
Ver fichero
@@ -60,9 +60,9 @@ namespace python {
*pythonStreamPtr.get() << "from pyHIPIFY.constants import *\n\n";
*pythonStreamPtr.get() << "CUDA_RENAMES_MAP = collections.OrderedDict([\n";
const std::string sHIP_UNS = ", HIP_UNSUPPORTED";
for (int i = 0; i < NUM_CONV_TYPES; i++) {
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) {
for (auto &ma : CUDA_INCLUDE_MAP) {
if (i == ma.second.type) {
std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) {
@@ -74,7 +74,7 @@ namespace python {
}
}
else {
for (auto& ma : CUDA_RENAMES_MAP()) {
for (auto &ma : CUDA_RENAMES_MAP()) {
if (i == ma.second.type) {
std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) {
projects/hip/hipify-clang/src/HipifyAction.cpp
+19 -11
Ver fichero
@@ -150,7 +150,7 @@ void HipifyAction::RewriteToken(const clang::Token &t) {
void HipifyAction::FindAndReplace(StringRef name,
clang::SourceLocation sl,
const std::map<StringRef, hipCounter> &repMap,
const std::map<StringRef, hipCounter> &repMap,
bool bReplace) {
const auto found = repMap.find(name);
if (found == repMap.end()) {
@@ -198,7 +198,7 @@ clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRa
// If the range is contained within a macro, update the macro definition.
// Otherwise, use the file location and hope for the best.
if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) {
return {SM.getFileLoc(begin), SM.getFileLoc(end)};
return {SM.getExpansionLoc(begin), SM.getExpansionLoc(end)};
}
return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
}
@@ -390,16 +390,24 @@ bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result)
OS << readSourceText(*SM, {argStart, argEnd});
}
OS << ")";
clang::SourceRange replacementRange = getWriteRange(*SM, {llcompat::getBeginLoc(launchKernel), llcompat::getEndLoc(launchKernel)});
clang::SourceLocation launchStart = replacementRange.getBegin();
clang::SourceLocation launchKernelExprLocBeg = launchKernel->getExprLoc();
clang::SourceLocation launchKernelExprLocEnd = launchKernelExprLocBeg.isMacroID() ? llcompat::getEndOfExpansionRangeForLoc(*SM, launchKernelExprLocBeg) : llcompat::getEndLoc(launchKernel);
clang::SourceLocation launchKernelEnd = llcompat::getEndLoc(launchKernel);
clang::BeforeThanCompare<clang::SourceLocation> isBefore(*SM);
launchKernelExprLocEnd = isBefore(launchKernelEnd, launchKernelExprLocEnd) ? launchKernelExprLocEnd : launchKernelEnd;
clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernelExprLocBeg, launchKernelExprLocEnd});
clang::SourceLocation launchBeg = replacementRange.getBegin();
clang::SourceLocation launchEnd = replacementRange.getEnd();
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart);
ct::Replacement Rep(*SM, launchStart, length, OS.str());
clang::FullSourceLoc fullSL(launchStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str());
return true;
if (isBefore(launchBeg, launchEnd)) {
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchBeg);
ct::Replacement Rep(*SM, launchBeg, length, OS.str());
clang::FullSourceLoc fullSL(launchBeg, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str());
return true;
}
return false;
}
bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) {
projects/hip/hipify-clang/src/LLVMCompat.cpp
+15 -15
Ver fichero
@@ -40,7 +40,7 @@ void PrintStackTraceOnErrorSignal() {
#endif
}
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file) {
ct::Replacements &getReplacements(ct::RefactoringTool &Tool, StringRef file) {
#if LLVM_VERSION_MAJOR > 3
// getReplacements() now returns a map from filename to Replacements - so create an entry
// for this source file and return a reference to it.
@@ -50,7 +50,7 @@ ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file) {
#endif
}
void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep) {
void insertReplacement(ct::Replacements &replacements, const ct::Replacement &rep) {
#if LLVM_VERSION_MAJOR > 3
// New clang added error checking to Replacements, and *insists* that you explicitly check it.
llvm::consumeError(replacements.add(rep));
@@ -60,7 +60,7 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re
#endif
}
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
void EnterPreprocessorTokenStream(clang::Preprocessor &_pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
_pp.EnterTokenStream(start, len, false, DisableMacroExpansion);
#else
@@ -72,7 +72,7 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *
#endif
}
clang::SourceLocation getBeginLoc(const clang::Stmt* stmt) {
clang::SourceLocation getBeginLoc(const clang::Stmt *stmt) {
#if LLVM_VERSION_MAJOR < 8
return stmt->getLocStart();
#else
@@ -80,7 +80,7 @@ clang::SourceLocation getBeginLoc(const clang::Stmt* stmt) {
#endif
}
clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc) {
clang::SourceLocation getBeginLoc(const clang::TypeLoc &typeLoc) {
#if LLVM_VERSION_MAJOR < 8
return typeLoc.getLocStart();
#else
@@ -88,7 +88,7 @@ clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc) {
#endif
}
clang::SourceLocation getEndLoc(const clang::Stmt* stmt) {
clang::SourceLocation getEndLoc(const clang::Stmt *stmt) {
#if LLVM_VERSION_MAJOR < 8
return stmt->getLocEnd();
#else
@@ -96,7 +96,7 @@ clang::SourceLocation getEndLoc(const clang::Stmt* stmt) {
#endif
}
clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc) {
clang::SourceLocation getEndLoc(const clang::TypeLoc &typeLoc) {
#if LLVM_VERSION_MAJOR < 8
return typeLoc.getLocEnd();
#else
@@ -127,14 +127,6 @@ bool pragma_once_outside_header() {
#endif
}
bool canCompileHostAndDeviceInOneJob() {
#if LLVM_VERSION_MAJOR >= 9 && defined(_WIN32)
return true;
#else
return false;
#endif
}
void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI) {
#if LLVM_VERSION_MAJOR > 9
clang::PreprocessorOptions &PPOpts = CI.getPreprocessorOpts();
@@ -151,4 +143,12 @@ bool CheckCompatibility() {
return true;
}
clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager &SM, const clang::SourceLocation &loc) {
#if LLVM_VERSION_MAJOR > 6
return SM.getExpansionRange(loc).getEnd();
#else
return SM.getExpansionRange(loc).second;
#endif
}
} // namespace llcompat
projects/hip/hipify-clang/src/LLVMCompat.h
+9 -9
Ver fichero
@@ -49,11 +49,11 @@ namespace llcompat {
#define LLVM_DEBUG(X) DEBUG(X)
#endif
clang::SourceLocation getBeginLoc(const clang::Stmt* stmt);
clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc);
clang::SourceLocation getBeginLoc(const clang::Stmt *stmt);
clang::SourceLocation getBeginLoc(const clang::TypeLoc &typeLoc);
clang::SourceLocation getEndLoc(const clang::Stmt* stmt);
clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc);
clang::SourceLocation getEndLoc(const clang::Stmt *stmt);
clang::SourceLocation getEndLoc(const clang::TypeLoc &typeLoc);
void PrintStackTraceOnErrorSignal();
@@ -65,17 +65,17 @@ using namespace llvm;
* Older LLVM versions don't actually support multiple filenames, so everything all gets
* smushed together. It is the caller's responsibility to cope with this.
*/
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file);
ct::Replacements &getReplacements(ct::RefactoringTool &Tool, StringRef file);
/**
* Add a Replacement to a Replacements.
*/
void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep);
void insertReplacement(ct::Replacements &replacements, const ct::Replacement &rep);
/**
* Version-agnostic version of Preprocessor::EnterTokenStream().
*/
void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
void EnterPreprocessorTokenStream(clang::Preprocessor &_pp,
const clang::Token *start,
size_t len,
bool DisableMacroExpansion);
@@ -85,10 +85,10 @@ std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
bool pragma_once_outside_header();
bool canCompileHostAndDeviceInOneJob();
void RetainExcludedConditionalBlocks(clang::CompilerInstance &CI);
bool CheckCompatibility();
clang::SourceLocation getEndOfExpansionRangeForLoc(const clang::SourceManager &SM, const clang::SourceLocation &loc);
} // namespace llcompat
projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h
+3 -3
Ver fichero
@@ -36,15 +36,15 @@ namespace ct = clang::tooling;
*/
template <typename T>
class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory {
ct::Replacements* replacements;
ct::Replacements *replacements;
public:
explicit ReplacementsFrontendActionFactory(ct::Replacements* r):
explicit ReplacementsFrontendActionFactory(ct::Replacements *r):
ct::FrontendActionFactory(),
replacements(r) {}
#if LLVM_VERSION_MAJOR < 10
clang::FrontendAction* create() override {
clang::FrontendAction *create() override {
return new T(replacements);
}
#else
projects/hip/hipify-clang/src/Statistics.cpp
+17 -17
Ver fichero
@@ -163,51 +163,51 @@ void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string
} // Anonymous namespace
void StatCounter::incrementCounter(const hipCounter& counter, const std::string& name) {
void StatCounter::incrementCounter(const hipCounter &counter, const std::string &name) {
counters[name]++;
apiCounters[(int) counter.apiType]++;
convTypeCounters[(int) counter.type]++;
}
void StatCounter::add(const StatCounter& other) {
for (const auto& p : other.counters) {
void StatCounter::add(const StatCounter &other) {
for (const auto &p : other.counters) {
counters[p.first] += p.second;
}
for (int i = 0; i < NUM_API_TYPES; i++) {
for (int i = 0; i < NUM_API_TYPES; ++i) {
apiCounters[i] += other.apiCounters[i];
}
for (int i = 0; i < NUM_CONV_TYPES; i++) {
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
convTypeCounters[i] += other.convTypeCounters[i];
}
}
int StatCounter::getConvSum() {
int acc = 0;
for (const int& i : convTypeCounters) {
for (const int &i : convTypeCounters) {
acc += i;
}
return acc;
}
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix) {
for (int i = 0; i < NUM_CONV_TYPES; i++) {
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string &prefix) {
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (convTypeCounters[i] > 0) {
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
break;
}
}
for (int i = 0; i < NUM_CONV_TYPES; i++) {
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
if (convTypeCounters[i] > 0) {
printStat(csv, printOut, counterNames[i], convTypeCounters[i]);
}
}
for (int i = 0; i < NUM_API_TYPES; i++) {
for (int i = 0; i < NUM_API_TYPES; ++i) {
if (apiCounters[i] > 0) {
conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n");
break;
}
}
for (int i = 0; i < NUM_API_TYPES; i++) {
for (int i = 0; i < NUM_API_TYPES; ++i) {
if (apiCounters[i] > 0) {
printStat(csv, printOut, apiNames[i], apiCounters[i]);
}
@@ -220,7 +220,7 @@ void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const st
}
}
Statistics::Statistics(const std::string& name): fileName(name) {
Statistics::Statistics(const std::string &name): fileName(name) {
// Compute the total bytes/lines in the input file.
std::ifstream src_file(name, std::ios::binary | std::ios::ate);
src_file.clear();
@@ -235,7 +235,7 @@ Statistics::Statistics(const std::string& name): fileName(name) {
///////// Counter update routines //////////
void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) {
void Statistics::incrementCounter(const hipCounter &counter, const std::string &name) {
if (Statistics::isUnsupported(counter)) {
unsupported.incrementCounter(counter, name);
} else {
@@ -308,7 +308,7 @@ void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut)
Statistics globalStats = getAggregate();
// A file is considered "converted" if we made any changes to it.
int convertedFiles = 0;
for (const auto& p : stats) {
for (const auto &p : stats) {
if (p.second.touchedLines && p.second.totalBytes &&
p.second.totalLines && !p.second.hasErrors) {
convertedFiles++;
@@ -326,18 +326,18 @@ void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut)
Statistics Statistics::getAggregate() {
Statistics globalStats("GLOBAL");
for (const auto& p : stats) {
for (const auto &p : stats) {
globalStats.add(p.second);
}
return globalStats;
}
Statistics& Statistics::current() {
Statistics &Statistics::current() {
assert(Statistics::currentStatistics);
return *Statistics::currentStatistics;
}
void Statistics::setActive(const std::string& name) {
void Statistics::setActive(const std::string &name) {
stats.emplace(std::make_pair(name, Statistics{name}));
Statistics::currentStatistics = &stats.at(name);
}
projects/hip/hipify-clang/src/Statistics.h
+7 -7
Ver fichero
@@ -172,11 +172,11 @@ private:
int convTypeCounters[NUM_CONV_TYPES] = {};
public:
void incrementCounter(const hipCounter& counter, const std::string& name);
void incrementCounter(const hipCounter &counter, const std::string &name);
// Add the counters from `other` onto the counters of this object.
void add(const StatCounter& other);
void add(const StatCounter &other);
int getConvSum();
void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix);
void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string &prefix);
};
/**
@@ -195,8 +195,8 @@ class Statistics {
chr::steady_clock::time_point completionTime;
public:
Statistics(const std::string& name);
void incrementCounter(const hipCounter &counter, const std::string& name);
Statistics(const std::string &name);
void incrementCounter(const hipCounter &counter, const std::string &name);
// Add the counters from `other` onto the counters of this object.
void add(const Statistics &other);
void lineTouched(int lineNumber);
@@ -226,12 +226,12 @@ public:
* processing one file at a time, this allows us to simply expose the stats for the current file globally,
* simplifying things.
*/
static Statistics& current();
static Statistics &current();
/**
* Set the active Statistics object to the named one, creating it if necessary, and write the completion
* timestamp into the currently active one.
*/
static void setActive(const std::string& name);
static void setActive(const std::string &name);
// Check the counter and option TranslateToRoc whether it should be translated to Roc or not.
static bool isToRoc(const hipCounter &counter);
// Check whether the counter is HIP_UNSUPPORTED or not.
projects/hip/hipify-clang/src/StringUtils.cpp
+4 -4
Ver fichero
@@ -33,14 +33,14 @@ llvm::StringRef unquoteStr(llvm::StringRef s) {
return s;
}
void removePrefixIfPresent(std::string &s, const std::string& prefix) {
void removePrefixIfPresent(std::string &s, const std::string &prefix) {
if (s.find(prefix) != 0) {
return;
}
s.erase(0, prefix.size());
}
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
std::string getAbsoluteFilePath(const std::string &sFile, std::error_code &EC) {
if (sFile.empty()) {
return sFile;
}
@@ -59,8 +59,8 @@ std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
return fileAbsPath.c_str();
}
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC,
const std::string& sDirType, bool bCreateDir) {
std::string getAbsoluteDirectoryPath(const std::string &sDir, std::error_code &EC,
const std::string &sDirType, bool bCreateDir) {
if (sDir.empty()) {
return sDir;
}
projects/hip/hipify-clang/src/StringUtils.h
+4 -4
Ver fichero
@@ -33,16 +33,16 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
/**
* If `s` starts with `prefix`, remove it. Otherwise, does nothing.
*/
void removePrefixIfPresent(std::string &s, const std::string& prefix);
void removePrefixIfPresent(std::string &s, const std::string &prefix);
/**
* Returns Absolute File Path based on filename, otherwise - error.
*/
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC);
std::string getAbsoluteFilePath(const std::string &sFile, std::error_code &EC);
/**
* Returns Absolute Directory Path based on directory name, otherwise - error;
* by default the directory is temporary and created.
*/
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC,
const std::string& sDirType = "temporary", bool bCreateDir = true);
std::string getAbsoluteDirectoryPath(const std::string &sDir, std::error_code &EC,
const std::string &sDirType = "temporary", bool bCreateDir = true);
projects/hip/hipify-clang/src/main.cpp
+20 -24
Ver fichero
@@ -194,19 +194,17 @@ int main(int argc, const char **argv) {
ct::RefactoringTool Tool(OptionsParser.getCompilations(), std::string(tmpFile.c_str()));
ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str());
ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str();
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
if (llcompat::canCompileHostAndDeviceInOneJob()) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-compile-host-device", ct::ArgumentInsertPosition::BEGIN));
} else {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
if (!IncludeDirs.empty()) {
for (std::string s : IncludeDirs) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::BEGIN));
}
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN));
if (!CudaPath.empty()) {
std::string sCudaPath = "--cuda-path=" + CudaPath;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN));
if (!MacroNames.empty()) {
for (std::string s : MacroNames) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::BEGIN));
}
}
// Includes for clang's CUDA wrappers for using by packaged hipify-clang
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN));
@@ -219,28 +217,26 @@ int main(int argc, const char **argv) {
stdCpp = "-std=c++14";
#endif
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN));
std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str();
#if defined(HIPIFY_CLANG_RES)
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES, ct::ArgumentInsertPosition::BEGIN));
#endif
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-fno-delayed-template-parsing", ct::ArgumentInsertPosition::BEGIN));
if (llcompat::pragma_once_outside_header()) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN));
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
if (!CudaGpuArch.empty()) {
std::string sCudaGpuArch = "--cuda-gpu-arch=" + CudaGpuArch;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaGpuArch.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
if (!MacroNames.empty()) {
for (std::string s : MacroNames) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::END));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END));
}
}
if (!IncludeDirs.empty()) {
for (std::string s : IncludeDirs) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::END));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END));
}
if (!CudaPath.empty()) {
std::string sCudaPath = "--cuda-path=" + CudaPath;
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sCudaPath.c_str(), ct::ArgumentInsertPosition::BEGIN));
}
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
if (Verbose) {
Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-v", ct::ArgumentInsertPosition::END));
}
projects/hip/include/hip/hcc_detail/hip_fp16.h
+6 -12
Ver fichero
@@ -1126,8 +1126,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data ==
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1135,8 +1134,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data !=
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1144,8 +1142,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data <=
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1153,8 +1150,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data >=
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1162,8 +1158,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data <
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1171,8 +1166,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data >
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+3 -1
Ver fichero
@@ -2994,6 +2994,7 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
* When using this API, start the profiler with profiling disabled. (--startdisabled)
* @warning : hipProfilerStart API is under development.
*/
DEPRECATED("use roctracer/rocTX instead")
hipError_t hipProfilerStart();
@@ -3002,6 +3003,7 @@ hipError_t hipProfilerStart();
* When using this API, start the profiler with profiling disabled. (--startdisabled)
* @warning : hipProfilerStop API is under development.
*/
DEPRECATED("use roctracer/rocTX instead")
hipError_t hipProfilerStop();
@@ -3313,7 +3315,7 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d
hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size_t* offset,
const void* devPtr, const struct hipChannelFormatDesc* desc,
size_t width, size_t height, textureReference* tex);
size_t width, size_t height, textureReference* tex, size_t pitch);
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& tex,
projects/hip/include/hip/hcc_detail/hip_vector_types.h
+10 -39
Ver fichero
@@ -37,11 +37,6 @@ THE SOFTWARE.
#if !defined(_MSC_VER) || __clang__
#if defined(__clang__)
#define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
#elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax.
#define __ROUND_UP_TO_NEXT_POT__(x) \
(1 << (31 - __builtin_clz(x) + (x > (1 << (31 - __builtin_clz(x))))))
#define __NATIVE_VECTOR__(n, T) \
__attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T))))
#endif
#if defined(__cplusplus) && defined(__clang__)
@@ -932,47 +927,23 @@ THE SOFTWARE.
using CUDA_name##4 = HIP_vector_type<T, 4>;
#else
#define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
typedef T CUDA_name##_impl1 __NATIVE_VECTOR__(1, T);\
typedef T CUDA_name##_impl2 __NATIVE_VECTOR__(2, T);\
typedef T CUDA_name##_impl3 __NATIVE_VECTOR__(3, T);\
typedef T CUDA_name##_impl4 __NATIVE_VECTOR__(4, T);\
typedef struct {\
union {\
CUDA_name##_impl1 data;\
struct {\
T x;\
};\
};\
T x;\
} CUDA_name##1;\
typedef struct {\
union {\
CUDA_name##_impl2 data;\
struct {\
T x;\
T y;\
};\
};\
T x;\
T y;\
} CUDA_name##2;\
typedef struct {\
union {\
T data[3];\
struct {\
T x;\
T y;\
T z;\
};\
};\
T x;\
T y;\
T z;\
} CUDA_name##3;\
typedef struct {\
union {\
CUDA_name##_impl4 data;\
struct {\
T x;\
T y;\
T z;\
T w;\
};\
};\
T x;\
T y;\
T z;\
T w;\
} CUDA_name##4;
#endif
projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp
+2 -2
Ver fichero
@@ -26,7 +26,7 @@ THE SOFTWARE.
#include "helpers.hpp"
#include "hc.hpp"
#include "hip/hip_hcc.h"
#include "hip/hip_ext.h"
#include "hip_runtime.h"
#include <functional>
@@ -795,4 +795,4 @@ requires(Domain<K> == {Ts...}) inline std::enable_if_t<
hipLaunchKernelGGL(kernel_name, num_blocks, dim_blocks, group_mem_bytes, stream, \
hipLaunchParm{}, ##__VA_ARGS__); \
} while (0)
} // namespace hip_impl
} // namespace hip_impl
projects/hip/include/hip/hip_ext.h
+118
Ver fichero
@@ -0,0 +1,118 @@
/*
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifndef HIP_INCLUDE_HIP_HIP_EXT_H
#define HIP_INCLUDE_HIP_HIP_EXT_H
#include "hip/hip_runtime_api.h"
#ifdef __HCC__
// Forward declarations:
namespace hc {
class accelerator;
class accelerator_view;
}; // namespace hc
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup HCC-specific features
* @warning These APIs provide access to special features of HCC compiler and are not available
*through the CUDA path.
* @{
*/
/**
* @brief Return hc::accelerator associated with the specified deviceId
* @return #hipSuccess, #hipErrorInvalidDevice
*/
HIP_PUBLIC_API
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
/**
* @brief Return hc::accelerator_view associated with the specified stream
*
* If stream is 0, the accelerator_view for the default stream is returned.
* @return #hipSuccess
*/
HIP_PUBLIC_API
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av);
#endif // #ifdef __HCC__
/**
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed
to kernelparams or extra
*
* @param [in[ f Kernel to launch.
* @param [in] gridDimX X grid dimension specified in work-items
* @param [in] gridDimY Y grid dimension specified in work-items
* @param [in] gridDimZ Z grid dimension specified in work-items
* @param [in] blockDimX X block dimensions specified in work-items
* @param [in] blockDimY Y grid dimension specified in work-items
* @param [in] blockDimZ Z grid dimension specified in work-items
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The
kernel can access this with HIP_DYNAMIC_SHARED.
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th
default stream is used with associated synchronization rules.
* @param [in] kernelParams
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
must be in the memory layout and alignment expected by the kernel.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
the kernel launch. The event must be created before calling this API.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please
refer to hip_porting_driver_api.md for sample usage.
* HIP/ROCm actually updates the start event when the associated kernel completes.
*/
HIP_PUBLIC_API
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent = nullptr,
hipEvent_t stopEvent = nullptr,
uint32_t flags = 0);
HIP_PUBLIC_API
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent = nullptr,
hipEvent_t stopEvent = nullptr)
__attribute__((deprecated("use hipExtModuleLaunchKernel instead")));
// doxygen end AMD-specific features
/**
* @}
*/
#endif // #ifdef HIP_INCLUDE_HIP_HIP_EXT_H
projects/hip/include/hip/hip_hcc.h
+2 -97
Ver fichero
@@ -1,16 +1,13 @@
/*
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
@@ -22,98 +19,6 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_HIP_HCC_H
#define HIP_INCLUDE_HIP_HIP_HCC_H
#ifdef __HCC__
#include "hip/hip_runtime_api.h"
// Forward declarations:
namespace hc {
class accelerator;
class accelerator_view;
}; // namespace hc
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup HCC-specific features
* @warning These APIs provide access to special features of HCC compiler and are not available
*through the CUDA path.
* @{
*/
/**
* @brief Return hc::accelerator associated with the specified deviceId
* @return #hipSuccess, #hipErrorInvalidDevice
*/
HIP_PUBLIC_API
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
/**
* @brief Return hc::accelerator_view associated with the specified stream
*
* If stream is 0, the accelerator_view for the default stream is returned.
* @return #hipSuccess
*/
HIP_PUBLIC_API
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av);
/**
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed
to kernelparams or extra
*
* @param [in[ f Kernel to launch.
* @param [in] gridDimX X grid dimension specified in work-items
* @param [in] gridDimY Y grid dimension specified in work-items
* @param [in] gridDimZ Z grid dimension specified in work-items
* @param [in] blockDimX X block dimensions specified in work-items
* @param [in] blockDimY Y grid dimension specified in work-items
* @param [in] blockDimZ Z grid dimension specified in work-items
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The
kernel can access this with HIP_DYNAMIC_SHARED.
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th
default stream is used with associated synchronization rules.
* @param [in] kernelParams
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
must be in the memory layout and alignment expected by the kernel.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
the kernel launch. The event must be created before calling this API.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please
refer to hip_porting_driver_api.md for sample usage.
* HIP/ROCm actually updates the start event when the associated kernel completes.
*/
HIP_PUBLIC_API
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent = nullptr,
hipEvent_t stopEvent = nullptr,
uint32_t flags = 0);
HIP_PUBLIC_API
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent = nullptr,
hipEvent_t stopEvent = nullptr)
__attribute__((deprecated("use hipExtModuleLaunchKernel instead")));
// doxygen end HCC-specific features
/**
* @}
*/
#endif // #ifdef __HCC__
#warning "hip/hip_hcc.h is deprecated, please use hip/hip_ext.h"
#include "hip/hip_ext.h"
#endif // #ifdef HIP_INCLUDE_HIP_HIP_HCC_H
projects/hip/include/hip/hip_profile.h
+3
Ver fichero
@@ -23,11 +23,14 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_HIP_PROFILE_H
#define HIP_INCLUDE_HIP_HIP_PROFILE_H
#warning "HIP Profiling through markers is deprecated, please check roctrace/rocTX support."
#if not defined(ENABLE_HIP_PROFILE)
#define ENABLE_HIP_PROFILE 1
#endif
#if defined(__HIP_PLATFORM_HCC__) and (ENABLE_HIP_PROFILE == 1)
#warning "HIP Markers are deprecated and would be removed soon."
#include <CXLActivityLogger.h>
#define HIP_SCOPED_MARKER(markerName, group) \
amdtScopedMarker __scopedMarker(markerName, group, nullptr);
projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp
+1 -1
Ver fichero
@@ -27,7 +27,7 @@ THE SOFTWARE.
#include <vector>
#ifdef __HIP_PLATFORM_HCC__
#include <hip/hip_hcc.h>
#include <hip/hip_ext.h>
#endif
#define LEN 64
projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
+1 -1
Ver fichero
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include <iostream>
#include <fstream>
#include <vector>
#include <hip/hip_hcc.h>
#include <hip/hip_ext.h>
#define LEN 64
#define SIZE LEN * sizeof(float)
projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
-2
Ver fichero
@@ -21,11 +21,9 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
//#include "hip/hip_runtime_api.h"
#include <iostream>
#include <fstream>
#include <vector>
//#include <hip/hip_hcc.h>
#define fileName "tex2dKernel.code"
projects/hip/src/hip_event.cpp
+83 -129
Ver fichero
@@ -44,28 +44,22 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future* cf, hipS
}
std::pair<hipEventStatus_t, uint64_t> ihipEvent_t::refreshEventStatus() {
auto ecd = locked_copyCrit();
if (ecd._state == hipEventStatusRecording) {
bool isReady1 = ecd._stream->locked_eventIsReady(this);
if (isReady1) {
LockedAccessor_EventCrit_t eCrit(_criticalData);
if ((eCrit->_eventData._type == hipEventTypeIndependent) ||
(eCrit->_eventData._type == hipEventTypeStopCommand)) {
eCrit->_eventData._timestamp = eCrit->_eventData.marker().get_end_tick();
} else if (eCrit->_eventData._type == hipEventTypeStartCommand) {
eCrit->_eventData._timestamp = eCrit->_eventData.marker().get_begin_tick();
} else {
eCrit->_eventData._timestamp = 0;
assert(0); // TODO - move to debug assert
}
eCrit->_eventData._state = hipEventStatusComplete;
return std::pair<hipEventStatus_t, uint64_t>(eCrit->_eventData._state,
eCrit->_eventData._timestamp);
static std::pair<hipEventStatus_t, uint64_t> refreshEventStatus(ihipEventData_t &ecd) {
if (ecd._state == hipEventStatusRecording && ecd.marker().is_ready()) {
if ((ecd._type == hipEventTypeIndependent) ||
(ecd._type == hipEventTypeStopCommand)) {
ecd._timestamp = ecd.marker().get_end_tick();
} else if (ecd._type == hipEventTypeStartCommand) {
ecd._timestamp = ecd.marker().get_begin_tick();
} else {
ecd._timestamp = 0;
assert(0); // TODO - move to debug assert
}
ecd._state = hipEventStatusComplete;
return std::pair<hipEventStatus_t, uint64_t>(ecd._state,
ecd._timestamp);
}
// Not complete path here:
@@ -109,51 +103,28 @@ hipError_t hipEventCreate(hipEvent_t* event) {
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream);
hipError_t status;
if (event){
auto ecd = event->locked_copyCrit();
if( ecd._state != hipEventStatusUnitialized) {
stream = ihipSyncAndResolveStream(stream);
if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
// TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
//
// If default stream , then wait on all queues.
ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true, true);
{
LockedAccessor_EventCrit_t eCrit(event->criticalData());
eCrit->_eventData.marker(hc::completion_future()); // reset event
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = hc::get_system_ticks();
eCrit->_eventData._state = hipEventStatusComplete;
}
status = hipSuccess;
} else {
// Record the event in the stream:
// Keep a copy outside the critical section so we lock stream first, then event - to
// avoid deadlock
hc::completion_future cf = stream->locked_recordEvent(event);
{
LockedAccessor_EventCrit_t eCrit(event->criticalData());
eCrit->_eventData.marker(cf);
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = 0;
eCrit->_eventData._state = hipEventStatusRecording;
}
status = hipSuccess;
}
} else {
status = hipErrorInvalidResourceHandle;
}
} else {
status = hipErrorInvalidResourceHandle;
if (!event) return ihipLogStatus(hipErrorInvalidResourceHandle);
stream = ihipSyncAndResolveStream(stream);
LockedAccessor_EventCrit_t eCrit(event->criticalData());
if (eCrit->_eventData._state == hipEventStatusUnitialized) return ihipLogStatus(hipErrorInvalidResourceHandle);
if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
// TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
// If default stream , then wait on all queues.
ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true, true);
eCrit->_eventData.marker(hc::completion_future()); // reset event
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = hc::get_system_ticks();
eCrit->_eventData._state = hipEventStatusComplete;
}
return ihipLogStatus(status);
else {
// Record the event in the stream:
eCrit->_eventData.marker(stream->locked_recordEvent(event));
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = 0;
eCrit->_eventData._state = hipEventStatusRecording;
}
return ihipLogStatus(hipSuccess);
}
@@ -203,84 +174,67 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
HIP_INIT_API(hipEventElapsedTime, ms, start, stop);
hipError_t status = hipSuccess;
if (ms == nullptr) return ihipLogStatus(hipErrorInvalidValue);
if ((start == nullptr) || (stop == nullptr)) return ihipLogStatus(hipErrorInvalidResourceHandle);
if (ms == nullptr) {
status = hipErrorInvalidValue;
*ms = 0.0f;
auto startEcd = start->locked_copyCrit();
auto stopEcd = stop->locked_copyCrit();
if ((start->_flags & hipEventDisableTiming) ||
(startEcd._state == hipEventStatusUnitialized) ||
(startEcd._state == hipEventStatusCreated) ||
(stop->_flags & hipEventDisableTiming) ||
(stopEcd._state == hipEventStatusUnitialized) ||
(stopEcd._state == hipEventStatusCreated)) {
// Both events must be at least recorded else return hipErrorInvalidResourceHandle
return ihipLogStatus(hipErrorInvalidResourceHandle);
}
else if ((start == nullptr) || (stop == nullptr)) {
status = hipErrorInvalidResourceHandle;
} else {
*ms = 0.0f;
auto startEcd = start->locked_copyCrit();
auto stopEcd = stop->locked_copyCrit();
if ((start->_flags & hipEventDisableTiming) ||
(startEcd._state == hipEventStatusUnitialized) ||
(startEcd._state == hipEventStatusCreated) || (stop->_flags & hipEventDisableTiming) ||
(stopEcd._state == hipEventStatusUnitialized) ||
(stopEcd._state == hipEventStatusCreated)) {
// Both events must be at least recorded else return hipErrorInvalidResourceHandle
// Refresh status, if still recording...
status = hipErrorInvalidResourceHandle;
auto startStatus = refreshEventStatus(startEcd); // pair < state, timestamp >
auto stopStatus = refreshEventStatus(stopEcd); // pair < state, timestamp >
if ((startStatus.first == hipEventStatusComplete) &&
(stopStatus.first == hipEventStatusComplete)) {
// Common case, we have good information for both events. 'second' is the timestamp:
int64_t tickDiff = (stopStatus.second - startStatus.second);
uint64_t freqHz;
hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
if (freqHz) {
*ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f;
return ihipLogStatus(hipSuccess);
} else {
// Refresh status, if still recording...
auto startStatus = start->refreshEventStatus(); // pair < state, timestamp >
auto stopStatus = stop->refreshEventStatus(); // pair < state, timestamp >
if ((startStatus.first == hipEventStatusComplete) &&
(stopStatus.first == hipEventStatusComplete)) {
// Common case, we have good information for both events. 'second" is the
// timestamp:
int64_t tickDiff = (stopStatus.second - startStatus.second);
uint64_t freqHz;
hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
if (freqHz) {
*ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f;
status = hipSuccess;
} else {
*ms = 0.0f;
status = hipErrorInvalidValue;
}
} else if ((startStatus.first == hipEventStatusRecording) ||
(stopStatus.first == hipEventStatusRecording)) {
status = hipErrorNotReady;
} else {
assert(0);
}
*ms = 0.0f;
return ihipLogStatus(hipErrorInvalidValue);
}
} else if ((startStatus.first == hipEventStatusRecording) ||
(stopStatus.first == hipEventStatusRecording)) {
return ihipLogStatus(hipErrorNotReady);
} else {
assert(0); // TODO should we return hipErrorUnknown ?
}
return ihipLogStatus(status);
return ihipLogStatus(hipSuccess);
}
hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event);
hipError_t status = hipSuccess;
if (!event) return ihipLogStatus(hipErrorInvalidResourceHandle);
if ( NULL == event)
{
status = hipErrorInvalidResourceHandle;
} else {
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
"hipEventQuery on event without system-scope fence ; consider creating with "
"hipEventReleaseToSystem\n");
}
auto ecd = event->locked_copyCrit();
if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
status = hipErrorNotReady;
} else {
status = hipSuccess;
}
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
"hipEventQuery on event without system-scope fence ; consider creating with "
"hipEventReleaseToSystem\n");
}
return ihipLogStatus(status);
auto ecd = event->locked_copyCrit();
if (ecd._state == hipEventStatusRecording && !ecd.marker().is_ready()) {
return ihipLogStatus(hipErrorNotReady);
}
return ihipLogStatus(hipSuccess);
}
projects/hip/src/hip_hcc.cpp
+10 -15
Ver fichero
@@ -46,7 +46,7 @@ THE SOFTWARE.
#include "hsa/hsa_ext_image.h"
#include "hip/hip_runtime.h"
#include "hip_hcc_internal.h"
#include "hip/hip_hcc.h"
#include "hip/hip_ext.h"
#include "trace_helper.h"
#include "env.h"
@@ -325,14 +325,6 @@ void ihipStream_t::locked_streamWaitEvent(ihipEventData_t& ecd) {
}
// Causes current stream to wait for specified event to complete:
// Note this does not provide any kind of host serialization.
bool ihipStream_t::locked_eventIsReady(hipEvent_t event) {
LockedAccessor_EventCrit_t ecrit(event->criticalData());
return (ecrit->_eventData.marker().is_ready());
}
// Create a marker in this stream.
// Save state in the event so it can track the status of the event.
hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event) {
@@ -1505,18 +1497,21 @@ hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) {
return e;
}
void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) {
bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs) {
hipError_t e = hipSuccess;
// Synchronize stream
tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n",
ToString(cb->_stream).c_str());
GET_TLS();
e = ihipStreamSynchronize(tls, cb->_stream);
ihipStreamCallback_t* cb = static_cast<ihipStreamCallback_t*> (cbArgs);
if(cb->comFuture.valid())
cb->comFuture.wait();
// Call registered callback function
cb->_callback(cb->_stream, e, cb->_userData);
hsa_signal_store_screlease(cb->_signal,0);
delete cb;
return false;
}
//---
projects/hip/src/hip_hcc_internal.h
+3 -5
Ver fichero
@@ -583,8 +583,6 @@ class ihipStream_t {
void locked_streamWaitEvent(ihipEventData_t& event);
hc::completion_future locked_recordEvent(hipEvent_t event);
bool locked_eventIsReady(hipEvent_t event);
ihipStreamCritical_t& criticalData() { return _criticalData; };
//---
@@ -652,7 +650,9 @@ class ihipStreamCallback_t {
: _stream(stream), _callback(callback), _userData(userData) {
};
hipStream_t _stream;
hsa_signal_t _signal;
hipStreamCallback_t _callback;
hc::completion_future comFuture;
void* _userData;
};
@@ -724,8 +724,6 @@ class ihipEvent_t {
explicit ihipEvent_t(unsigned flags);
void attachToCompletionFuture(const hc::completion_future* cf, hipStream_t stream,
ihipEventType_t eventType);
std::pair<hipEventStatus_t, uint64_t> refreshEventStatus(); // returns pair <state, timestamp>
// Return a copy of the critical state. The critical data is locked during the copy.
ihipEventData_t locked_copyCrit() {
@@ -972,7 +970,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod,
hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0);
hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream);
void ihipStreamCallbackHandler(ihipStreamCallback_t* cb);
bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs);
// Stream printf functions:
inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) {
projects/hip/src/hip_memory.cpp
+396 -619
Ver fichero
La diferencia del archivo ha sido suprimido porque es demasiado grande Cargar Diff
projects/hip/src/hip_module.cpp
+1 -1
Ver fichero
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include "hip/hcc_detail/hsa_helpers.hpp"
#include "hip/hcc_detail/program_state.hpp"
#include "hip_hcc_internal.h"
#include "hip/hip_hcc.h"
#include "hip/hip_ext.h"
#include "program_state.inl"
#include "trace_helper.h"
projects/hip/src/hip_stream.cpp
+125 -2
Ver fichero
@@ -254,14 +254,137 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
//---
void setCallbackPacket(hsa_queue_t* queue,
uint64_t& index, uint64_t& nextIndex,
hsa_barrier_and_packet_t** barrier1,
hsa_barrier_and_packet_t** barrier2){
uint64_t tempIndex = 0;
uint32_t mask = queue->size - 1;
hsa_barrier_and_packet_t* tempBarrier;
// Check for empty packets
do{
tempIndex = hsa_queue_load_write_index_scacquire(queue);
tempBarrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[tempIndex & mask]);
}while(!(tempBarrier->header & HSA_PACKET_TYPE_INVALID));
// Reserve two packets for two barriers
index = hsa_queue_add_write_index_scacquire(queue, 2);
if(index > mask)
{
index = 0;
nextIndex = 1;
}
else if(index == mask)
nextIndex = 0;
else
nextIndex = index + 1;
tempBarrier = new hsa_barrier_and_packet_t;
memset(tempBarrier, 0, sizeof(hsa_barrier_and_packet_t));
tempBarrier->header = HSA_PACKET_TYPE_INVALID;
// Barrier 1
*barrier1 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index & mask]);
memcpy(*barrier1,tempBarrier,sizeof(hsa_barrier_and_packet_t));
// Barrier 2
*barrier2 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[nextIndex & mask]);
memcpy(*barrier2,tempBarrier,sizeof(hsa_barrier_and_packet_t));
delete tempBarrier;
}
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
unsigned int flags) {
HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
hipError_t e = hipSuccess;
// Create a thread in detached mode to handle callback
if(stream == hipStreamNull)
{
ihipCtx_t* device = ihipGetTlsDefaultCtx();
stream = device->_defaultStream;
}
stream = ihipSyncAndResolveStream(stream);
// Lock the stream
LockedAccessor_StreamCrit_t crit(stream->criticalData());
// Device synchronization
hc::completion_future marker = crit->_av.create_marker(hc::system_scope);
// 1. Lock the queue
hsa_queue_t* lockedQ = static_cast<hsa_queue_t*> (crit->_av.acquire_locked_hsa_queue());
if(lockedQ == nullptr)
{
// No queue attached to stream hence exiting early
return ihipLogStatus(hipErrorMissingConfiguration);
}
// 2. Allocate a singals
hsa_signal_t signal;
hsa_status_t status = hsa_signal_create(1, 0, NULL, &signal);
if(status != HSA_STATUS_SUCCESS)
{
crit->_av.release_locked_hsa_queue();
return ihipLogStatus(hipErrorInvalidValue);
}
hsa_signal_t depSignal;
status = hsa_signal_create(1, 0, NULL, &depSignal);
if(status != HSA_STATUS_SUCCESS)
{
crit->_av.release_locked_hsa_queue();
return ihipLogStatus(hipErrorInvalidValue);
}
// 3. Store callback details, will destroy allocation in callback handler
ihipStreamCallback_t* cb = new ihipStreamCallback_t(stream, callback, userData);
std::thread(ihipStreamCallbackHandler, cb).detach();
if(cb == nullptr)
{
crit->_av.release_locked_hsa_queue();
return ihipLogStatus(hipErrorMemoryAllocation);
}
cb->_signal = depSignal;
cb->comFuture = marker ;
// 4. Create barrier packets
uint64_t index ;
uint64_t nextIndex;
hsa_barrier_and_packet_t* barrier;
hsa_barrier_and_packet_t* depBarrier;
setCallbackPacket(lockedQ, index, nextIndex, &barrier, &depBarrier);
barrier->completion_signal = signal;
depBarrier->dep_signal[0] = depSignal;
uint16_t header = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE)| 1 << HSA_PACKET_HEADER_BARRIER;
// 5. Update packet header,
// Intentionally updated second barrier header before first in order to avoid race
depBarrier->header = header;
barrier->header = header;
// 6. Trigger the doorbell
nextIndex = nextIndex + 1;
hsa_queue_store_write_index_screlease(lockedQ, nextIndex);
hsa_signal_store_relaxed(lockedQ->doorbell_signal, index+1);
// 7. Release queue
crit->_av.release_locked_hsa_queue();
// 8. Register signal callback
hsa_amd_signal_async_handler(signal, HSA_SIGNAL_CONDITION_EQ, 0, ihipStreamCallbackHandler, cb);
return ihipLogStatus(e);
}
projects/hip/src/hip_texture.cpp
+14 -6
Ver fichero
@@ -222,7 +222,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
void* devPtr = nullptr;
size_t pitch = 0;
switch (pResDesc->resType) {
case hipResourceTypeArray:
devPtr = pResDesc->res.array.array->data;
@@ -279,6 +279,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
pitch = pResDesc->res.pitch2D.pitchInBytes;
getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode,
&channelOrder, &channelType);
break;
@@ -296,7 +297,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS !=
hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return ihipLogStatus(hipErrorRuntimeOther);
@@ -456,7 +457,7 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev
hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, size_t* offset,
const void* devPtr, const struct hipChannelFormatDesc* desc,
size_t width, size_t height, textureReference* tex) {
size_t width, size_t height, textureReference* tex, size_t pitch) {
hipError_t hip_status = hipSuccess;
enum hipTextureAddressMode addressMode = tex->addressMode[0];
enum hipTextureFilterMode filterMode = tex->filterMode;
@@ -504,7 +505,7 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS !=
hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return hipErrorRuntimeOther;
@@ -522,8 +523,12 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d
size_t pitch) {
HIP_INIT_API(hipBindTexture2D, offset, tex, devPtr, desc, width, height, pitch);
hipError_t hip_status = hipSuccess;
//TODO: Fix when HSA accepts user defined pitch
if(pitch % 64) pitch =0;
hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, offset, devPtr,
desc, width, height, tex);
desc, width, height, tex, pitch);
return ihipLogStatus(hip_status);
}
@@ -797,7 +802,10 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT
size_t offset;
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
//TODO: Fix when HSA accepts user defined pitch
if(pitch % 64) pitch =0;
hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, &offset, devPtr,
NULL, desc->Width, desc->Height, tex);
NULL, desc->Width, desc->Height, tex, pitch);
return ihipLogStatus(hip_status);
}
projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu
+46
Ver fichero
@@ -0,0 +1,46 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
// ToDo: change lit testing in order to parse the output.
#define LEN 1024
#define SIZE LEN * sizeof(float)
#define ITER 1024*1024
// CHECK: #include <hip/hip_runtime.h>
#include <algorithm>
#define CUDA_LAUNCH(cuda_call,dimGrid,dimBlock, ...) \
cuda_call<<<dimGrid,dimBlock>>>(__VA_ARGS__);
__global__ void Inc1(float *Ad, float *Bd) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if (tx < 1) {
for (int i = 0; i < ITER; ++i) {
Ad[tx] = Ad[tx] + 1.0f;
for (int j = 0; j < 256; ++j) {
Bd[tx] = Ad[tx];
}
}
}
}
int main() {
float *A, *Ad, *Bd;
A = new float[LEN];
for (int i = 0; i < LEN; ++i) {
A[i] = 0.0f;
}
// CHECK: hipError_t status;
cudaError_t status;
// CHECK: status = hipHostRegister(A, SIZE, hipHostRegisterMapped);
status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped);
// CHECK: hipHostGetDevicePointer(&Ad, A, 0);
cudaHostGetDevicePointer(&Ad, A, 0);
// CHECK: hipMalloc((void**)&Bd, SIZE);
cudaMalloc((void**)&Bd, SIZE);
dim3 dimGrid(LEN / 512, 1, 1);
dim3 dimBlock(512, 1, 1);
// CHECK: hipLaunchKernelGGL(Inc1, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd);
CUDA_LAUNCH(Inc1, dimGrid, dimBlock, Ad, Bd);
}
projects/hip/tests/src/hipHcc.cpp
+1 -1
Ver fichero
@@ -30,7 +30,7 @@ THE SOFTWARE.
#include <stdio.h>
#include <iostream>
#include "hip/hip_runtime.h"
#include "hip/hip_hcc.h"
#include "hip/hip_ext.h"
#include "test_common.h"
#define CHECK(error) \
projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+2 -2
Ver fichero
@@ -108,11 +108,11 @@ int main() {
dim3 dimBlock = dim3(1);
dim3 dimGrid = dim3(1);
int numBlocks = 0;
uint workgroups[2] = {32, 64};
uint workgroups[4] = {32, 64, 128, 256};
system_clock::time_point start = system_clock::now();
for (uint i = 0; i < 2; ++i) {
for (uint i = 0; i < 4; ++i) {
dimBlock.x = workgroups[i];
// Calculate the device occupancy to know how many blocks can be run concurrently
projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
+1 -1
Ver fichero
@@ -32,7 +32,7 @@ THE SOFTWARE.
#include <iostream>
#include <fstream>
#include <vector>
#include <hip/hip_hcc.h>
#include <hip/hip_ext.h>
#define LEN 64
#define SIZE LEN * sizeof(float)
projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
-2
Ver fichero
@@ -27,11 +27,9 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
//#include "hip/hip_runtime_api.h"
#include <iostream>
#include <fstream>
#include <vector>
//#include <hip/hip_hcc.h>
#define fileName "tex2d_kernel.code"
projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp
+145
Ver fichero
@@ -0,0 +1,145 @@
#include <stdio.h>
#include <hip/hip_runtime.h>
#include <unistd.h>
#include "test_common.h"
#include <atomic>
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
enum class ExecState
{
EXEC_NOT_STARTED,
EXEC_STARTED,
EXEC_CB_STARTED,
EXEC_CB_FINISHED,
EXEC_FINISHED
};
struct UserData
{
size_t size;
int* ptr;
};
// Global variable to check exection order
std::atomic<ExecState> gData(ExecState::EXEC_NOT_STARTED);
void myCallback(hipStream_t stream, hipError_t status, void* user_data)
{
if(gData.load() != ExecState::EXEC_STARTED)
return; // Error hence return early
gData.store(ExecState::EXEC_CB_STARTED);
UserData* data = reinterpret_cast<UserData*>(user_data);
printf("Callback started\n");
sleep(1);
printf("Callback ending.\n");
gData.store(ExecState::EXEC_CB_FINISHED);
}
bool test(int count)
{
printf("\n============ Test iteration %d =============\n",count);
// Stream
hipStream_t stream;
bool result = true;
gData.store(ExecState::EXEC_STARTED);
HIPCHECK(hipStreamCreate(&stream));
// Array size
size_t size = 10000;
// Device array
int *data = NULL;
HIPCHECK(hipMalloc((void**)&data, sizeof(int) * size));
// Initialize device array to -1
HIPCHECK(hipMemset(data, -1, sizeof(int) * size));
// Host array
int *host = NULL;
HIPCHECK(hipHostMalloc((void**)&host, sizeof(int) * size));
// Print host ptr address
printf("In main thread\n");
// Initialize user_data for callback
UserData arg;
arg.size = size;
arg.ptr = host;
// Synchronize device
HIPCHECK(hipDeviceSynchronize());
// Asynchronous copy from device to host
HIPCHECK(hipMemcpyAsync(host, data, sizeof(int) * size, hipMemcpyDeviceToHost, stream));
// Asynchronous memset on device
HIPCHECK(hipMemsetAsync(data, 0, sizeof(int) * size, stream));
// Add callback - should happen after hipMemsetAsync()
HIPCHECK(hipStreamAddCallback(stream, myCallback, &arg, 0));
printf("Will wait in main thread until callback completes\n");
//This should synchronize the stream (including the callback)
HIPCHECK(hipStreamSynchronize(stream));
if(gData.load() != ExecState::EXEC_CB_FINISHED)
{
std::cout<<"Callback is not finished\n";
return false;
}
printf("Callback completed will resume main thread execution\n");
if(host[size/2] != -1)
{
// Print some host data that just got copied
printf("Pseudo host data printing (should be -1): %d\n", host[size/2]);
result = false;
}
HIPCHECK(hipMemcpy(host, data, sizeof(int)*size, hipMemcpyDeviceToHost));
if(host[size-1] != 0)
{
printf("Pseudo host data printing (should be 0): %d\n", host[size-1]);
result = false;
}
HIPCHECK(hipFree(data));
HIPCHECK(hipHostFree(host));
HIPCHECK(hipStreamDestroy(stream));
gData.store(ExecState::EXEC_FINISHED);
return result;
}
int main()
{
// Test involves multithreading hence running multiple times
// to make sure consitency in the behavior
bool status = true;
for(int i=0; i < 10; i++){
status = test(i+1);
if(status == false)
{
failed("Test Failed!\n");
break;
}
}
if(status == true) passed();
return 0;
}
projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp
+133
Ver fichero
@@ -0,0 +1,133 @@
/*
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#include <mutex>
#include <condition_variable>
#include <stdlib.h>
//Globals
const int workloadCount = 1000000;
std::mutex gMutx;
bool callbackCompleted = false;
std::condition_variable condVar;
// Device function
__global__ void increment(int *data,int N)
{
int i = blockIdx.x*blockDim.x + threadIdx.x;
if(i < N)
data[i] = 1 + data[i];
}
struct USER_DATA
{
int* result; // Data received from device
int* copyOfOriginalData; // Copy of initial data which will be used for validation
};
// Callback
void callback(hipStream_t event, hipError_t status, void *userData)
{
USER_DATA *data = (USER_DATA *)userData;
if(!(data == nullptr || data->result == nullptr || data->copyOfOriginalData == nullptr))
{
for(int i=0;i<workloadCount;i++)
{
if(data->result[i] != data->copyOfOriginalData[i]+1)
{
std::cout<<"Error value : "<<data->result[i]<<"| Expected value :"<<data->copyOfOriginalData[i]+1<<std::endl;
break;
}
}
}
callbackCompleted = true;
condVar.notify_all();
}
int main()
{
int *hData = nullptr;
int *dData = nullptr;
int *hResultData = nullptr;
int devCount = 0;
int size = workloadCount * sizeof(int);
// query device count
HIPCHECK(hipGetDeviceCount(&devCount));
// Allocate
// Host allocation
hData = (int *)malloc(size);
hResultData = (int *)malloc(size);
if(hData == nullptr || hResultData == nullptr)
{
HIPCHECK(hipErrorInvalidValue);
}
// Initialize host data
for(int i =0; i<workloadCount; i++)
{
hData[i] = rand()%workloadCount;
}
// Device allocation
HIPCHECK(hipMalloc(&dData,size));
HIPCHECK(hipMemcpyAsync(dData,hData,size,hipMemcpyHostToDevice,0));
dim3 block(256);
dim3 grid((workloadCount + block.x-1) / block.x);
hipLaunchKernelGGL(increment, grid, block, 0, 0, dData, workloadCount);
HIPCHECK(hipMemcpyAsync(hResultData, dData, size, hipMemcpyDeviceToHost, 0));
hipStream_t stream;
USER_DATA *inputParam = (USER_DATA*)malloc(sizeof(USER_DATA));
if(inputParam == nullptr) return 0;
inputParam->result = hResultData;
inputParam->copyOfOriginalData = hData;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipStreamAddCallback(stream,callback,inputParam,0));
// Wait for stream add callback to complete
std::unique_lock<std::mutex> l(gMutx);
while(!callbackCompleted)
condVar.wait(l);
// Will destroy device memory hence explicite hipFree is not needed
HIPCHECK(hipDeviceReset());
free(hData);
free(hResultData);
passed();
}
projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp
+178
Ver fichero
@@ -0,0 +1,178 @@
/*
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
#include <thread>
#include <unistd.h>
#include <mutex>
#include <condition_variable>
#include "test_common.h"
// Will indicate completion of callback added as part of hipStreamAddCallback
struct signal
{
int completedThreads;
std::mutex mu;
std::condition_variable cv;
};
struct workload
{
int _workloadId;
int _deviceID;
int *copyOf_hData; // copy of host data which will be used to validation
int *hData; // will contain host data
int *dData; // device data will be stored
hipStream_t _stream;// stream on which data will be processed
bool success; // start will be stored
};
// Global data
int numWorkloads = 8;
const int perWorkloadSize = 1000000;
signal completionSignal;
// Kernel run on device
__global__ void increment(int *data, int N)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N)
data[i] = data[i]+1;
}
/*
* Method validates processed data array with saved copy and notifies via conditional variable to all waiting threads
*/
void Analyze(hipStream_t event, hipError_t status, void *data)
{
HIPCHECK(status);
workload *W = (workload *) data;
if(W != NULL)
{
W->success = true;
for (int i=0; i< perWorkloadSize; ++i)
{
W->success &= (W->copyOf_hData[i] == (W->hData[i]+1));
if(!W->success)
{
std::cout<<"\nExpected Data :"<<(W->hData[i]+1)<<" Current Data :"<<W->copyOf_hData[i]<<std::endl;
break;
}
}
}
std::lock_guard<std::mutex> guard(completionSignal.mu);
completionSignal.completedThreads += 1;
completionSignal.cv.notify_all();
}
/*
* Thread routine to launch workloads into separate stream
*/
void LaunchWorkload(void *args)
{
workload *W = (workload *) args;
if (W == nullptr) return;
std::srand(std::time(nullptr));
HIPCHECK(hipSetDevice(W->_deviceID));
// Allocate memory
HIPCHECK(hipMalloc(&W->dData,perWorkloadSize*sizeof(int)));
size_t s = perWorkloadSize*sizeof(int);
HIPCHECK(hipMemset(W->dData,0,s))
W->hData = (int *) malloc(perWorkloadSize*sizeof(int));
W->copyOf_hData = (int *) malloc(perWorkloadSize*sizeof(int));
// Initialize host array
for(int i =0;i<perWorkloadSize;i++)
{
W->hData[i] = W->copyOf_hData[i] = std::rand() % perWorkloadSize;
}
HIPCHECK(hipStreamCreate(&W->_stream));
dim3 block(256);
dim3 grid((perWorkloadSize + block.x-1) / block.x);
HIPCHECK(hipMemcpyAsync(W->dData,W->hData,perWorkloadSize*sizeof(int),hipMemcpyHostToDevice,W->_stream));
hipLaunchKernelGGL(increment, grid, block,0, W->_stream,W->dData, perWorkloadSize);
HIPCHECK(hipMemcpyAsync(W->copyOf_hData,W->dData,perWorkloadSize*sizeof(int),hipMemcpyDeviceToHost,W->_stream));
HIPCHECK(hipStreamAddCallback(W->_stream, Analyze, W,0))
}
int main(int argc, char* argv[]) {
int numDevice = 0;
HIPCHECK(hipGetDeviceCount(&numDevice));
std::thread workerThread[numWorkloads];
workload *workloads;
workloads = (workload *) malloc(numWorkloads * sizeof(workload));
for(int i =0;i<numWorkloads; i++)
{
workloads[i]._workloadId = i;
workloads[i]._deviceID = i%numDevice;
// launch threads
workerThread[i] = std::thread(LaunchWorkload,&workloads[i]);
}
// should wait for joining all threads
for(int i =0; i<numWorkloads; i++)
{
workerThread[i].join();
}
// wait util all callbacks are done
std::unique_lock<std::mutex> l(completionSignal.mu);
while(completionSignal.completedThreads != numWorkloads)
{
completionSignal.cv.wait(l);
}
//clean-up
hipDeviceReset();
free(workloads);
passed();
return 0;
}
projects/hip/tests/src/test_common.h
+30
Ver fichero
@@ -153,6 +153,36 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg);
unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N);
template<typename T> // pointer type
void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth)
{
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
int offset = i*width*height + j*width + k;
if (hData[offset] != hOutputData[offset]) {
std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" << hOutputData[offset]<<" ";
failed("mistmatch at:%d %d %d",i,j,k);
}
}
}
}
}
template<typename T>
void checkArray(T input, T output, size_t height, size_t width)
{
for(int i=0; i<height; i++ ){
for(int j=0; j<width; j++ ){
int offset = i*width + j;
if( input[offset] != output[offset] ){
std::cerr << '[' << i << ',' << j << ',' << "]:" << input[offset] << "----" << output[offset]<<" ";
failed("mistmatch at:%d %d",i,j);
}
}
}
}
template <typename T>
__global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) {
projects/hip/tests/src/texture/hipBindTex2DPitch.cpp
+87
Ver fichero
@@ -0,0 +1,87 @@
/*
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*HIT_START
* BUILD: %t %s ../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE_H 8
#define SIZE_W 12
#define TYPE_t float
texture<TYPE_t, 2, hipReadModeElementType> tex;
// texture object is a kernel argument
__global__ void texture2dCopyKernel( TYPE_t* dst) {
int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y;
if ( (x< SIZE_W) && (y< SIZE_H) ){
dst[SIZE_W*y+x] = tex2D(tex, x, y);
}
}
int main (void)
{
TYPE_t* B;
TYPE_t* A;
TYPE_t* devPtrB;
TYPE_t* devPtrA;
B = new TYPE_t[SIZE_H*SIZE_W];
A = new TYPE_t[SIZE_H*SIZE_W];
for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){
A[i-1] = i;
}
size_t devPitchA, tex_ofs;
HIPCHECK(hipMallocPitch((void**)&devPtrA, &devPitchA ,SIZE_W*sizeof(TYPE_t), SIZE_H)) ;
HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t),
SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice));
tex.normalized = false;
#if defined(__HIP_PLATFORM_NVCC__)
cudaError_t status = cudaBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc,
SIZE_W, SIZE_H, devPitchA);
if (status != cudaSuccess) {
printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, cudaGetErrorString(status),
status,__FILE__, __LINE__, KNRM);
failed("API returned error code.");
}
#else
HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc,
SIZE_W, SIZE_H, devPitchA));
#endif
HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;
hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB);
hipDeviceSynchronize();
HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t),
SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost));
HipTest::checkArray(A, B, SIZE_H, SIZE_W);
delete []A;
delete []B;
hipFree(devPtrA);
hipFree(devPtrB);
passed();
}
projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp
+12 -21
Ver fichero
@@ -26,15 +26,6 @@ THE SOFTWARE.
* HIT_END
*/
#define CHECK(cmd) \
{\
hipError_t error = cmd;\
if(error != hipSuccess) {\
fprintf(stderr, "error: '%s' (%d) at %s:%d\n", hipGetErrorString(error), error, __FILE__, __LINE__);\
exit(EXIT_FAILURE);\
}\
}
#include "hip/hip_runtime.h"
#include "test_common.h"
#define SIZE 10
@@ -54,31 +45,31 @@ bool textureTest(enum hipArray_Format texFormat)
{
T hData[] = {65, 66, 67, 68, 69, 70, 71, 72,73,74};
T *dData = NULL;
CHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE));
CHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE));
HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE));
HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE));
textureReference* texRef = &textureNormalizedVal_1D;
CHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp));
CHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp));
CHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint));
CHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES));
CHECK(hipTexRefSetFormat(texRef, texFormat, 1));
HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp));
HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp));
HIPCHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint));
HIPCHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES));
HIPCHECK(hipTexRefSetFormat(texRef, texFormat, 1));
HIP_ARRAY_DESCRIPTOR desc;
desc.Width = SIZE;
desc.Height = 1;
desc.Format = texFormat;
desc.NumChannels = 1;
CHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE));
HIPCHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE));
bool testResult = true;
float *dOutputData = NULL;
CHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE));
HIPCHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE));
hipLaunchKernelGGL(HIP_KERNEL_NAME(normalizedValTextureTest), dim3(1,1,1), dim3(SIZE,1,1), 0, 0, SIZE, dOutputData);
float *hOutputData = new float[SIZE];
CHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE)));
HIPCHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE)));
for(int i = 0; i < SIZE; i++)
{
@@ -100,9 +91,9 @@ int main(int argc, char** argv)
{
int device = 0;
bool status = true;
CHECK(hipSetDevice(device));
HIPCHECK(hipSetDevice(device));
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, device));
HIPCHECK(hipGetDeviceProperties(&props, device));
std::cout << "Device :: " << props.name << std::endl;
#ifdef __HIP_PLATFORM_HCC__
std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl;
projects/hip/tests/src/texture/hipTexObjPitch.cpp
+106
Ver fichero
@@ -0,0 +1,106 @@
/*
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*HIT_START
* BUILD: %t %s ../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
#define SIZE_H 20
#define SIZE_W 179
// texture object is a kernel argument
template <typename TYPE_t>
__global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) {
for(int i =0;i<SIZE_H;i++)
for(int j = 0;j<SIZE_W;j++)
dst[SIZE_W*i+j] = tex2D<TYPE_t>(texObj, j, i);
__syncthreads();
}
template <typename TYPE_t>
void texture2Dtest()
{
TYPE_t* B;
TYPE_t* A;
TYPE_t* devPtrB;
TYPE_t* devPtrA;
B = new TYPE_t[SIZE_H*SIZE_W];
A = new TYPE_t[SIZE_H*SIZE_W];
for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){
A[i-1] = i;
}
size_t devPitchA;
HIPCHECK(hipMallocPitch((void**)&devPtrA, &devPitchA ,SIZE_W*sizeof(TYPE_t), SIZE_H)) ;
HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t),
SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice));
// Use the texture object
hipResourceDesc texRes;
hipMemset(&texRes, 0, sizeof(texRes));
texRes.resType = hipResourceTypePitch2D;
texRes.res.pitch2D.devPtr = devPtrA;
texRes.res.pitch2D.height = SIZE_H;
texRes.res.pitch2D.width = SIZE_W;
texRes.res.pitch2D.pitchInBytes = devPitchA;
texRes.res.pitch2D.desc = hipCreateChannelDesc<TYPE_t>();
hipTextureDesc texDescr;
hipMemset(&texDescr, 0, sizeof(texDescr));
texDescr.normalizedCoords = false;
texDescr.filterMode = hipFilterModePoint;
texDescr.mipmapFilterMode = hipFilterModePoint;
texDescr.addressMode[0] = hipAddressModeClamp;
texDescr.addressMode[1] = hipAddressModeClamp;
texDescr.addressMode[2] = hipAddressModeClamp;
texDescr.readMode = hipReadModeElementType;
hipTextureObject_t texObj;
hipResourceViewDesc resDesc;
HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, &resDesc));
HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ;
hipLaunchKernelGGL(texture2dCopyKernel, dim3(1,1,1), dim3(1,1,1), 0, 0,
texObj, devPtrB, devPtrA);
HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t),
SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost));
HipTest::checkArray(A, B, SIZE_H, SIZE_W);
delete []A;
delete []B;
hipFree(devPtrA);
hipFree(devPtrB);
}
int main()
{
texture2Dtest<float>();
texture2Dtest<int>();
texture2Dtest<unsigned char>();
texture2Dtest<short>();
texture2Dtest<char>();
texture2Dtest<unsigned int>();
passed();
}
projects/hip/tests/src/texture/simpleTexture2DLayered.cpp
+106
Ver fichero
@@ -0,0 +1,106 @@
/*
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
typedef float T;
// Texture reference for 2D Layered texture
texture<float, hipTextureType2DLayered> tex2DL;
__global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer)
{
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
outputData[layer*width*height + y*width + x] = tex2DLayered(tex2DL, x, y, layer);
}
////////////////////////////////////////////////////////////////////////////////
void runTest(int width,int height,int num_layers,texture<T, hipTextureType2DLayered> *tex)
{
unsigned int size = width * height * num_layers * sizeof(T);
T* hData = (T*) malloc(size);
memset(hData, 0, size);
for (unsigned int layer = 0; layer < num_layers; layer++){
for (int i = 0; i < (int)(width * height); i++){
hData[layer*width*height + i] = i;
}
}
hipChannelFormatDesc channelDesc;
// Allocate array and copy image data
channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindFloat);
hipArray *arr;
HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, num_layers), hipArrayLayered));
hipMemcpy3DParms myparms = {0};
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = arr;
myparms.extent = make_hipExtent(width, height, num_layers);
//myparms.kind = hipMemcpyHostToDevice;
HIPCHECK(hipMemcpy3D(&myparms));
// set texture parameters
tex->addressMode[0] = hipAddressModeWrap;
tex->addressMode[1] = hipAddressModeWrap;
tex->filterMode = hipFilterModePoint;
tex->normalized = false;
// Bind the array to the texture
HIPCHECK(hipBindTextureToArray(*tex, arr, channelDesc));
// Allocate device memory for result
T* dData = NULL;
hipMalloc((void **) &dData, size);
dim3 dimBlock(8, 8, 1);
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
for (unsigned int layer = 0; layer < num_layers; layer++)
hipLaunchKernelGGL(simpleKernelLayeredArray, dimGrid, dimBlock, 0, 0, dData, width, height, layer);
HIPCHECK(hipDeviceSynchronize());
// Allocate mem for the result on host side
T *hOutputData = (T*) malloc(size);
memset(hOutputData, 0, size);
// copy result from device to host
HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
HipTest::checkArray(hData,hOutputData,width,height,num_layers);
hipFree(dData);
hipFreeArray(arr);
free(hData);
free(hOutputData);
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
runTest(512,512,5,&tex2DL);
passed();
}
projects/hip/tests/src/texture/simpleTexture3D.cpp
+133
Ver fichero
@@ -0,0 +1,133 @@
/*
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t
* HIT_END
*/
#include "test_common.h"
//typedef char T;
const char *sampleName = "simpleTexture3D";
// Texture reference for 3D texture
texture<float, hipTextureType3D, hipReadModeElementType> texf;
texture<int, hipTextureType3D, hipReadModeElementType> texi;
texture<char, hipTextureType3D, hipReadModeElementType> texc;
template <typename T>
__global__ void simpleKernel3DArray(T* outputData,
int width,
int height,int depth)
{
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
if(std::is_same<T, float>::value)
outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i);
else if(std::is_same<T, int>::value)
outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i);
else if(std::is_same<T, char>::value)
outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i);
}
}
}
}
////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for tex3D
////////////////////////////////////////////////////////////////////////////////
template <typename T>
void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipReadModeElementType> *tex)
{
unsigned int size = width * height * depth * sizeof(T);
T* hData = (T*) malloc(size);
memset(hData, 0, size);
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
hData[i*width*height + j*width +k] = i*width*height + j*width + k;
}
}
}
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindSigned);
hipArray *arr;
HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayCubemap));
hipMemcpy3DParms myparms = {0};
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = arr;
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyHostToDevice;
HIPCHECK(hipMemcpy3D(&myparms));
// set texture parameters
tex->addressMode[0] = hipAddressModeWrap;
tex->addressMode[1] = hipAddressModeWrap;
tex->filterMode = hipFilterModePoint;
tex->normalized = false;
// Bind the array to the texture
HIPCHECK(hipBindTextureToArray(*tex, arr, channelDesc));
// Allocate device memory for result
T* dData = NULL;
hipMalloc((void **) &dData, size);
hipLaunchKernelGGL(simpleKernel3DArray, dim3(1,1,1), dim3(1,1,1), 0, 0, dData, width, height, depth);
HIPCHECK(hipDeviceSynchronize());
// Allocate mem for the result on host side
T *hOutputData = (T*) malloc(size);
memset(hOutputData, 0, size);
// copy result from device to host
HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
HipTest::checkArray(hData,hOutputData,width,height,depth);
hipFree(dData);
hipFreeArray(arr);
free(hData);
free(hOutputData);
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
printf("%s starting...\n", sampleName);
for(int i=1;i<25;i++)
{
runTest<float>(i,i,i,&texf);
runTest<int>(i+1,i,i,&texi);
runTest<char>(i,i+1,i,&texc);
}
passed();
}