diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 7aeb415bcc..b26ee22d62 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -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)
diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md
index da5abc19da..e8c3d0807b 100644
--- a/projects/hip/hipify-clang/README.md
+++ b/projects/hip/hipify-clang/README.md
@@ -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
```
### 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
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp
index cd5c952145..d74ba70f21 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp
@@ -117,7 +117,7 @@ namespace perl {
{sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam}
};
- void generateHeader(unique_ptr& streamPtr) {
+ void generateHeader(unique_ptr &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& streamPtr) {
+ void generateStatFunctions(unique_ptr &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& streamPtr) {
+ void generateSimpleSubstitutions(unique_ptr &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& streamPtr) {
+ void generateExternShared(unique_ptr &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& streamPtr) {
+ void generateKernelLaunch(unique_ptr &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& streamPtr) {
+ void generateCubNamespace(unique_ptr &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& streamPtr) {
+ void generateHostFunctions(unique_ptr &streamPtr) {
*streamPtr.get() << endl << sub << "transformHostFunctions" << " {" << endl_tab << my_k << endl;
set &funcSet = DeviceSymbolFunctions0;
const string s0 = "$k += s/(?second.hipName.str() << "\"";
@@ -291,12 +291,12 @@ namespace perl {
*streamPtr.get() << tab << return_k << "}" << endl;
}
- void generateDeviceFunctions(unique_ptr& streamPtr) {
+ void generateDeviceFunctions(unique_ptr &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++;
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp
index 5ffd84e0b0..fec138915d 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp
@@ -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)) {
diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp
index 510d91978a..dee9a25d49 100644
--- a/projects/hip/hipify-clang/src/HipifyAction.cpp
+++ b/projects/hip/hipify-clang/src/HipifyAction.cpp
@@ -150,7 +150,7 @@ void HipifyAction::RewriteToken(const clang::Token &t) {
void HipifyAction::FindAndReplace(StringRef name,
clang::SourceLocation sl,
- const std::map &repMap,
+ const std::map &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 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) {
diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp
index daebd214d1..604841bcd9 100644
--- a/projects/hip/hipify-clang/src/LLVMCompat.cpp
+++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp
@@ -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
diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h
index c1d671f513..48e008d40d 100644
--- a/projects/hip/hipify-clang/src/LLVMCompat.h
+++ b/projects/hip/hipify-clang/src/LLVMCompat.h
@@ -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 &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
diff --git a/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h b/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h
index 77d23ff6b2..92d77655af 100644
--- a/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h
+++ b/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h
@@ -36,15 +36,15 @@ namespace ct = clang::tooling;
*/
template
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
diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp
index 8dff998dfa..686b51da41 100644
--- a/projects/hip/hipify-clang/src/Statistics.cpp
+++ b/projects/hip/hipify-clang/src/Statistics.cpp
@@ -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);
}
diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h
index 9b9889d0e5..030ea6b2d6 100644
--- a/projects/hip/hipify-clang/src/Statistics.h
+++ b/projects/hip/hipify-clang/src/Statistics.h
@@ -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 ¤t();
/**
* 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.
diff --git a/projects/hip/hipify-clang/src/StringUtils.cpp b/projects/hip/hipify-clang/src/StringUtils.cpp
index 78d0f8113e..31ab331174 100644
--- a/projects/hip/hipify-clang/src/StringUtils.cpp
+++ b/projects/hip/hipify-clang/src/StringUtils.cpp
@@ -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;
}
diff --git a/projects/hip/hipify-clang/src/StringUtils.h b/projects/hip/hipify-clang/src/StringUtils.h
index 39ce726d1b..ecbca5e832 100644
--- a/projects/hip/hipify-clang/src/StringUtils.h
+++ b/projects/hip/hipify-clang/src/StringUtils.h
@@ -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);
diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp
index 3c7424fbca..b4200bd109 100644
--- a/projects/hip/hipify-clang/src/main.cpp
+++ b/projects/hip/hipify-clang/src/main.cpp
@@ -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 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));
}
diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h
index 18ad566121..52abc1a004 100644
--- a/projects/hip/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h
@@ -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__
diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
index 5a61219f4a..ba7f2f4ade 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
@@ -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
hipError_t hipBindTexture2D(size_t* offset, struct texture& tex,
diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h
index 75c79d422c..3a10a06372 100644
--- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h
+++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h
@@ -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;
#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
diff --git a/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp
index ca52e5b614..96d449b213 100644
--- a/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp
+++ b/projects/hip/include/hip/hcc_detail/macro_based_grid_launch.hpp
@@ -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
@@ -795,4 +795,4 @@ requires(Domain == {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
\ No newline at end of file
+} // namespace hip_impl
diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h
new file mode 100644
index 0000000000..3df911ee7a
--- /dev/null
+++ b/projects/hip/include/hip/hip_ext.h
@@ -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
diff --git a/projects/hip/include/hip/hip_hcc.h b/projects/hip/include/hip/hip_hcc.h
index a82a8c622b..e7e27fc721 100644
--- a/projects/hip/include/hip/hip_hcc.h
+++ b/projects/hip/include/hip/hip_hcc.h
@@ -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
diff --git a/projects/hip/include/hip/hip_profile.h b/projects/hip/include/hip/hip_profile.h
index 7474839258..95224af4a3 100644
--- a/projects/hip/include/hip/hip_profile.h
+++ b/projects/hip/include/hip/hip_profile.h
@@ -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
#define HIP_SCOPED_MARKER(markerName, group) \
amdtScopedMarker __scopedMarker(markerName, group, nullptr);
diff --git a/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp
index 4ff000d1af..38cf0d414c 100644
--- a/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp
+++ b/projects/hip/samples/0_Intro/module_api/launchKernelHcc.cpp
@@ -27,7 +27,7 @@ THE SOFTWARE.
#include
#ifdef __HIP_PLATFORM_HCC__
-#include
+#include
#endif
#define LEN 64
diff --git a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
index 7aa7f7ba2d..3a2804b7a2 100644
--- a/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
+++ b/projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
@@ -25,7 +25,7 @@ THE SOFTWARE.
#include
#include
#include
-#include
+#include
#define LEN 64
#define SIZE LEN * sizeof(float)
diff --git a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
index 7c22e19d5d..b42ac86ad1 100755
--- a/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+++ b/projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
@@ -21,11 +21,9 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
-//#include "hip/hip_runtime_api.h"
#include
#include
#include
-//#include
#define fileName "tex2dKernel.code"
diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp
index 73e296dae3..dcab059774 100644
--- a/projects/hip/src/hip_event.cpp
+++ b/projects/hip/src/hip_event.cpp
@@ -44,28 +44,22 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future* cf, hipS
}
-std::pair 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(eCrit->_eventData._state,
- eCrit->_eventData._timestamp);
+static std::pair 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(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);
}
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 7150140e12..377ce0aad4 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -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 (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;
}
//---
diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h
index 658cfbf576..6337edd370 100644
--- a/projects/hip/src/hip_hcc_internal.h
+++ b/projects/hip/src/hip_hcc_internal.h
@@ -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 refreshEventStatus(); // returns pair
-
// 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) {
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index 009c2db505..266f9b51d6 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -41,22 +41,23 @@ hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKi
// Return success if number of bytes to copy is 0
if (sizeBytes == 0) return e;
+ if (!dst || !src) return hipErrorInvalidValue;
- stream = ihipSyncAndResolveStream(stream);
-
- if ((dst == NULL) || (src == NULL)) {
- e = hipErrorInvalidValue;
- } else if (stream) {
- try {
- stream->locked_copyAsync(dst, src, sizeBytes, kind);
- } catch (ihipException& ex) {
- e = ex._code;
- }
- } else {
- e = hipErrorInvalidValue;
+ if (!(stream = ihipSyncAndResolveStream(stream))) {
+ return hipErrorInvalidValue;
}
- return e;
+ try {
+ stream->locked_copyAsync(dst, src, sizeBytes, kind);
+ }
+ catch (ihipException& ex) {
+ e = ex._code;
+ }
+ catch (...) {
+ return hipErrorUnknown;
+ }
+
+ return hipSuccess;
}
// return 0 on success or -1 on error:
@@ -429,58 +430,62 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
return hipHostMalloc(ptr, sizeBytes, flags);
};
+hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width, int height, int depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType,void ** ptr, hsa_ext_image_data_info_t &imageInfo, int array_size __dparm(0)) {
+ auto ctx = ihipGetTlsDefaultCtx();
+ if (ctx) {
+ hc::accelerator acc = ctx->getDevice()->_acc;
+ hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
+ if (!agent)
+ return hipErrorInvalidResourceHandle;
+ size_t allocGranularity = 0;
+ hsa_amd_memory_pool_t* allocRegion = static_cast(acc.get_hsa_am_region());
+ hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
+
+ hsa_ext_image_descriptor_t imageDescriptor;
+ imageDescriptor.geometry = geometry;
+ imageDescriptor.width = width;
+ imageDescriptor.height = height;
+ imageDescriptor.depth = depth;
+ imageDescriptor.array_size = array_size;
+ imageDescriptor.format.channel_order = channelOrder;
+ imageDescriptor.format.channel_type = channelType;
+
+ hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+ hsa_status_t status =
+ hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
+ if(imageInfo.size == 0 || HSA_STATUS_SUCCESS != status){
+ return hipErrorRuntimeOther;
+ }
+ size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
+ const unsigned am_flags = 0;
+ *ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx,
+ false /*shareWithAll*/, am_flags, 0, alignment);
+ if (*ptr == NULL) {
+ return hipErrorMemoryAllocation;
+ }
+ return hipSuccess;
+ }
+ else {
+ return hipErrorMemoryAllocation;
+ }
+}
+
// width in bytes
hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) {
hipError_t hip_status = hipSuccess;
- if(ptr==NULL || pitch == NULL)
- {
- hip_status=hipErrorInvalidValue;
- return hip_status;
- }
- // hardcoded 128 bytes
- *pitch = ((((int)width - 1) / 128) + 1) * 128;
- const size_t sizeBytes = (*pitch) * height * ((depth==0) ? 1 : depth);
-
- auto ctx = ihipGetTlsDefaultCtx();
-
- if (ctx) {
- hc::accelerator acc = ctx->getDevice()->_acc;
- hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
-
- size_t allocGranularity = 0;
- hsa_amd_memory_pool_t* allocRegion =
- static_cast(acc.get_hsa_am_region());
- hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
- &allocGranularity);
-
- hsa_ext_image_descriptor_t imageDescriptor;
- imageDescriptor.width = *pitch;
- imageDescriptor.height = height;
- imageDescriptor.depth = depth;
- imageDescriptor.array_size = 0;
- if (depth == 0)
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
- else
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D;
- imageDescriptor.format.channel_order = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
- imageDescriptor.format.channel_type = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
-
- hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
- hsa_ext_image_data_info_t imageInfo;
- hsa_status_t status =
- hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
- size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
-
- const unsigned am_flags = 0;
- *ptr = hip_internal::allocAndSharePtr("device_pitch", sizeBytes, ctx,
- false /*shareWithAll*/, am_flags, 0, alignment);
-
- if (sizeBytes && (*ptr == NULL)) {
- hip_status = hipErrorMemoryAllocation;
- }
- } else {
- hip_status = hipErrorMemoryAllocation;
+ if(ptr==NULL || pitch == NULL){
+ return hipErrorInvalidValue;
}
+ hsa_ext_image_data_info_t imageInfo;
+ if (depth == 0)
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,width,height,0,HSA_EXT_IMAGE_CHANNEL_ORDER_R,
+ HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo);
+ else
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,width,height,depth,HSA_EXT_IMAGE_CHANNEL_ORDER_R,
+ HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo);
+
+ if(hip_status == hipSuccess)
+ *pitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
return hip_status;
}
@@ -500,9 +505,9 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height, unsigned int elementSizeBytes){
HIP_INIT_SPECIAL_API(hipMemAllocPitch, (TRACE_MEM), dptr, pitch, widthInBytes, height,elementSizeBytes);
HIP_SET_DEVICE();
-
+
if (widthInBytes == 0 || height == 0) return ihipLogStatus(hipErrorInvalidValue);
-
+
return ihipLogStatus(ihipMallocPitch(tls, dptr, pitch, widthInBytes, height, 0));
}
@@ -541,18 +546,74 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType);
+hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height, int depth, hipChannelFormatDesc desc, hsa_ext_image_data_info_t &imageInfo,int array_size __dparm(0))
+{
+ hsa_ext_image_descriptor_t imageDescriptor;
+ imageDescriptor.geometry = geometry;
+ imageDescriptor.width = width;
+ imageDescriptor.height = height;
+ imageDescriptor.depth = depth;
+ imageDescriptor.array_size = array_size;
+ hsa_ext_image_channel_order_t channelOrder;
+ hsa_ext_image_channel_type_t channelType;
+ getChannelOrderAndType(desc, hipReadModeElementType, &channelOrder, &channelType);
+ imageDescriptor.format.channel_order = channelOrder;
+ imageDescriptor.format.channel_type = channelType;
+
+ hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+ // Get the current device agent.
+ hc::accelerator acc;
+ hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
+ if (!agent)
+ return hipErrorInvalidResourceHandle;
+ hsa_status_t status =
+ hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
+ if(HSA_STATUS_SUCCESS != status){
+ return hipErrorRuntimeOther;
+ }
+
+ return hipSuccess;
+}
+
+hipError_t ihipArrayToImageFormat(hipArray_Format format,hsa_ext_image_channel_type_t &channelType) {
+ switch (format) {
+ case HIP_AD_FORMAT_UNSIGNED_INT8:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
+ break;
+ case HIP_AD_FORMAT_UNSIGNED_INT16:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
+ break;
+ case HIP_AD_FORMAT_UNSIGNED_INT32:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
+ break;
+ case HIP_AD_FORMAT_SIGNED_INT8:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
+ break;
+ case HIP_AD_FORMAT_SIGNED_INT16:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
+ break;
+ case HIP_AD_FORMAT_SIGNED_INT32:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
+ break;
+ case HIP_AD_FORMAT_HALF:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
+ break;
+ case HIP_AD_FORMAT_FLOAT:
+ channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
+ break;
+ default:
+ return hipErrorUnknown;
+ break;
+ }
+ return hipSuccess;
+}
+
hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
HIP_INIT_SPECIAL_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if (pAllocateArray->Width > 0) {
- auto ctx = ihipGetTlsDefaultCtx();
*array = (hipArray*)malloc(sizeof(hipArray));
- HIP_ARRAY3D_DESCRIPTOR array3D;
- array3D.Width = pAllocateArray->Width;
- array3D.Height = pAllocateArray->Height;
- array3D.Format = pAllocateArray->Format;
- array3D.NumChannels = pAllocateArray->NumChannels;
array[0]->width = pAllocateArray->Width;
array[0]->height = pAllocateArray->Height;
array[0]->Format = pAllocateArray->Format;
@@ -560,100 +621,26 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat
array[0]->isDrv = true;
array[0]->textureType = hipTextureType2D;
void** ptr = &array[0]->data;
- if (ctx) {
- const unsigned am_flags = 0;
- size_t size = pAllocateArray->Width;
- if (pAllocateArray->Height > 0) {
- size = size * pAllocateArray->Height;
- }
- hsa_ext_image_channel_type_t channelType;
- size_t allocSize = 0;
- switch (pAllocateArray->Format) {
- case HIP_AD_FORMAT_UNSIGNED_INT8:
- allocSize = size * sizeof(uint8_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
- break;
- case HIP_AD_FORMAT_UNSIGNED_INT16:
- allocSize = size * sizeof(uint16_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
- break;
- case HIP_AD_FORMAT_UNSIGNED_INT32:
- allocSize = size * sizeof(uint32_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
- break;
- case HIP_AD_FORMAT_SIGNED_INT8:
- allocSize = size * sizeof(int8_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
- break;
- case HIP_AD_FORMAT_SIGNED_INT16:
- allocSize = size * sizeof(int16_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
- break;
- case HIP_AD_FORMAT_SIGNED_INT32:
- allocSize = size * sizeof(int32_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
- break;
- case HIP_AD_FORMAT_HALF:
- allocSize = size * sizeof(int16_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
- break;
- case HIP_AD_FORMAT_FLOAT:
- allocSize = size * sizeof(float);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
- break;
- default:
- hip_status = hipErrorUnknown;
- break;
- }
- hc::accelerator acc = ctx->getDevice()->_acc;
- hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
+ hsa_ext_image_channel_type_t channelType;
+ hsa_ext_image_channel_order_t channelOrder;
- size_t allocGranularity = 0;
- hsa_amd_memory_pool_t* allocRegion =
- static_cast(acc.get_hsa_am_region());
- hsa_amd_memory_pool_get_info(
- *allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
+ hip_status = ihipArrayToImageFormat(pAllocateArray->Format,channelType);
+ if(hipSuccess != hip_status)
+ return ihipLogStatus(hip_status);
- hsa_ext_image_descriptor_t imageDescriptor;
-
- imageDescriptor.width = pAllocateArray->Width;
- imageDescriptor.height = pAllocateArray->Height;
- imageDescriptor.depth = 0;
- imageDescriptor.array_size = 0;
-
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
-
- hsa_ext_image_channel_order_t channelOrder;
-
- if (pAllocateArray->NumChannels == 4) {
- channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
- } else if (pAllocateArray->NumChannels == 2) {
- channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
- } else if (pAllocateArray->NumChannels == 1) {
- channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
- }
- imageDescriptor.format.channel_order = channelOrder;
- imageDescriptor.format.channel_type = channelType;
-
- hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
- hsa_ext_image_data_info_t imageInfo;
- hsa_status_t status =
- hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
- size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
-
- *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx,
- false /*shareWithAll*/, am_flags, 0, alignment);
- if (size && (*ptr == NULL)) {
- hip_status = hipErrorMemoryAllocation;
- }
- } else {
- hip_status = hipErrorMemoryAllocation;
+ if (pAllocateArray->NumChannels == 4) {
+ channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
+ } else if (pAllocateArray->NumChannels == 2) {
+ channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
+ } else if (pAllocateArray->NumChannels == 1) {
+ channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
+ hsa_ext_image_data_info_t imageInfo;
+ return ihipLogStatus(allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,pAllocateArray->Width,
+ pAllocateArray->Height,0,channelOrder,channelType,ptr,imageInfo));
} else {
- hip_status = hipErrorInvalidValue;
+ return ihipLogStatus(hipErrorInvalidValue);
}
-
- return ihipLogStatus(hip_status);
}
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width,
@@ -662,8 +649,6 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if (width > 0) {
- auto ctx = ihipGetTlsDefaultCtx();
-
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = flags;
array[0]->width = width;
@@ -674,67 +659,25 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si
array[0]->textureType = hipTextureType2D;
void** ptr = &array[0]->data;
- if (ctx) {
- const unsigned am_flags = 0;
- size_t size = width;
- if (height > 0) {
- size = size * height;
- }
-
- const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8);
-
- hc::accelerator acc = ctx->getDevice()->_acc;
- hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
-
- size_t allocGranularity = 0;
- hsa_amd_memory_pool_t* allocRegion =
- static_cast(acc.get_hsa_am_region());
- hsa_amd_memory_pool_get_info(
- *allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
-
- hsa_ext_image_descriptor_t imageDescriptor;
-
- imageDescriptor.width = width;
- imageDescriptor.height = height;
- imageDescriptor.depth = 0;
- imageDescriptor.array_size = 0;
- switch (flags) {
- case hipArrayLayered:
- case hipArrayCubemap:
- case hipArraySurfaceLoadStore:
- case hipArrayTextureGather:
- assert(0);
- break;
- case hipArrayDefault:
- default:
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
- break;
- }
- hsa_ext_image_channel_order_t channelOrder;
- hsa_ext_image_channel_type_t channelType;
- getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
- imageDescriptor.format.channel_order = channelOrder;
- imageDescriptor.format.channel_type = channelType;
-
- hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
- hsa_ext_image_data_info_t imageInfo;
- hsa_status_t status =
- hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
- size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
-
- *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx,
- false /*shareWithAll*/, am_flags, 0, alignment);
- if (size && (*ptr == NULL)) {
- hip_status = hipErrorMemoryAllocation;
- }
-
- } else {
- hip_status = hipErrorMemoryAllocation;
+ hsa_ext_image_channel_order_t channelOrder;
+ hsa_ext_image_channel_type_t channelType;
+ getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
+ hsa_ext_image_data_info_t imageInfo;
+ switch (flags) {
+ case hipArrayLayered:
+ case hipArrayCubemap:
+ case hipArraySurfaceLoadStore:
+ case hipArrayTextureGather:
+ assert(0);
+ break;
+ case hipArrayDefault:
+ default:
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,width,height,0,channelOrder,channelType,ptr,imageInfo);
+ break;
}
} else {
hip_status = hipErrorInvalidValue;
}
-
return ihipLogStatus(hip_status);
}
@@ -742,8 +685,6 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
HIP_INIT_SPECIAL_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray);
hipError_t hip_status = hipSuccess;
- auto ctx = ihipGetTlsDefaultCtx();
-
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = pAllocateArray->Flags;
array[0]->width = pAllocateArray->Width;
@@ -752,111 +693,37 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
array[0]->Format = pAllocateArray->Format;
array[0]->NumChannels = pAllocateArray->NumChannels;
array[0]->isDrv = true;
- array[0]->textureType = hipTextureType3D;
void** ptr = &array[0]->data;
- if (ctx) {
- const unsigned am_flags = 0;
- const size_t size = pAllocateArray->Width * pAllocateArray->Height * pAllocateArray->Depth;
+ hsa_ext_image_channel_order_t channelOrder;
+ if (pAllocateArray->NumChannels == 4) {
+ channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
+ } else if (pAllocateArray->NumChannels == 2) {
+ channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
+ } else if (pAllocateArray->NumChannels == 1) {
+ channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
+ }
- size_t allocSize = 0;
- hsa_ext_image_channel_type_t channelType;
- switch (pAllocateArray->Format) {
- case HIP_AD_FORMAT_UNSIGNED_INT8:
- allocSize = size * sizeof(uint8_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
- break;
- case HIP_AD_FORMAT_UNSIGNED_INT16:
- allocSize = size * sizeof(uint16_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
- break;
- case HIP_AD_FORMAT_UNSIGNED_INT32:
- allocSize = size * sizeof(uint32_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
- break;
- case HIP_AD_FORMAT_SIGNED_INT8:
- allocSize = size * sizeof(int8_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
- break;
- case HIP_AD_FORMAT_SIGNED_INT16:
- allocSize = size * sizeof(int16_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
- break;
- case HIP_AD_FORMAT_SIGNED_INT32:
- allocSize = size * sizeof(int32_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
- break;
- case HIP_AD_FORMAT_HALF:
- allocSize = size * sizeof(int16_t);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
- break;
- case HIP_AD_FORMAT_FLOAT:
- allocSize = size * sizeof(float);
- channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
- break;
- default:
- hip_status = hipErrorUnknown;
- break;
- }
-
- hc::accelerator acc = ctx->getDevice()->_acc;
- hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
-
- size_t allocGranularity = 0;
- hsa_amd_memory_pool_t* allocRegion =
- static_cast(acc.get_hsa_am_region());
- hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
- &allocGranularity);
-
- hsa_ext_image_descriptor_t imageDescriptor;
- imageDescriptor.width = pAllocateArray->Width;
- imageDescriptor.height = pAllocateArray->Height;
- imageDescriptor.depth = 0;
- imageDescriptor.array_size = 0;
- switch (pAllocateArray->Flags) {
- case hipArrayLayered:
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA;
- imageDescriptor.array_size = pAllocateArray->Depth;
- break;
- case hipArraySurfaceLoadStore:
- case hipArrayTextureGather:
- case hipArrayDefault:
- assert(0);
- break;
- case hipArrayCubemap:
- default:
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D;
- imageDescriptor.depth = pAllocateArray->Depth;
- break;
- }
- hsa_ext_image_channel_order_t channelOrder;
-
- // getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
- if (pAllocateArray->NumChannels == 4) {
- channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
- } else if (pAllocateArray->NumChannels == 2) {
- channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
- } else if (pAllocateArray->NumChannels == 1) {
- channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
- }
- imageDescriptor.format.channel_order = channelOrder;
- imageDescriptor.format.channel_type = channelType;
-
- hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
- hsa_ext_image_data_info_t imageInfo;
- hsa_status_t status =
- hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
- size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
-
- *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0,
- alignment);
-
- if (size && (*ptr == NULL)) {
- hip_status = hipErrorMemoryAllocation;
- }
-
- } else {
- hip_status = hipErrorMemoryAllocation;
+ hsa_ext_image_channel_type_t channelType;
+ hip_status = ihipArrayToImageFormat(pAllocateArray->Format,channelType);
+ hsa_ext_image_data_info_t imageInfo;
+ switch (pAllocateArray->Flags) {
+ case hipArrayLayered:
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2DA,pAllocateArray->Width,pAllocateArray->Height,0,
+ channelOrder,channelType,ptr,imageInfo,pAllocateArray->Depth);
+ array[0]->textureType = hipTextureType2DLayered;
+ break;
+ case hipArraySurfaceLoadStore:
+ case hipArrayTextureGather:
+ assert(0);
+ break;
+ case hipArrayDefault:
+ case hipArrayCubemap:
+ default:
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,pAllocateArray->Width,pAllocateArray->Height,
+ pAllocateArray->Depth,channelOrder,channelType,ptr,imageInfo);
+ array[0]->textureType = hipTextureType3D;
+ break;
}
return ihipLogStatus(hip_status);
@@ -865,19 +732,13 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc,
struct hipExtent extent, unsigned int flags) {
-
-
HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
- if(array==NULL )
- {
- hip_status=hipErrorInvalidValue;
- return ihipLogStatus(hip_status);
+ if(array==NULL ){
+ return ihipLogStatus(hipErrorInvalidValue);
}
- auto ctx = ihipGetTlsDefaultCtx();
-
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = flags;
array[0]->width = extent.width;
@@ -885,68 +746,27 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
array[0]->depth = extent.depth;
array[0]->desc = *desc;
array[0]->isDrv = false;
- array[0]->textureType = hipTextureType3D;
void** ptr = &array[0]->data;
-
- if (ctx) {
- const unsigned am_flags = 0;
- const size_t size = extent.width * extent.height * extent.depth;
-
- const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8);
-
- hc::accelerator acc = ctx->getDevice()->_acc;
- hsa_agent_t* agent = static_cast(acc.get_hsa_agent());
-
- size_t allocGranularity = 0;
- hsa_amd_memory_pool_t* allocRegion =
- static_cast(acc.get_hsa_am_region());
- hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
- &allocGranularity);
-
- hsa_ext_image_descriptor_t imageDescriptor;
- imageDescriptor.width = extent.width;
- imageDescriptor.height = extent.height;
- imageDescriptor.depth = extent.depth;
- imageDescriptor.array_size = 0;
- switch (flags) {
- case hipArrayLayered:
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA;
- imageDescriptor.array_size = extent.depth;
- break;
- case hipArraySurfaceLoadStore:
- case hipArrayTextureGather:
- case hipArrayDefault:
- assert(0);
- break;
- case hipArrayCubemap:
- default:
- imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D;
- imageDescriptor.depth = extent.depth;
- break;
- }
- hsa_ext_image_channel_order_t channelOrder;
- hsa_ext_image_channel_type_t channelType;
- getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
- imageDescriptor.format.channel_order = channelOrder;
- imageDescriptor.format.channel_type = channelType;
-
- hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
- hsa_ext_image_data_info_t imageInfo;
- hsa_status_t status =
- hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
- size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
-
- *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0,
- alignment);
-
- if (size && (*ptr == NULL)) {
- hip_status = hipErrorMemoryAllocation;
- }
-
- } else {
- hip_status = hipErrorMemoryAllocation;
+ hsa_ext_image_channel_order_t channelOrder;
+ hsa_ext_image_channel_type_t channelType;
+ getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
+ hsa_ext_image_data_info_t imageInfo;
+ switch (flags) {
+ case hipArrayLayered:
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2DA,extent.width,extent.height,0,channelOrder,channelType,ptr,imageInfo,extent.depth);
+ array[0]->textureType = hipTextureType2DLayered;
+ break;
+ case hipArraySurfaceLoadStore:
+ case hipArrayTextureGather:
+ assert(0);
+ break;
+ case hipArrayDefault:
+ case hipArrayCubemap:
+ default:
+ hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,extent.width,extent.height,extent.depth,channelOrder,channelType,ptr,imageInfo);
+ array[0]->textureType = hipTextureType3D;
+ break;
}
-
return ihipLogStatus(hip_status);
}
@@ -1208,7 +1028,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
-
+
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
@@ -1246,7 +1066,6 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
return ihipLogStatus(e);
}
-
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
@@ -1275,7 +1094,7 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
-
+
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
@@ -1283,7 +1102,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
-
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false);
} catch (ihipException& ex) {
@@ -1328,8 +1146,6 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
- hc::completion_future marker;
-
hipError_t e = hipSuccess;
size_t byteSize;
@@ -1380,8 +1196,6 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
- hc::completion_future marker;
-
hipError_t e = hipSuccess;
try {
@@ -1399,8 +1213,6 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
- hc::completion_future marker;
-
hipError_t e = hipSuccess;
try {
@@ -1417,8 +1229,6 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
- hc::completion_future marker;
-
hipError_t e = hipSuccess;
try {
stream->locked_copySync((char*)dstArray->data + dstOffset, srcHost, count,
@@ -1435,8 +1245,6 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
- hc::completion_future marker;
-
hipError_t e = hipSuccess;
try {
@@ -1452,15 +1260,9 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) {
hipError_t e = hipSuccess;
if(p) {
- size_t byteSize;
- size_t depth;
- size_t height;
- size_t widthInBytes;
- size_t srcPitch;
- size_t dstPitch;
- void* srcPtr;
- void* dstPtr;
- size_t ySize;
+ size_t byteSize, width, height, depth, widthInBytes, srcPitch, dstPitch, ySize;
+ hipChannelFormatDesc desc;
+ void* srcPtr;void* dstPtr;
if (p->dstArray != nullptr) {
if (p->dstArray->isDrv == false) {
switch (p->dstArray->desc.f) {
@@ -1482,22 +1284,32 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo
}
depth = p->extent.depth;
height = p->extent.height;
+ width = p->extent.width;
widthInBytes = p->extent.width * byteSize;
srcPitch = p->srcPtr.pitch;
srcPtr = p->srcPtr.ptr;
ySize = p->srcPtr.ysize;
- dstPitch = p->dstArray->width * byteSize;
+ desc = p->dstArray->desc;
dstPtr = p->dstArray->data;
} else {
depth = p->Depth;
height = p->Height;
widthInBytes = p->WidthInBytes;
- dstPitch = p->dstArray->width * 4;
+ width = p->dstArray->width;
+ desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
srcPitch = p->srcPitch;
srcPtr = (void*)p->srcHost;
ySize = p->srcHeight;
dstPtr = p->dstArray->data;
}
+ hsa_ext_image_data_info_t imageInfo;
+ if(hipTextureType2DLayered == p->dstArray->textureType)
+ GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth);
+ else
+ GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo);
+
+ dstPitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
+
} else {
// Non array destination
depth = p->extent.depth;
@@ -1509,8 +1321,8 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo
ySize = p->srcPtr.ysize;
dstPitch = p->dstPtr.pitch;
}
+
stream = ihipSyncAndResolveStream(stream);
- hc::completion_future marker;
try {
if((widthInBytes == dstPitch) && (widthInBytes == srcPitch)) {
if(isAsync)
@@ -1556,16 +1368,23 @@ hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms* p, hipStream_t stream
}
namespace {
-template
+template
__global__ void hip_fill_n(RandomAccessIterator f, N n, T value) {
- const uint32_t grid_dim = gridDim.x * blockDim.x;
+ const auto grid_dim = gridDim.x * blockDim.x * items_per_lane;
+ const auto gidx = blockIdx.x * block_dim + threadIdx.x;
- size_t idx = blockIdx.x * block_dim + threadIdx.x;
- while (idx < n) {
- __builtin_memcpy(reinterpret_cast(&f[idx]), reinterpret_cast(&value),
- sizeof(T));
+ size_t idx = gidx * items_per_lane;
+ while (idx + items_per_lane <= n) {
+ for (auto i = 0u; i != items_per_lane; ++i) {
+ __builtin_nontemporal_store(value, &f[idx + i]);
+ }
idx += grid_dim;
}
+
+ if (gidx < n % grid_dim) {
+ __builtin_nontemporal_store(value, &f[n - gidx - 1]);
+ }
}
template {}>::type* = nullptr>
@@ -1620,11 +1439,14 @@ hipError_t ihipMemPtrGetInfo(void* ptr, size_t* size) {
template
void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t count) {
static constexpr uint32_t block_dim = 256;
+ static constexpr uint32_t max_write_width = 4 * sizeof(std::uint32_t); // 4 DWORDs
+ static constexpr uint32_t items_per_lane = max_write_width / sizeof(T);
- const uint32_t grid_dim = clamp_integer(count / block_dim, 1, UINT32_MAX);
+ const uint32_t grid_dim = clamp_integer(
+ count / (block_dim * items_per_lane), 1, UINT32_MAX);
- hipLaunchKernelGGL(hip_fill_n, dim3(grid_dim), dim3{block_dim}, 0u, stream, ptr,
- count, std::move(val));
+ hipLaunchKernelGGL(hip_fill_n, dim3(grid_dim),
+ dim3{block_dim}, 0u, stream, ptr, count, std::move(val));
}
template
@@ -1643,62 +1465,133 @@ typedef enum ihipMemsetDataType {
ihipMemsetDataTypeInt = 2
}ihipMemsetDataType;
-hipError_t ihipMemset(void* dst, int value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType)
-{
- hipError_t e = hipSuccess;
+hipError_t ihipMemsetAsync(void* dst, int value, size_t count, hipStream_t stream, enum ihipMemsetDataType copyDataType) {
+ if (count == 0) return hipSuccess;
+ if (!dst) return hipErrorInvalidValue;
- if (count == 0) return e;
-
- size_t allocSize = 0;
- bool isInbound = (ihipMemPtrGetInfo(dst, &allocSize) == hipSuccess);
- isInbound &= (allocSize >= count);
-
- if (stream && (dst != NULL) && isInbound) {
- if(copyDataType == ihipMemsetDataTypeChar){
+ try {
+ if (copyDataType == ihipMemsetDataTypeChar) {
if ((count & 0x3) == 0) {
// use a faster dword-per-workitem copy:
- try {
- value = value & 0xff;
- uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
- ihipMemsetKernel (stream, static_cast (dst), value32, count/sizeof(uint32_t));
- }
- catch (std::exception &ex) {
- e = hipErrorInvalidValue;
- }
- } else {
+ value = value & 0xff;
+ uint32_t value32 = (value << 24) | (value << 16) | (value << 8) | (value) ;
+ ihipMemsetKernel (stream, static_cast (dst), value32, count/sizeof(uint32_t));
+ } else {
// use a slow byte-per-workitem copy:
- try {
- ihipMemsetKernel (stream, static_cast (dst), value, count);
- }
- catch (std::exception &ex) {
- e = hipErrorInvalidValue;
- }
- }
- } else {
- if(copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value
- try {
- ihipMemsetKernel (stream, static_cast (dst), value, count);
- } catch (std::exception &ex) {
- e = hipErrorInvalidValue;
- }
- } else if(copyDataType == ihipMemsetDataTypeShort) {
- try {
- value = value & 0xffff;
- ihipMemsetKernel (stream, static_cast (dst), value, count);
- } catch (std::exception &ex) {
- e = hipErrorInvalidValue;
- }
+ ihipMemsetKernel (stream, static_cast (dst), value, count);
}
+ } else if (copyDataType == ihipMemsetDataTypeInt) { // 4 Bytes value
+ ihipMemsetKernel (stream, static_cast (dst), value, count);
+ } else if (copyDataType == ihipMemsetDataTypeShort) {
+ value = value & 0xffff;
+ ihipMemsetKernel (stream, static_cast (dst), value, count);
}
- if (HIP_API_BLOCKING) {
- tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str());
- stream->locked_wait();
- }
- } else {
- e = hipErrorInvalidValue;
+ } catch (...) {
+ return hipErrorInvalidValue;
}
- return e;
-};
+
+ if (HIP_API_BLOCKING) {
+ tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str());
+ stream->locked_wait();
+ }
+
+ return hipSuccess;
+}
+
+namespace {
+ template
+ void handleHeadTail(T* dst, std::size_t n_head, std::size_t n_body,
+ std::size_t n_tail, hipStream_t stream, int value) {
+ struct Cleaner {
+ static
+ __global__
+ void clean(T* p, std::size_t nh, std::size_t nb, int x) noexcept {
+ p[(threadIdx.x < nh) ? threadIdx.x : (threadIdx.x - nh + nb)] = x;
+ }
+ };
+
+ hipLaunchKernelGGL(Cleaner::clean, 1, n_head + n_tail, 0, stream,
+ dst, n_head,
+ n_body * sizeof(std::uint32_t) / sizeof(T), value);
+
+ }
+} // Anonymous namespace.
+
+hipError_t ihipMemsetSync(void* dst, int value, size_t count, hipStream_t stream, ihipMemsetDataType copyDataType) {
+ if (count == 0) return hipSuccess;
+ if (!dst) return hipErrorInvalidValue;
+
+ try {
+ size_t n = count;
+ auto aligned_dst{(copyDataType == ihipMemsetDataTypeInt) ? dst :
+ reinterpret_cast(
+ hip_impl::round_up_to_next_multiple_nonnegative(
+ reinterpret_cast(dst), 4ul))};
+ size_t n_head{};
+ size_t n_tail{};
+ int original_value = value;
+
+ switch (copyDataType) {
+ case ihipMemsetDataTypeChar:
+ value &= 0xff;
+ value = (value << 24) | (value << 16) | (value << 8) | value;
+ n_head = static_cast(aligned_dst) -
+ static_cast(dst);
+ n -= n_head;
+ n /= sizeof(std::uint32_t);
+ n_tail = count % sizeof(std::uint32_t);
+ break;
+ case ihipMemsetDataTypeShort:
+ value &= 0xffff;
+ value = (value << 16) | value;
+ n_head = static_cast(aligned_dst) -
+ static_cast(dst);
+ n = (count - n_head) *
+ sizeof(std::uint16_t) / sizeof(std::uint32_t);
+ n_tail = ((count - n_head) *
+ sizeof(std::uint16_t)) % sizeof(std::uint32_t);
+ break;
+ default: break;
+ }
+
+ // queue the memset kernel for the remainder of the buffer before the HSA call below
+ if (aligned_dst != dst || n_tail != 0) {
+ switch (copyDataType) {
+ case ihipMemsetDataTypeChar:
+ handleHeadTail(static_cast(dst), n_head, n,
+ n_tail, stream, value & 0xff);
+ break;
+ case ihipMemsetDataTypeShort:
+ handleHeadTail(static_cast(dst), n_head, n,
+ n_tail, stream, value & 0xffff);
+ break;
+ default: break;
+ }
+ }
+
+ // The stream must be locked from all other op insertions to guarantee
+ // that the following HSA call can complete before any other ops.
+ // Flush the stream while locked. Once the stream is empty, we can safely perform
+ // the out-of-band HSA call. Lastly, the stream will unlock via RAII.
+ if (!stream) stream = ihipSyncAndResolveStream(stream);
+ if (!stream) return hipErrorInvalidValue;
+
+ LockedAccessor_StreamCrit_t crit(stream->criticalData());
+ crit->_av.wait(stream->waitMode());
+ const auto s = hsa_amd_memory_fill(aligned_dst, value, n);
+ if (s != HSA_STATUS_SUCCESS) return hipErrorInvalidValue;
+ }
+ catch (...) {
+ return hipErrorInvalidValue;
+ }
+
+ if (HIP_API_BLOCKING) {
+ tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetSync.\n", ToString(stream).c_str());
+ stream->locked_wait();
+ }
+
+ return hipSuccess;
+}
hipError_t getLockedPointer(void *hostPtr, size_t dataLen, void **devicePtrPtr)
{
@@ -1716,7 +1609,7 @@ hipError_t getLockedPointer(void *hostPtr, size_t dataLen, void **devicePtrPtr)
return(hipSuccess);
};
return(hipErrorHostMemoryNotRegistered);
-};
+}
// TODO - review and optimize
hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
@@ -1955,182 +1848,66 @@ hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream);
-
- hipError_t e = hipSuccess;
-
- stream = ihipSyncAndResolveStream(stream);
-
- e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
-
- return ihipLogStatus(e);
-};
+ return ihipLogStatus(ihipMemsetAsync(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar));
+}
hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipMemsetD32Async, (TRACE_MCMD), dst, value, count, stream);
-
- hipError_t e = hipSuccess;
-
- stream = ihipSyncAndResolveStream(stream);
-
- e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeInt);
-
- return ihipLogStatus(e);
-};
+ return ihipLogStatus(ihipMemsetAsync(dst, value, count, stream, ihipMemsetDataTypeInt));
+}
hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemset, (TRACE_MCMD), dst, value, sizeBytes);
-
- hipError_t e = hipSuccess;
-
- hipStream_t stream = hipStreamNull;
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
- stream->locked_wait();
- } else {
- e = hipErrorInvalidValue;
- }
- return ihipLogStatus(e);
+ return ihipLogStatus(ihipMemsetSync(dst, value, sizeBytes, nullptr, ihipMemsetDataTypeChar));
}
hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
HIP_INIT_SPECIAL_API(hipMemset2D, (TRACE_MCMD), dst, pitch, value, width, height);
-
- hipError_t e = hipSuccess;
-
- hipStream_t stream = hipStreamNull;
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- size_t sizeBytes = pitch * height;
- e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
- stream->locked_wait();
- } else {
- e = hipErrorInvalidValue;
- }
-
- return ihipLogStatus(e);
+ size_t sizeBytes = pitch * height;
+ return ihipLogStatus(ihipMemsetSync(dst, value, sizeBytes, nullptr, ihipMemsetDataTypeChar));
}
-hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream )
-{
+hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream ) {
HIP_INIT_SPECIAL_API(hipMemset2DAsync, (TRACE_MCMD), dst, pitch, value, width, height, stream);
-
- hipError_t e = hipSuccess;
-
- stream = ihipSyncAndResolveStream(stream);
-
- if (stream) {
- size_t sizeBytes = pitch * height;
- e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar);
- } else {
- e = hipErrorInvalidValue;
- }
-
- return ihipLogStatus(e);
-};
+ size_t sizeBytes = pitch * height;
+ return ihipLogStatus(ihipMemsetAsync(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar));
+}
hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) {
HIP_INIT_SPECIAL_API(hipMemsetD8, (TRACE_MCMD), dst, value, count);
-
- hipError_t e = hipSuccess;
-
- hipStream_t stream = hipStreamNull;
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeChar);
- stream->locked_wait();
- } else {
- e = hipErrorInvalidValue;
- }
- return ihipLogStatus(e);
+ return ihipLogStatus(ihipMemsetSync(dst, value, count, nullptr, ihipMemsetDataTypeChar));
}
hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t count , hipStream_t stream ) {
HIP_INIT_SPECIAL_API(hipMemsetD8Async, (TRACE_MCMD), dst, value, count, stream);
-
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- return ihipLogStatus(ihipMemset(dst, value, count, stream, ihipMemsetDataTypeChar));
- } else {
- return ihipLogStatus(hipErrorInvalidValue);
- }
+ return ihipLogStatus(ihipMemsetAsync(dst, value, count, stream, ihipMemsetDataTypeChar));
}
hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t count){
HIP_INIT_SPECIAL_API(hipMemsetD16, (TRACE_MCMD), dst, value, count);
- hipError_t e = hipSuccess;
- hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
- if (stream) {
- e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeShort);
- if(hipSuccess == e)
- stream->locked_wait();
- } else {
- e = hipErrorInvalidValue;
- }
- return ihipLogStatus(e);
+ return ihipLogStatus(ihipMemsetSync(dst, value, count, nullptr, ihipMemsetDataTypeShort));
}
hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t count, hipStream_t stream ){
HIP_INIT_SPECIAL_API(hipMemsetD16Async, (TRACE_MCMD), dst, value, count, stream);
-
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- return ihipLogStatus(ihipMemset(dst, value, count, stream, ihipMemsetDataTypeShort));
- } else {
- return ihipLogStatus(hipErrorInvalidValue);
- }
+ return ihipLogStatus(ihipMemsetAsync(dst, value, count, stream, ihipMemsetDataTypeShort));
}
hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) {
HIP_INIT_SPECIAL_API(hipMemsetD32, (TRACE_MCMD), dst, value, count);
-
- hipError_t e = hipSuccess;
-
- hipStream_t stream = hipStreamNull;
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- e = ihipMemset(dst, value, count, stream, ihipMemsetDataTypeInt);
- stream->locked_wait();
- } else {
- e = hipErrorInvalidValue;
- }
- return ihipLogStatus(e);
+ return ihipLogStatus(ihipMemsetSync(dst, value, count, nullptr, ihipMemsetDataTypeInt));
}
-hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent )
-{
+hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) {
HIP_INIT_SPECIAL_API(hipMemset3D, (TRACE_MCMD), &pitchedDevPtr, value, &extent);
- hipError_t e = hipSuccess;
-
- hipStream_t stream = hipStreamNull;
- // TODO - call an ihip memset so HIP_TRACE is correct.
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
- e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar);
- stream->locked_wait();
- } else {
- e = hipErrorInvalidValue;
- }
-
- return ihipLogStatus(e);
+ size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
+ return ihipLogStatus(ihipMemsetSync(pitchedDevPtr.ptr, value, sizeBytes, nullptr, ihipMemsetDataTypeChar));
}
-hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream )
-{
+hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream ) {
HIP_INIT_SPECIAL_API(hipMemset3DAsync, (TRACE_MCMD), &pitchedDevPtr, value, &extent);
- hipError_t e = hipSuccess;
-
- // TODO - call an ihip memset so HIP_TRACE is correct.
- stream = ihipSyncAndResolveStream(stream);
- if (stream) {
- size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
- e = ihipMemset(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar);
- } else {
- e = hipErrorInvalidValue;
- }
-
- return ihipLogStatus(e);
+ size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
+ return ihipLogStatus(ihipMemsetAsync(pitchedDevPtr.ptr, value, sizeBytes, stream, ihipMemsetDataTypeChar));
}
hipError_t hipMemGetInfo(size_t* free, size_t* total) {
@@ -2146,20 +1923,20 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) {
} else {
e = hipErrorInvalidValue;
}
-
+
if (free) {
if (!device->_driver_node_id) return ihipLogStatus(hipErrorInvalidDevice);
-
- std::string fileName = std::string("/sys/class/kfd/kfd/topology/nodes/") + std::to_string(device->_driver_node_id) + std::string("/mem_banks/0/used_memory");
+
+ std::string fileName = std::string("/sys/class/kfd/kfd/topology/nodes/") + std::to_string(device->_driver_node_id) + std::string("/mem_banks/0/used_memory");
std::ifstream file;
file.open(fileName);
if (!file) return ihipLogStatus(hipErrorFileNotFound);
-
- std::string deviceSize;
+
+ std::string deviceSize;
size_t deviceMemSize;
-
+
file >> deviceSize;
- file.close();
+ file.close();
if ((deviceMemSize=strtol(deviceSize.c_str(),NULL,10))){
*free = device->_props.totalGlobalMem - deviceMemSize;
// Deduct the amount of memory from the free memory reported from the system
diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp
index 2afbabf0a8..498ee5bb03 100644
--- a/projects/hip/src/hip_module.cpp
+++ b/projects/hip/src/hip_module.cpp
@@ -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"
diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp
index e3e4975b7e..bf3289bb49 100644
--- a/projects/hip/src/hip_stream.cpp
+++ b/projects/hip/src/hip_stream.cpp
@@ -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 (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);
}
diff --git a/projects/hip/src/hip_texture.cpp b/projects/hip/src/hip_texture.cpp
index 8820878fc0..7ea8da8624 100644
--- a/projects/hip/src/hip_texture.cpp
+++ b/projects/hip/src/hip_texture.cpp
@@ -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);
}
diff --git a/projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu b/projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu
new file mode 100644
index 0000000000..3795d3c799
--- /dev/null
+++ b/projects/hip/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu
@@ -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
+#include
+
+#define CUDA_LAUNCH(cuda_call,dimGrid,dimBlock, ...) \
+ cuda_call<<>>(__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);
+}
diff --git a/projects/hip/tests/src/hipHcc.cpp b/projects/hip/tests/src/hipHcc.cpp
index e0a843d4f5..2d4523ecd6 100644
--- a/projects/hip/tests/src/hipHcc.cpp
+++ b/projects/hip/tests/src/hipHcc.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
#include
#include
#include "hip/hip_runtime.h"
-#include "hip/hip_hcc.h"
+#include "hip/hip_ext.h"
#include "test_common.h"
#define CHECK(error) \
diff --git a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
index 89d003ea94..c76685fa89 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
@@ -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
diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
index 956ba4e50a..f98d5e4ec4 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
@@ -32,7 +32,7 @@ THE SOFTWARE.
#include
#include
#include
-#include
+#include
#define LEN 64
#define SIZE LEN * sizeof(float)
diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
index e678c35d6e..9ae5883608 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
@@ -27,11 +27,9 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
-//#include "hip/hip_runtime_api.h"
#include
#include
#include
-//#include
#define fileName "tex2d_kernel.code"
diff --git a/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp b/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp
new file mode 100644
index 0000000000..e6492c7ce2
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp
@@ -0,0 +1,145 @@
+#include
+#include
+#include
+#include "test_common.h"
+#include
+
+/* 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 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(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;
+}
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp
new file mode 100644
index 0000000000..b2fb168fbc
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp
@@ -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
+#include
+#include
+
+
+//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;iresult[i] != data->copyOfOriginalData[i]+1)
+ {
+ std::cout<<"Error value : "<result[i]<<"| Expected value :"<copyOfOriginalData[i]+1<result = hResultData;
+ inputParam->copyOfOriginalData = hData;
+
+ HIPCHECK(hipStreamCreate(&stream));
+ HIPCHECK(hipStreamAddCallback(stream,callback,inputParam,0));
+
+ // Wait for stream add callback to complete
+ std::unique_lock l(gMutx);
+
+ while(!callbackCompleted)
+ condVar.wait(l);
+
+ // Will destroy device memory hence explicite hipFree is not needed
+ HIPCHECK(hipDeviceReset());
+
+ free(hData);
+ free(hResultData);
+ passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp
new file mode 100644
index 0000000000..3babf5869e
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp
@@ -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
+#include
+#include
+#include
+#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 :"<copyOf_hData[i]< 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;ihData[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 l(completionSignal.mu);
+ while(completionSignal.completedThreads != numWorkloads)
+ {
+ completionSignal.cv.wait(l);
+ }
+
+ //clean-up
+ hipDeviceReset();
+ free(workloads);
+
+ passed();
+ return 0;
+}
diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h
index 67a8e5e60a..92528e38e3 100644
--- a/projects/hip/tests/src/test_common.h
+++ b/projects/hip/tests/src/test_common.h
@@ -153,6 +153,36 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg);
unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N);
+template // 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
+void checkArray(T input, T output, size_t height, size_t width)
+{
+ for(int i=0; i
__global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) {
diff --git a/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp b/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp
new file mode 100644
index 0000000000..f463942329
--- /dev/null
+++ b/projects/hip/tests/src/texture/hipBindTex2DPitch.cpp
@@ -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 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();
+}
diff --git a/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp b/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp
index fd2e8ccf0f..0a9c879376 100644
--- a/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp
+++ b/projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp
@@ -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;
diff --git a/projects/hip/tests/src/texture/hipTexObjPitch.cpp b/projects/hip/tests/src/texture/hipTexObjPitch.cpp
new file mode 100644
index 0000000000..47648d5e73
--- /dev/null
+++ b/projects/hip/tests/src/texture/hipTexObjPitch.cpp
@@ -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
+__global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) {
+
+ for(int i =0;i(texObj, j, i);
+ __syncthreads();
+}
+
+template
+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();
+
+ 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();
+ texture2Dtest();
+ texture2Dtest();
+ texture2Dtest();
+ texture2Dtest();
+ texture2Dtest();
+ passed();
+}
diff --git a/projects/hip/tests/src/texture/simpleTexture2DLayered.cpp b/projects/hip/tests/src/texture/simpleTexture2DLayered.cpp
new file mode 100644
index 0000000000..e3ba6d9afe
--- /dev/null
+++ b/projects/hip/tests/src/texture/simpleTexture2DLayered.cpp
@@ -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 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 *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();
+}
+
diff --git a/projects/hip/tests/src/texture/simpleTexture3D.cpp b/projects/hip/tests/src/texture/simpleTexture3D.cpp
new file mode 100644
index 0000000000..741b1671af
--- /dev/null
+++ b/projects/hip/tests/src/texture/simpleTexture3D.cpp
@@ -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 texf;
+texture texi;
+texture texc;
+
+template
+__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::value)
+ outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i);
+ else if(std::is_same::value)
+ outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i);
+ else if(std::is_same::value)
+ outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i);
+ }
+ }
+ }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+//! Run a simple test for tex3D
+////////////////////////////////////////////////////////////////////////////////
+template
+void runTest(int width,int height,int depth,texture *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(i,i,i,&texf);
+ runTest(i+1,i,i,&texi);
+ runTest(i,i+1,i,&texc);
+ }
+ passed();
+}
+