Merge 'master' into 'amd-master'
Change-Id: Ibb1dbe78aa52c8aeb77851ee1471e1edaec4c654
[ROCm/hip commit: cbe6b65e2c]
Этот коммит содержится в:
Jenkins
@@ -55,6 +55,14 @@ add_to_config(_versionInfo HIP_VERSION_MAJOR)
|
||||
add_to_config(_versionInfo HIP_VERSION_MINOR)
|
||||
add_to_config(_versionInfo HIP_VERSION_PATCH)
|
||||
|
||||
if(CMAKE_CXX_COMPILER MATCHES ".*hcc")
|
||||
set(HIP_COMPILER "hcc" CACHE STRING "HIP Compiler")
|
||||
set(HIP_PLATFORM "hcc" CACHE STRING "HIP Platform")
|
||||
get_filename_component(CXX_PATH ${CMAKE_CXX_COMPILER} DIRECTORY)
|
||||
get_filename_component(CXX_PATH ${CXX_PATH} DIRECTORY)
|
||||
set(HCC_HOME "${CXX_PATH}" CACHE PATH "Path to which HCC has been installed")
|
||||
endif()
|
||||
|
||||
# overwrite HIP_VERSION_PATCH for packaging
|
||||
if(DEFINED ENV{ROCM_BUILD_ID})
|
||||
set(HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}.${HIP_VERSION_GITCOUNT}-$ENV{ROCM_BUILD_ID}-${HIP_VERSION_GITHASH})
|
||||
|
||||
+1506
-282
Разница между файлами не показана из-за своего большого размера
Загрузить разницу
@@ -11,6 +11,7 @@ and provides practical suggestions on how to port CUDA code and work through com
|
||||
* [General Tips](#general-tips)
|
||||
* [Scanning existing CUDA code to scope the porting effort](#scanning-existing-cuda-code-to-scope-the-porting-effort)
|
||||
* [Converting a project "in-place"](#converting-a-project-in-place)
|
||||
* [CUDA to HIP Math Library Equivalents](#library-equivalents)
|
||||
- [Distinguishing Compiler Modes](#distinguishing-compiler-modes)
|
||||
* [Identifying HIP Target Platform](#identifying-hip-target-platform)
|
||||
* [Identifying the Compiler: hcc, hip-clang, or nvcc](#identifying-the-compiler-hcc-hip-clang-or-nvcc)
|
||||
@@ -46,7 +47,6 @@ and provides practical suggestions on how to port CUDA code and work through com
|
||||
+ [/usr/include/c++/v1/memory:5172:15: error: call to implicitly deleted default constructor of 'std::__1::bad_weak_ptr' throw bad_weak_ptr();](#usrincludecv1memory517215-error-call-to-implicitly-deleted-default-constructor-of-std__1bad_weak_ptr-throw-bad_weak_ptr)
|
||||
* [HIP Environment Variables](#hip-environment-variables)
|
||||
* [Editor Highlighting](#editor-highlighting)
|
||||
* [CUDA to HIP Math Library Equivalents](#library-equivalents)
|
||||
|
||||
|
||||
<!-- tocstop -->
|
||||
@@ -124,7 +124,21 @@ directory names.
|
||||
> hipconvertinplace.sh MY_SRC_DIR
|
||||
```
|
||||
|
||||
### Library Equivalents
|
||||
|
||||
| CUDA Library | ROCm Library | Comment |
|
||||
|------- | --------- | ----- |
|
||||
| cuBLAS | rocBLAS | Basic Linear Algebra Subroutines
|
||||
| cuFFT | rocFFT | Fast Fourier Transfer Library
|
||||
| cuSPARSE | rocSPARSE | Sparse BLAS + SPMV
|
||||
| cuSolver | rocSolver | Lapack library
|
||||
| AMG-X | rocALUTION | Sparse iterative solvers and preconditioners with Geometric and Algebraic MultiGrid
|
||||
| Thrust | hipThrust | C++ parallel algorithms library
|
||||
| CUB | rocPRIM | Low Level Optimized Parallel Primitives
|
||||
| cuDNN | MIOpen | Deep learning Solver Library
|
||||
| cuRAND | rocRAND | Random Number Generator Library
|
||||
| EIGEN | EIGEN – HIP port | C++ template library for linear algebra: matrices, vectors, numerical solvers,
|
||||
| NCCL | RCCL | Communications Primitives Library based on the MPI equivalents
|
||||
|
||||
|
||||
|
||||
@@ -559,18 +573,3 @@ HIP_VISIBLE_DEVICES = 0 : Only devices whose index is present in the
|
||||
See the utils/vim or utils/gedit directories to add handy highlighting to hip files.
|
||||
|
||||
|
||||
### Library Equivalents
|
||||
|
||||
| CUDA Library | ROCm Library | Comment |
|
||||
|------- | --------- | ----- |
|
||||
| cuBLAS | rocBLAS | Basic Linear Algebra Subroutines
|
||||
| cuFFT | rocFFT | Fast Fourier Transfer Library
|
||||
| cuSPARSE | rocSPARSE | Sparse BLAS + SPMV
|
||||
| cuSolver | rocSolver | Lapack library
|
||||
| AMG-X | rocALUTION | Sparse iterative solvers and preconditioners with Geometric and Algebraic MultiGrid
|
||||
| Thrust | hipThrust | C++ parallel algorithms library
|
||||
| CUB | rocPRIM | Low Level Optimized Parallel Primitives
|
||||
| cuDNN | MIOpen | Deep learning Solver Library
|
||||
| cuRAND | rocRAND | Random Number Generator Library
|
||||
| EIGEN | EIGEN – HIP port | C++ template library for linear algebra: matrices, vectors, numerical solvers,
|
||||
| NCCL | RCCL | Communications Primitives Library based on the MPI equivalents
|
||||
|
||||
@@ -1,14 +1,7 @@
|
||||
if (CUDA_VERSION VERSION_GREATER "9.2")
|
||||
cmake_minimum_required(VERSION 3.12.3)
|
||||
else()
|
||||
if (CUDA_VERSION VERSION_LESS "9.0")
|
||||
cmake_minimum_required(VERSION 3.5.0)
|
||||
else()
|
||||
cmake_minimum_required(VERSION 3.6.0)
|
||||
endif()
|
||||
if (MSVC AND MSVC_VERSION VERSION_GREATER "1900")
|
||||
cmake_minimum_required(VERSION 3.7.2)
|
||||
endif()
|
||||
cmake_minimum_required(VERSION 3.7.2)
|
||||
endif()
|
||||
|
||||
project(hipify-clang)
|
||||
|
||||
@@ -143,9 +143,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.0-v7.6.2.24`
|
||||
- Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.3.30`
|
||||
|
||||
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.2.24`
|
||||
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.3.30`
|
||||
|
||||
5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
|
||||
|
||||
@@ -179,9 +179,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.2.24
|
||||
Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.3.30
|
||||
|
||||
Ubuntu 16-18: LLVM 8.0.0 - 8.0.1, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.6.2.24
|
||||
Ubuntu 16-18: LLVM 8.0.0 - 8.0.1, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.6.3.30
|
||||
|
||||
Build system for the above configurations:
|
||||
|
||||
@@ -196,7 +196,7 @@ cmake
|
||||
-DCMAKE_INSTALL_PREFIX=../dist \
|
||||
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/8.0.1/dist \
|
||||
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0 \
|
||||
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.2.24 \
|
||||
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.3.30 \
|
||||
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.1/build/bin/llvm-lit \
|
||||
..
|
||||
```
|
||||
@@ -248,67 +248,71 @@ Running HIPify regression tests
|
||||
CUDA 10.0 - will be used for testing
|
||||
LLVM 8.0.1 - will be used for testing
|
||||
x86_64 - Platform architecture
|
||||
Linux 5.2.0-rc1-kfd-compute-roc-master-int-hipclang-623 - Platform OS
|
||||
Linux 5.2.0 - Platform OS
|
||||
64 - hipify-clang binary bitness
|
||||
64 - python 2.7.12 binary bitness
|
||||
========================================
|
||||
-- Testing: 54 tests, 12 threads --
|
||||
PASS: hipify :: unit_tests/headers/headers_test_03.cu (1 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_02.cu (2 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_01.cu (3 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_05.cu (4 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_11.cu (5 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_10.cu (6 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_07.cu (7 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_06.cu (8 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (11 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (12 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (13 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (14 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (16 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (17 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (18 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (19 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (20 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (21 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (22 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (23 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (24 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (25 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (27 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (28 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (29 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (30 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (31 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (32 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (33 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (34 of 54)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_09.cu (35 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (36 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (37 of 54)
|
||||
PASS: hipify :: unit_tests/samples/allocators.cu (38 of 54)
|
||||
PASS: hipify :: unit_tests/samples/coalescing.cu (39 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (40 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (41 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (42 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (43 of 54)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (44 of 54)
|
||||
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (45 of 54)
|
||||
PASS: hipify :: unit_tests/samples/axpy.cu (46 of 54)
|
||||
PASS: hipify :: unit_tests/samples/intro.cu (47 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (48 of 54)
|
||||
PASS: hipify :: unit_tests/samples/cudaRegister.cu (49 of 54)
|
||||
PASS: hipify :: unit_tests/samples/vec_add.cu (50 of 54)
|
||||
PASS: hipify :: unit_tests/samples/square.cu (51 of 54)
|
||||
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (52 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (53 of 54)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (54 of 54)
|
||||
Testing Time: 2.92s
|
||||
Expected Passes : 54
|
||||
-- Testing: 58 tests, 12 threads --
|
||||
PASS: hipify :: unit_tests/headers/headers_test_03.cu (1 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_02.cu (2 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_10.cu (3 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_05.cu (4 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_01.cu (5 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_11.cu (6 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_06.cu (7 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (12 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (13 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (14 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (15 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (16 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (17 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (18 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (19 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (20 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (21 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (22 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (23 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (24 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (25 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (27 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (28 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (29 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (30 of 58)
|
||||
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (31 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (32 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (33 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (34 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (35 of 58)
|
||||
PASS: hipify :: unit_tests/headers/headers_test_09.cu (36 of 58)
|
||||
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (37 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (38 of 58)
|
||||
PASS: hipify :: unit_tests/samples/MallocManaged.cpp (39 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (40 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (41 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (42 of 58)
|
||||
PASS: hipify :: unit_tests/samples/allocators.cu (43 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (44 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (45 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (46 of 58)
|
||||
PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (47 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (48 of 58)
|
||||
PASS: hipify :: unit_tests/samples/coalescing.cu (49 of 58)
|
||||
PASS: hipify :: unit_tests/samples/square.cu (50 of 58)
|
||||
PASS: hipify :: unit_tests/samples/vec_add.cu (51 of 58)
|
||||
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (52 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (53 of 58)
|
||||
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (54 of 58)
|
||||
PASS: hipify :: unit_tests/samples/intro.cu (55 of 58)
|
||||
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (56 of 58)
|
||||
PASS: hipify :: unit_tests/samples/axpy.cu (57 of 58)
|
||||
PASS: hipify :: unit_tests/samples/cudaRegister.cu (58 of 58)
|
||||
Testing Time: 2.81s
|
||||
Expected Passes : 58
|
||||
[100%] Built target test-hipify
|
||||
```
|
||||
### <a name="windows"></a >Windows
|
||||
@@ -317,13 +321,13 @@ On Windows 10 the following configurations are tested:
|
||||
|
||||
LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-5.1.10 - cudnn-7.1.4.18
|
||||
|
||||
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.2.24
|
||||
LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.3.30
|
||||
|
||||
LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.2.24
|
||||
LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.3.30
|
||||
|
||||
Build system for the above configurations:
|
||||
|
||||
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.0).
|
||||
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.3).
|
||||
|
||||
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
|
||||
|
||||
@@ -336,7 +340,7 @@ cmake
|
||||
-DCMAKE_PREFIX_PATH=f:/LLVM/6.0.1/dist \
|
||||
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \
|
||||
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \
|
||||
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.2.24 \
|
||||
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.3.30 \
|
||||
-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py \
|
||||
-Thost=x64
|
||||
..
|
||||
@@ -347,7 +351,7 @@ cmake
|
||||
-- - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm
|
||||
-- - Include path : F:/LLVM/6.0.1/dist/include
|
||||
-- - Binary path : F:/LLVM/6.0.1/dist/bin
|
||||
-- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.7.4", minimum required is "3.6")
|
||||
-- Found PythonInterp: C:/Program Files/Python37/python.exe (found suitable version "3.7.4", minimum required is "3.6")
|
||||
-- Found lit: C:/Program Files/Python36/Scripts/lit.exe
|
||||
-- Found FileCheck: F:/LLVM/6.0.1/dist/bin/FileCheck.exe
|
||||
-- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0")
|
||||
|
||||
@@ -65,6 +65,8 @@ extern const std::map<llvm::StringRef, hipCounter> CUDA_SPARSE_FUNCTION_MAP;
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_CAFFE2_TYPE_NAME_MAP;
|
||||
// Maps the names of CUDA CAFFE2 API functions to the corresponding HIP functions
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_CAFFE2_FUNCTION_MAP;
|
||||
// Maps the names of CUDA Device functions to the corresponding HIP functions
|
||||
extern const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP;
|
||||
|
||||
/**
|
||||
* The union of all the above maps, except includes.
|
||||
|
||||
@@ -0,0 +1,48 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include "CUDA2HIP.h"
|
||||
|
||||
// Maps CUDA header names to HIP header names
|
||||
const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP{
|
||||
{"umin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"llmin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"ullmin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"umax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"llmax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"ullmax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__isinff", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__isnanf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__finite", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__finitef", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__signbit", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__isnan", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__isinf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__signbitf", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__signbitl", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__finitel", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__isinfl", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__isnanl", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"_ldsign", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"_fdsign", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"_Pow_int", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
};
|
||||
@@ -67,7 +67,6 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
|
||||
* Otherwise, the source file is updated with the corresponding hipification.
|
||||
*/
|
||||
void HipifyAction::RewriteToken(const clang::Token& t) {
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
// String literals containing CUDA references need fixing.
|
||||
if (t.is(clang::tok::string_literal)) {
|
||||
StringRef s(t.getLiteralData(), t.getLength());
|
||||
@@ -78,13 +77,19 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
|
||||
return;
|
||||
}
|
||||
StringRef name = t.getRawIdentifier();
|
||||
const auto found = CUDA_RENAMES_MAP().find(name);
|
||||
if (found == CUDA_RENAMES_MAP().end()) {
|
||||
clang::SourceLocation sl = t.getLocation();
|
||||
FindAndReplace(name, sl, CUDA_RENAMES_MAP());
|
||||
}
|
||||
|
||||
void HipifyAction::FindAndReplace(llvm::StringRef name,
|
||||
clang::SourceLocation sl,
|
||||
const std::map<llvm::StringRef, hipCounter>& repMap) {
|
||||
const auto found = repMap.find(name);
|
||||
if (found == repMap.end()) {
|
||||
// So it's an identifier, but not CUDA? Boring.
|
||||
return;
|
||||
}
|
||||
Statistics::current().incrementCounter(found->second, name.str());
|
||||
clang::SourceLocation sl = t.getLocation();
|
||||
clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
|
||||
// Warn the user about unsupported identifier.
|
||||
if (Statistics::isUnsupported(found->second)) {
|
||||
@@ -96,6 +101,7 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
|
||||
return;
|
||||
}
|
||||
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
|
||||
clang::SourceManager& SM = getCompilerInstance().getSourceManager();
|
||||
ct::Replacement Rep(SM, sl, name.size(), repName.str());
|
||||
clang::FullSourceLoc fullSL(sl, SM);
|
||||
insertReplacement(Rep, fullSL);
|
||||
@@ -372,6 +378,14 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
|
||||
return true;
|
||||
}
|
||||
|
||||
bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaDeviceFuncCall")) {
|
||||
const clang::FunctionDecl *funcDcl = call->getDirectCallee();
|
||||
FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP);
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) {
|
||||
llcompat::insertReplacement(*replacements, rep);
|
||||
if (PrintStats) {
|
||||
@@ -395,7 +409,22 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
|
||||
).bind("cudaSharedIncompleteArrayVar"),
|
||||
this
|
||||
);
|
||||
// Ownership is transferred to the caller...
|
||||
Finder->addMatcher(
|
||||
mat::callExpr(
|
||||
mat::isExpansionInMainFile(),
|
||||
mat::callee(
|
||||
mat::functionDecl(
|
||||
mat::anyOf(
|
||||
mat::hasAttr(clang::attr::CUDADevice),
|
||||
mat::hasAttr(clang::attr::CUDAGlobal)
|
||||
),
|
||||
mat::unless(mat::hasAttr(clang::attr::CUDAHost))
|
||||
)
|
||||
)
|
||||
).bind("cudaDeviceFuncCall"),
|
||||
this
|
||||
);
|
||||
// Ownership is transferred to the caller.
|
||||
return Finder->newASTConsumer();
|
||||
}
|
||||
|
||||
@@ -517,4 +546,5 @@ void HipifyAction::ExecuteAction() {
|
||||
void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
|
||||
if (cudaLaunchKernel(Result)) return;
|
||||
if (cudaSharedIncompleteArrayVar(Result)) return;
|
||||
if (cudaDeviceFuncCall(Result)) return;
|
||||
}
|
||||
|
||||
@@ -70,6 +70,7 @@ public:
|
||||
bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
|
||||
// Called by the preprocessor for each include directive during the non-raw lexing pass.
|
||||
void InclusionDirective(clang::SourceLocation hash_loc,
|
||||
const clang::Token &include_token,
|
||||
@@ -99,4 +100,5 @@ protected:
|
||||
void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override;
|
||||
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
|
||||
bool Exclude(const hipCounter & hipToken);
|
||||
void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap);
|
||||
};
|
||||
|
||||
@@ -43,7 +43,13 @@ public:
|
||||
ct::FrontendActionFactory(),
|
||||
replacements(r) {}
|
||||
|
||||
#if LLVM_VERSION_MAJOR < 10
|
||||
clang::FrontendAction* create() override {
|
||||
return new T(replacements);
|
||||
}
|
||||
#else
|
||||
std::unique_ptr <clang::FrontendAction> create() override {
|
||||
return std::unique_ptr<clang::FrontendAction>(new T(replacements));
|
||||
}
|
||||
#endif
|
||||
};
|
||||
|
||||
@@ -57,6 +57,7 @@ const char *counterNames[NUM_CONV_TYPES] = {
|
||||
"complex", // CONV_COMPLEX
|
||||
"library", // CONV_LIB_FUNC
|
||||
"device_library", // CONV_LIB_DEVICE_FUNC
|
||||
"device_function", // CONV_DEVICE_FUNC
|
||||
"include", // CONV_INCLUDE
|
||||
"include_cuda_main_header", // CONV_INCLUDE_CUDA_MAIN_H
|
||||
"type", // CONV_TYPE
|
||||
|
||||
@@ -112,6 +112,7 @@ enum ConvTypes {
|
||||
CONV_COMPLEX,
|
||||
CONV_LIB_FUNC,
|
||||
CONV_LIB_DEVICE_FUNC,
|
||||
CONV_DEVICE_FUNC,
|
||||
CONV_INCLUDE,
|
||||
CONV_INCLUDE_CUDA_MAIN_H,
|
||||
CONV_TYPE,
|
||||
|
||||
@@ -970,7 +970,7 @@ static void __barrier(int n)
|
||||
|
||||
__device__
|
||||
inline
|
||||
__attribute__((noduplicate))
|
||||
__attribute__((convergent))
|
||||
void __syncthreads()
|
||||
{
|
||||
__barrier(__CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
@@ -331,11 +331,6 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
|
||||
|
||||
typedef int hipLaunchParm;
|
||||
|
||||
#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
|
||||
do { \
|
||||
kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(hipLaunchParm{}, ##__VA_ARGS__); \
|
||||
} while (0)
|
||||
|
||||
#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
|
||||
do { \
|
||||
kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \
|
||||
@@ -452,6 +447,7 @@ hc_get_workitem_absolute_id(int dim)
|
||||
#endif
|
||||
|
||||
// Support std::complex.
|
||||
#ifndef _OPENMP
|
||||
#pragma push_macro("__CUDA__")
|
||||
#define __CUDA__
|
||||
#include <__clang_cuda_math_forward_declares.h>
|
||||
@@ -461,6 +457,7 @@ hc_get_workitem_absolute_id(int dim)
|
||||
#include <cuda_wrappers/new>
|
||||
#undef __CUDA__
|
||||
#pragma pop_macro("__CUDA__")
|
||||
#endif // ndef _OPENMP
|
||||
|
||||
#if __HIP_VDI__
|
||||
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
|
||||
|
||||
@@ -220,8 +220,17 @@ TidInfo::TidInfo() : _apiSeqNum(0) {
|
||||
tid_ss_num << std::this_thread::get_id();
|
||||
tid_ss << std::hex << std::stoull(tid_ss_num.str());
|
||||
|
||||
tprintf(DB_API, "HIP initialized short_tid#%d (maps to full_tid: 0x%s)\n", _shortTid,
|
||||
tid_ss.str().c_str());
|
||||
// cannot use tprintf here since it will recurse back into TlsData constructor
|
||||
#if COMPILE_HIP_DB
|
||||
if (HIP_DB & (1 << DB_API)) {
|
||||
char msgStr[1000];
|
||||
snprintf(msgStr, sizeof(msgStr),
|
||||
"HIP initialized short_tid#%d (maps to full_tid: 0x%s)\n",
|
||||
tid(), tid_ss.str().c_str());
|
||||
fprintf(stderr, " %ship-%s pid:%d tid:%d:%s%s", dbName[DB_API]._color,
|
||||
dbName[DB_API]._shortName, pid(), tid(), msgStr, KNRM);
|
||||
}
|
||||
#endif
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
@@ -1557,13 +1557,6 @@ hipError_t ihipMemPtrGetInfo(void* ptr, size_t* size) {
|
||||
|
||||
template <typename T>
|
||||
void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t count) {
|
||||
// Just Use count, instead of dividing by 4, the calling API already does it
|
||||
if (sizeof(T) == sizeof(uint32_t) && (count % sizeof(uint32_t) == 0) &&
|
||||
!hsa_amd_memory_fill(ptr, reinterpret_cast<const std::uint32_t&>(val), count)) {
|
||||
// Only return if the execution completes without error
|
||||
// if error occured, try the normal version
|
||||
return;
|
||||
}
|
||||
static constexpr uint32_t block_dim = 256;
|
||||
|
||||
const uint32_t grid_dim = clamp_integer<size_t>(count / block_dim, 1, UINT32_MAX);
|
||||
|
||||
@@ -0,0 +1,46 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
|
||||
// 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)
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <algorithm>
|
||||
|
||||
namespace my {
|
||||
unsigned int umin(unsigned int arg1, unsigned int arg2) {
|
||||
return (arg1 < arg2) ? arg1 : arg2;
|
||||
}
|
||||
unsigned int umax(unsigned int arg1, unsigned int arg2) {
|
||||
return (arg1 > arg2) ? arg1 : arg2;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsigned int u2)
|
||||
{
|
||||
unsigned int _umin = umin(u1, u2);
|
||||
unsigned int _umax = umax(u1, u2);
|
||||
int i = threadIdx.x;
|
||||
A[i] = i + _umin;
|
||||
B[i] = i + _umax;
|
||||
C[i] = A[i] + B[i];
|
||||
}
|
||||
|
||||
int main() {
|
||||
unsigned int u1 = 33;
|
||||
unsigned int u2 = 34;
|
||||
unsigned int _min = my::umin(u1, u2);
|
||||
unsigned int _max = my::umax(u1, u2);
|
||||
float *A, *B, *C;
|
||||
// CHECK: hipMalloc((void**)&A, SIZE);
|
||||
cudaMalloc((void**)&A, SIZE);
|
||||
// CHECK: hipMalloc((void**)&B, SIZE);
|
||||
cudaMalloc((void**)&B, SIZE);
|
||||
// CHECK: hipMalloc((void**)&C, SIZE);
|
||||
cudaMalloc((void**)&C, SIZE);
|
||||
dim3 dimGrid(LEN / 512, 1, 1);
|
||||
dim3 dimBlock(512, 1, 1);
|
||||
// CHECK: hipLaunchKernelGGL(uint_arithm, dim3(dimGrid), dim3(dimBlock), 0, 0, A, B, C, u1, u2);
|
||||
uint_arithm<<<dimGrid, dimBlock>>>(A, B, C, u1, u2);
|
||||
return _min < _max;
|
||||
}
|
||||
+198
@@ -0,0 +1,198 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
|
||||
/*
|
||||
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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// CHECK: #include "hip/hip_runtime.h"
|
||||
#include "cuda_runtime.h"
|
||||
#include <iostream>
|
||||
#define NUM 1000000
|
||||
|
||||
// CHECK: if (status != hipSuccess) {
|
||||
#define CUDA_CHECK(status) \
|
||||
if (status != cudaSuccess) { \
|
||||
std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \
|
||||
exit(0); \
|
||||
}
|
||||
|
||||
// Device (Kernel) function
|
||||
__global__ void multiply(float* C, float* A, float* B, int N) {
|
||||
int tx = blockDim.x*blockIdx.x+threadIdx.x;
|
||||
if (tx < N) {
|
||||
C[tx] = A[tx] * B[tx];
|
||||
}
|
||||
}
|
||||
|
||||
// CPU implementation
|
||||
void multiplyCPU(float* C, float* A, float* B, int N) {
|
||||
for(unsigned int i=0; i<N; i++) {
|
||||
C[i] = A[i] * B[i];
|
||||
}
|
||||
}
|
||||
|
||||
void launchKernel(float* C, float* A, float* B, bool manual) {
|
||||
// CHECK: hipDeviceProp_t devProp;
|
||||
cudaDeviceProp devProp;
|
||||
// CHECK: CUDA_CHECK(hipGetDeviceProperties(&devProp, 0));
|
||||
CUDA_CHECK(cudaGetDeviceProperties(&devProp, 0));
|
||||
|
||||
// CHECK: hipEvent_t start, stop;
|
||||
cudaEvent_t start, stop;
|
||||
// CHECK: CUDA_CHECK(hipEventCreate(&start));
|
||||
CUDA_CHECK(cudaEventCreate(&start));
|
||||
// CHECK: CUDA_CHECK(hipEventCreate(&stop));
|
||||
CUDA_CHECK(cudaEventCreate(&stop));
|
||||
float eventMs = 1.0f;
|
||||
const unsigned threadsperblock = 32;
|
||||
const unsigned blocks = (NUM/threadsperblock) + 1;
|
||||
|
||||
int mingridSize = 0;
|
||||
int gridSize = 0;
|
||||
int blockSize = 0;
|
||||
|
||||
if (manual) {
|
||||
blockSize = threadsperblock;
|
||||
gridSize = blocks;
|
||||
std::cout << std::endl << "Manual Configuration with block size " << blockSize << std::endl;
|
||||
} else {
|
||||
// CHECK: CUDA_CHECK(hipOccupancyMaxPotentialBlockSize(&mingridSize, &blockSize, multiply, 0, 0));
|
||||
CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(&mingridSize, &blockSize, multiply, 0, 0));
|
||||
std::cout << std::endl << "Automatic Configuation based on hipOccupancyMaxPotentialBlockSize " << std::endl;
|
||||
std::cout << "Suggested blocksize is " << blockSize << ", Minimum gridsize is " << mingridSize << std::endl;
|
||||
gridSize = (NUM/blockSize)+1;
|
||||
}
|
||||
|
||||
// Record the start event
|
||||
// CHECK: CUDA_CHECK(hipEventRecord(start, NULL));
|
||||
CUDA_CHECK(cudaEventRecord(start, NULL));
|
||||
|
||||
// Launching the Kernel from Host
|
||||
// CHECK: hipLaunchKernelGGL(multiply, dim3(gridSize), dim3(blockSize), 0, 0, C, A, B, NUM);
|
||||
multiply <<<gridSize , blockSize>>> (C, A, B, NUM);
|
||||
|
||||
// Record the stop event
|
||||
// CHECK: CUDA_CHECK(hipEventRecord(stop, NULL));
|
||||
CUDA_CHECK(cudaEventRecord(stop, NULL));
|
||||
// CHECK: CUDA_CHECK(hipEventSynchronize(stop));
|
||||
CUDA_CHECK(cudaEventSynchronize(stop));
|
||||
|
||||
// CHECK: CUDA_CHECK(hipEventElapsedTime(&eventMs, start, stop));
|
||||
CUDA_CHECK(cudaEventElapsedTime(&eventMs, start, stop));
|
||||
printf("kernel Execution time = %6.3fms\n", eventMs);
|
||||
|
||||
// Calculate Occupancy
|
||||
int numBlock = 0;
|
||||
// CHECK: CUDA_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0));
|
||||
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0));
|
||||
|
||||
if(devProp.maxThreadsPerMultiProcessor) {
|
||||
std::cout << "Theoretical Occupancy is " << (double)numBlock* blockSize/devProp.maxThreadsPerMultiProcessor * 100 << "%" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
float *A, *B, *C0, *C1, *cpuC;
|
||||
float *Ad, *Bd, *C0d, *C1d;
|
||||
int errors=0;
|
||||
|
||||
// Initialize the input data
|
||||
A = (float*)malloc(NUM * sizeof(float));
|
||||
B = (float*)malloc(NUM * sizeof(float));
|
||||
C0 = (float*)malloc(NUM * sizeof(float));
|
||||
C1 = (float*)malloc(NUM * sizeof(float));
|
||||
cpuC = (float*)malloc(NUM * sizeof(float));
|
||||
|
||||
for(int i=0; i< NUM; i++) {
|
||||
A[i] = i;
|
||||
B[i] = i;
|
||||
}
|
||||
|
||||
// Allocate the memory on the device side
|
||||
// CHECK: CUDA_CHECK(hipMalloc((void**)&Ad, NUM * sizeof(float)));
|
||||
CUDA_CHECK(cudaMalloc((void**)&Ad, NUM * sizeof(float)));
|
||||
// CHECK: CUDA_CHECK(hipMalloc((void**)&Bd, NUM * sizeof(float)));
|
||||
CUDA_CHECK(cudaMalloc((void**)&Bd, NUM * sizeof(float)));
|
||||
// CHECK: CUDA_CHECK(hipMalloc((void**)&C0d, NUM * sizeof(float)));
|
||||
CUDA_CHECK(cudaMalloc((void**)&C0d, NUM * sizeof(float)));
|
||||
// CHECK: CUDA_CHECK(hipMalloc((void**)&C1d, NUM * sizeof(float)));
|
||||
CUDA_CHECK(cudaMalloc((void**)&C1d, NUM * sizeof(float)));
|
||||
|
||||
// Memory transfer from host to device
|
||||
// CHECK: CUDA_CHECK(hipMemcpy(Ad,A,NUM * sizeof(float), hipMemcpyHostToDevice));
|
||||
CUDA_CHECK(cudaMemcpy(Ad,A,NUM * sizeof(float), cudaMemcpyHostToDevice));
|
||||
// CHECK: CUDA_CHECK(hipMemcpy(Bd,B,NUM * sizeof(float), hipMemcpyHostToDevice));
|
||||
CUDA_CHECK(cudaMemcpy(Bd,B,NUM * sizeof(float), cudaMemcpyHostToDevice));
|
||||
|
||||
// Kernel launch with manual/default block size
|
||||
launchKernel(C0d, Ad, Bd, 1);
|
||||
|
||||
// Kernel launch with the block size suggested by cudaOccupancyMaxPotentialBlockSize
|
||||
launchKernel(C1d, Ad, Bd, 0);
|
||||
|
||||
// Memory transfer from device to host
|
||||
// CHECK: CUDA_CHECK(hipMemcpy(C0,C0d, NUM * sizeof(float), hipMemcpyDeviceToHost));
|
||||
CUDA_CHECK(cudaMemcpy(C0,C0d, NUM * sizeof(float), cudaMemcpyDeviceToHost));
|
||||
// CHECK: CUDA_CHECK(hipMemcpy(C1,C1d, NUM * sizeof(float), hipMemcpyDeviceToHost));
|
||||
CUDA_CHECK(cudaMemcpy(C1,C1d, NUM * sizeof(float), cudaMemcpyDeviceToHost));
|
||||
|
||||
// CPU computation
|
||||
multiplyCPU(cpuC, A, B, NUM);
|
||||
|
||||
// Verify the results
|
||||
double eps = 1.0E-6;
|
||||
|
||||
for (int i = 0; i < NUM; i++) {
|
||||
if (std::abs(C0[i] - cpuC[i]) > eps) {
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
|
||||
if (errors != 0) {
|
||||
printf("\nManual Test FAILED: %d errors\n", errors);
|
||||
errors=0;
|
||||
} else {
|
||||
printf("\nManual Test PASSED!\n");
|
||||
}
|
||||
|
||||
for (int i = 0; i < NUM; i++) {
|
||||
if (std::abs(C1[i] - cpuC[i]) > eps) {
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
|
||||
if (errors != 0) {
|
||||
printf("\n Automatic Test FAILED: %d errors\n", errors);
|
||||
} else {
|
||||
printf("\nAutomatic Test PASSED!\n");
|
||||
}
|
||||
|
||||
// CHECK: CUDA_CHECK(hipFree(Ad));
|
||||
CUDA_CHECK(cudaFree(Ad));
|
||||
// CHECK: CUDA_CHECK(hipFree(Bd));
|
||||
CUDA_CHECK(cudaFree(Bd));
|
||||
// CHECK: CUDA_CHECK(hipFree(C0d));
|
||||
CUDA_CHECK(cudaFree(C0d));
|
||||
// CHECK: CUDA_CHECK(hipFree(C1d));
|
||||
CUDA_CHECK(cudaFree(C1d));
|
||||
|
||||
free(A);
|
||||
free(B);
|
||||
free(C0);
|
||||
free(C1);
|
||||
free(cpuC);
|
||||
}
|
||||
@@ -19,6 +19,8 @@ __global__ void axpy(T a, T *x, T *y) {
|
||||
y[threadIdx.x] = a * x[threadIdx.x];
|
||||
}
|
||||
|
||||
__global__ void empty() {
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
const int kDataLen = 4;
|
||||
@@ -64,6 +66,15 @@ int main(int argc, char* argv[]) {
|
||||
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
|
||||
empty<<<1, kDataLen>>> ( );
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
|
||||
empty<<<1, kDataLen, 0>>>();
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
|
||||
empty<<<1, kDataLen, 0, 0>>>();
|
||||
|
||||
// CHECK: COMPLETE_LAUNCH;
|
||||
COMPLETE_LAUNCH;
|
||||
|
||||
|
||||
@@ -1,713 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2018 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 HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
// Includes HIP Runtime
|
||||
#include "hip/hip_runtime.h"
|
||||
#include <test_common.h>
|
||||
|
||||
#include <set>
|
||||
#include <unordered_map>
|
||||
|
||||
#define test_failed(test_name) \
|
||||
printf("%s %s FAILED!%s\n", KRED, test_name, KNRM);
|
||||
#define test_passed(test_name) \
|
||||
printf("%s %s PASSED!%s\n", KGRN, test_name, KNRM);
|
||||
|
||||
/* BEGIN DUPLICATION */
|
||||
|
||||
// TODO: The host code below is duplicated, and should ideally reside
|
||||
// in a common host-side library.
|
||||
|
||||
typedef enum {
|
||||
__OCKL_AS_PACKET_EMPTY = 0,
|
||||
__OCKL_AS_PACKET_READY = 1
|
||||
} __ockl_as_packet_type_t;
|
||||
|
||||
#define __OCKL_AS_PAYLOAD_ALIGNMENT 4
|
||||
#define __OCKL_AS_PAYLOAD_BYTES 48
|
||||
|
||||
typedef enum {
|
||||
__OCKL_AS_PACKET_HEADER_TYPE = 0, // corresponds to HSA_PACKET_HEADER_TYPE
|
||||
__OCKL_AS_PACKET_HEADER_RESERVED0 = 8,
|
||||
__OCKL_AS_PACKET_HEADER_FLAGS = 13,
|
||||
__OCKL_AS_PACKET_HEADER_BYTES = 16,
|
||||
__OCKL_AS_PACKET_HEADER_SERVICE = 24,
|
||||
} __ockl_as_packet_header_t;
|
||||
|
||||
typedef enum {
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_TYPE = 8,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_RESERVED0 = 5,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_FLAGS = 3,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_BYTES = 8,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_SERVICE = 8
|
||||
} __ockl_as_packet_header_width_t;
|
||||
|
||||
// A packet is 64 bytes long, and the payload starts at index 16.
|
||||
struct __ockl_as_packet_t {
|
||||
uint header;
|
||||
uint reserved1;
|
||||
ulong connection_id;
|
||||
|
||||
uchar payload[__OCKL_AS_PAYLOAD_BYTES];
|
||||
};
|
||||
|
||||
typedef enum {
|
||||
__OCKL_AS_STATUS_SUCCESS,
|
||||
__OCKL_AS_STATUS_INVALID_REQUEST,
|
||||
__OCKL_AS_STATUS_OUT_OF_RESOURCES,
|
||||
__OCKL_AS_STATUS_BUSY,
|
||||
__OCKL_AS_STATUS_UNKNOWN_ERROR
|
||||
} __ockl_as_status_t;
|
||||
|
||||
typedef enum {
|
||||
__OCKL_AS_CONNECTION_BEGIN = 1,
|
||||
__OCKL_AS_CONNECTION_END = 2,
|
||||
} __ockl_as_flag_t;
|
||||
|
||||
typedef enum { __OCKL_AS_FEATURE_ASYNCHRONOUS = 1 } __ockl_as_feature_t;
|
||||
|
||||
typedef struct {
|
||||
// Opaque handle. The value 0 is reserved.
|
||||
ulong handle;
|
||||
} __ockl_as_signal_t;
|
||||
|
||||
typedef struct __ockl_as_packet_t __ockl_as_packet_t;
|
||||
|
||||
typedef struct {
|
||||
ulong read_index;
|
||||
ulong write_index;
|
||||
__ockl_as_signal_t doorbell_signal;
|
||||
__ockl_as_packet_t *base_address;
|
||||
ulong size;
|
||||
} __ockl_as_stream_t;
|
||||
|
||||
#define ATTR_GLOBAL __attribute__((address_space(1)))
|
||||
|
||||
extern "C" __device__ __ockl_as_status_t
|
||||
__ockl_as_write_block(__ockl_as_stream_t ATTR_GLOBAL *stream, uchar service_id,
|
||||
uint64_t *connection_id, const uchar *str, uint32_t len,
|
||||
uchar flags);
|
||||
|
||||
__device__ __ockl_as_status_t
|
||||
__hip_as_write_block(__ockl_as_stream_t *stream, uchar service_id,
|
||||
uint64_t *connection_id, const uchar *str, uint32_t len,
|
||||
uchar flags)
|
||||
{
|
||||
__ockl_as_stream_t ATTR_GLOBAL *gstream =
|
||||
(__ockl_as_stream_t ATTR_GLOBAL *)(stream);
|
||||
return __ockl_as_write_block(gstream, service_id, connection_id, str, len,
|
||||
flags);
|
||||
}
|
||||
|
||||
/* END DUPLICATION */
|
||||
|
||||
static __ockl_as_stream_t *
|
||||
createStream(void *ptr, uint buffer_size, uint num_packets)
|
||||
{
|
||||
memset(ptr, 0, buffer_size);
|
||||
__ockl_as_stream_t *r = (__ockl_as_stream_t *)ptr;
|
||||
r->base_address = (__ockl_as_packet_t *)(&r[1]);
|
||||
r->doorbell_signal = {0};
|
||||
r->size = num_packets;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
static uint8_t
|
||||
get_header_field(uint32_t header, uint8_t offset, uint8_t size)
|
||||
{
|
||||
return (header >> offset) & ((1 << size) - 1);
|
||||
}
|
||||
|
||||
static uint8_t
|
||||
get_packet_type(uint32_t header)
|
||||
{
|
||||
return get_header_field(header, __OCKL_AS_PACKET_HEADER_TYPE,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_TYPE);
|
||||
}
|
||||
|
||||
static uint8_t
|
||||
get_packet_flags(uint32_t header)
|
||||
{
|
||||
return get_header_field(header, __OCKL_AS_PACKET_HEADER_FLAGS,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_FLAGS);
|
||||
}
|
||||
|
||||
static uint8_t
|
||||
get_packet_bytes(uint32_t header)
|
||||
{
|
||||
return get_header_field(header, __OCKL_AS_PACKET_HEADER_BYTES,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_BYTES);
|
||||
}
|
||||
|
||||
static uint8_t
|
||||
get_packet_service(uint32_t header)
|
||||
{
|
||||
return get_header_field(header, __OCKL_AS_PACKET_HEADER_SERVICE,
|
||||
__OCKL_AS_PACKET_HEADER_WIDTH_SERVICE);
|
||||
}
|
||||
|
||||
const unsigned int __OCKL_AS_PACKET_SIZE = sizeof(__ockl_as_packet_t);
|
||||
|
||||
/* END DUPLICATION */
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define STR_HELLO_WORLD "hello world"
|
||||
#define STRLEN_HELLO_WORLD 11
|
||||
|
||||
const unsigned int THREADS_PER_BLOCK = 123; // include a partial warp
|
||||
const unsigned int NUM_BLOCKS = 3; // because powers of two are too convenient
|
||||
const unsigned int NUM_THREADS = NUM_BLOCKS * THREADS_PER_BLOCK;
|
||||
const unsigned int NUM_PACKETS_INSUFFICIENT = NUM_THREADS - 23;
|
||||
const unsigned int NUM_PACKETS_LARGE = NUM_THREADS * 4;
|
||||
const unsigned int NUM_SERVICES = 7;
|
||||
const unsigned int TEST_SERVICE = 42;
|
||||
|
||||
unsigned int
|
||||
read_uint(const unsigned char *ptr)
|
||||
{
|
||||
unsigned int value = 0;
|
||||
|
||||
for (int ii = sizeof(unsigned int) - 1; ii >= 0; --ii) {
|
||||
value <<= 8;
|
||||
value |= ptr[ii];
|
||||
}
|
||||
|
||||
return value;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
singlePacketSingleProducer(__ockl_as_stream_t *stream)
|
||||
{
|
||||
uint len = STRLEN_HELLO_WORLD;
|
||||
|
||||
uint64_t connection_id;
|
||||
|
||||
__hip_as_write_block(stream, TEST_SERVICE, &connection_id,
|
||||
(const uint8_t *)STR_HELLO_WORLD, STRLEN_HELLO_WORLD,
|
||||
__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END);
|
||||
}
|
||||
|
||||
bool
|
||||
checkSinglePacketSingleProducer(__ockl_as_stream_t *stream)
|
||||
{
|
||||
if (stream->write_index != 1)
|
||||
return false;
|
||||
|
||||
__ockl_as_packet_t *packet = &stream->base_address[0];
|
||||
uint header = packet->header;
|
||||
if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
|
||||
return false;
|
||||
|
||||
if (get_packet_service(header) != TEST_SERVICE)
|
||||
return false;
|
||||
|
||||
if (get_packet_bytes(header) != STRLEN_HELLO_WORLD)
|
||||
return false;
|
||||
|
||||
if (get_packet_flags(header) !=
|
||||
(__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END))
|
||||
return false;
|
||||
|
||||
if (0 != strcmp(STR_HELLO_WORLD, (const char *)packet->payload))
|
||||
return false;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool
|
||||
ocklAsSinglePacketSingleProducer()
|
||||
{
|
||||
bool success = true;
|
||||
unsigned int numThreads = 1;
|
||||
unsigned int numBlocks = 1;
|
||||
|
||||
unsigned int numPackets = 1;
|
||||
unsigned int bufferSize =
|
||||
sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
|
||||
|
||||
void *buffer;
|
||||
HIPCHECK(hipHostMalloc(&buffer, bufferSize));
|
||||
|
||||
__ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
|
||||
|
||||
hipLaunchKernelGGL(singlePacketSingleProducer, dim3(numBlocks),
|
||||
dim3(numThreads), 0, 0, stream);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
if (!checkSinglePacketSingleProducer(stream)) {
|
||||
test_failed(__func__);
|
||||
success = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostFree(buffer));
|
||||
return success;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
multipleProducers(__ockl_as_stream_t *stream)
|
||||
{
|
||||
const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
unsigned char data = (unsigned char)tid;
|
||||
|
||||
uint64_t connection_id;
|
||||
|
||||
__hip_as_write_block(stream, tid % NUM_SERVICES, &connection_id,
|
||||
(const unsigned char *)&tid, sizeof(unsigned int),
|
||||
__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END);
|
||||
}
|
||||
|
||||
bool
|
||||
checkMultipleProducers(__ockl_as_stream_t *stream)
|
||||
{
|
||||
int data[NUM_SERVICES] = {
|
||||
0,
|
||||
};
|
||||
|
||||
if (stream->write_index != NUM_THREADS)
|
||||
return false;
|
||||
|
||||
for (int ii = 0; ii != NUM_THREADS; ++ii) {
|
||||
__ockl_as_packet_t *packet = &stream->base_address[ii];
|
||||
uint header = packet->header;
|
||||
if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
|
||||
return false;
|
||||
|
||||
if (get_packet_bytes(header) != sizeof(unsigned int))
|
||||
return false;
|
||||
|
||||
if (get_packet_flags(header) !=
|
||||
(__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END))
|
||||
return false;
|
||||
|
||||
unsigned char service = get_packet_service(header);
|
||||
unsigned int payload = read_uint(packet->payload);
|
||||
|
||||
if (service != payload % NUM_SERVICES)
|
||||
return false;
|
||||
data[service]++;
|
||||
}
|
||||
|
||||
int expected[NUM_SERVICES];
|
||||
for (int ii = 0; ii != NUM_SERVICES; ++ii) {
|
||||
expected[ii] = NUM_THREADS / NUM_SERVICES;
|
||||
if (ii < NUM_THREADS % NUM_SERVICES) {
|
||||
expected[ii]++;
|
||||
}
|
||||
}
|
||||
|
||||
for (int ii = 0; ii != NUM_SERVICES; ++ii) {
|
||||
if (data[ii] != expected[ii])
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool
|
||||
ocklAsMultipleProducers()
|
||||
{
|
||||
bool success = true;
|
||||
unsigned int numPackets = NUM_THREADS;
|
||||
unsigned int bufferSize =
|
||||
sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
|
||||
|
||||
void *buffer;
|
||||
HIPCHECK(hipHostMalloc(&buffer, bufferSize));
|
||||
|
||||
__ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
|
||||
|
||||
hipLaunchKernelGGL(multipleProducers, dim3(NUM_BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, stream);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
if (!checkMultipleProducers(stream)) {
|
||||
test_failed(__func__);
|
||||
success = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostFree(buffer));
|
||||
return success;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
dropPackets(__ockl_as_stream_t *stream, unsigned int *status)
|
||||
{
|
||||
const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
|
||||
uint64_t connection_id;
|
||||
|
||||
status[tid] =
|
||||
__hip_as_write_block(stream, TEST_SERVICE, &connection_id,
|
||||
(const unsigned char *)&tid, sizeof(unsigned int),
|
||||
__OCKL_AS_CONNECTION_BEGIN |
|
||||
__OCKL_AS_CONNECTION_END);
|
||||
}
|
||||
|
||||
bool
|
||||
checkDropPackets(__ockl_as_stream_t *stream, unsigned int *status)
|
||||
{
|
||||
unsigned int errorsExpected = NUM_THREADS - NUM_PACKETS_INSUFFICIENT;
|
||||
for (int ii = 0; ii != NUM_THREADS; ++ii) {
|
||||
switch (status[ii]) {
|
||||
case __OCKL_AS_STATUS_OUT_OF_RESOURCES:
|
||||
if (errorsExpected == 0)
|
||||
return false;
|
||||
--errorsExpected;
|
||||
break;
|
||||
case __OCKL_AS_STATUS_SUCCESS:
|
||||
break;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
}
|
||||
if (errorsExpected != 0)
|
||||
return false;
|
||||
|
||||
if (stream->write_index != NUM_PACKETS_INSUFFICIENT)
|
||||
return false;
|
||||
|
||||
for (int ii = 0; ii != NUM_PACKETS_INSUFFICIENT; ++ii) {
|
||||
__ockl_as_packet_t *packet = &stream->base_address[ii];
|
||||
uint header = packet->header;
|
||||
if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
|
||||
return false;
|
||||
|
||||
if (get_packet_service(header) != TEST_SERVICE)
|
||||
return false;
|
||||
|
||||
if (get_packet_bytes(header) != sizeof(unsigned int))
|
||||
return false;
|
||||
|
||||
if (get_packet_flags(header) !=
|
||||
(__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END))
|
||||
return false;
|
||||
|
||||
unsigned int payload = read_uint(packet->payload);
|
||||
if (payload >= NUM_THREADS)
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool
|
||||
ocklAsDropPackets()
|
||||
{
|
||||
bool success = true;
|
||||
|
||||
unsigned int numPackets = NUM_PACKETS_INSUFFICIENT;
|
||||
unsigned int bufferSize =
|
||||
sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
|
||||
void *buffer;
|
||||
HIPCHECK(hipHostMalloc(&buffer, bufferSize));
|
||||
|
||||
void *status;
|
||||
HIPCHECK(hipHostMalloc(&status, NUM_THREADS * sizeof(unsigned int)));
|
||||
|
||||
__ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
|
||||
|
||||
hipEvent_t event;
|
||||
HIPCHECK(hipEventCreate(&event));
|
||||
hipLaunchKernelGGL(dropPackets, dim3(NUM_BLOCKS), dim3(THREADS_PER_BLOCK),
|
||||
0, 0, stream, (unsigned int *)status);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
if (!checkDropPackets(stream, (unsigned int *)status)) {
|
||||
test_failed(__func__);
|
||||
success = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostFree(buffer));
|
||||
return success;
|
||||
}
|
||||
|
||||
#define STR30 "Cras nec volutpat mi, sed sed."
|
||||
#define STR47 "Lorem ipsum dolor sit amet, consectetur nullam."
|
||||
#define STR60 "Curabitur id maximus nibh. Donec quis porttitor nisl nullam."
|
||||
#define STR95 \
|
||||
"In mollis imperdiet nibh nec ullamcorper." \
|
||||
" Suspendisse placerat massa iaculis ipsum viverra sed."
|
||||
#define STR124 \
|
||||
"Proin ut diam sit amet erat mollis gravida ac non sem." \
|
||||
" Mauris viverra leo metus, id luctus metus feugiat sed. Morbi " \
|
||||
"posuere."
|
||||
|
||||
#define DECLARE_TEST_DATA() \
|
||||
const char *str30 = STR30; \
|
||||
const char *str60 = STR60; \
|
||||
const char *str47 = STR47; \
|
||||
const char *str95 = STR95; \
|
||||
const char *str124 = STR124; \
|
||||
const int numStr = 5; \
|
||||
const char *strArray[5] = {str30, str60, str47, str95, str124}; \
|
||||
unsigned char strLengths[5] = {30, 60, 47, 95, 124};
|
||||
|
||||
__global__ void
|
||||
mixedProducers(__ockl_as_stream_t *stream)
|
||||
{
|
||||
DECLARE_TEST_DATA();
|
||||
|
||||
const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
const unsigned int idx = tid % 5;
|
||||
uint64_t connection_id;
|
||||
|
||||
__hip_as_write_block(stream, idx, &connection_id,
|
||||
(unsigned const char *)strArray[idx], strLengths[idx],
|
||||
__OCKL_AS_CONNECTION_BEGIN | __OCKL_AS_CONNECTION_END);
|
||||
}
|
||||
|
||||
bool
|
||||
checkMixedProducers(__ockl_as_stream_t *stream)
|
||||
{
|
||||
typedef std::unordered_map<uint64_t, std::string> stream_buffer_map_t;
|
||||
stream_buffer_map_t buffers;
|
||||
std::unordered_map<std::string, int> strRecd;
|
||||
|
||||
DECLARE_TEST_DATA();
|
||||
|
||||
for (unsigned long read_index = 0; read_index != stream->write_index;
|
||||
++read_index) {
|
||||
__ockl_as_packet_t *packet = stream->base_address + read_index;
|
||||
|
||||
uint header = packet->header;
|
||||
if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
|
||||
return false;
|
||||
|
||||
uint bytes = get_packet_bytes(header);
|
||||
unsigned char flags = get_packet_flags(header);
|
||||
unsigned char service = get_packet_service(header);
|
||||
unsigned long connection_id = packet->connection_id;
|
||||
unsigned char *payload = packet->payload;
|
||||
|
||||
if ((flags & __OCKL_AS_CONNECTION_BEGIN) !=
|
||||
(buffers.count(connection_id) == 0))
|
||||
return false;
|
||||
|
||||
std::string &buf = buffers[connection_id];
|
||||
buf.insert(buf.end(), payload, payload + bytes);
|
||||
|
||||
if (flags & __OCKL_AS_CONNECTION_END) {
|
||||
if (buf != strArray[service])
|
||||
return false;
|
||||
strRecd[buf]++;
|
||||
buffers.erase(connection_id);
|
||||
}
|
||||
}
|
||||
|
||||
int expected_counts[numStr];
|
||||
for (int ii = 0; ii != numStr; ++ii) {
|
||||
expected_counts[ii] = NUM_THREADS / numStr;
|
||||
if (ii < (NUM_THREADS % numStr)) {
|
||||
++expected_counts[ii];
|
||||
}
|
||||
}
|
||||
|
||||
if (strRecd.size() != numStr)
|
||||
return false;
|
||||
|
||||
for (int ii = 0; ii != numStr; ++ii) {
|
||||
std::string mystr(strArray[ii]);
|
||||
if (strRecd[mystr] != expected_counts[ii])
|
||||
return false;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
bool
|
||||
ocklAsMixedProducers()
|
||||
{
|
||||
bool success = true;
|
||||
|
||||
unsigned int numPackets = NUM_PACKETS_LARGE;
|
||||
unsigned int bufferSize =
|
||||
sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
|
||||
|
||||
void *buffer;
|
||||
HIPCHECK(hipHostMalloc(&buffer, bufferSize));
|
||||
|
||||
__ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
|
||||
|
||||
hipLaunchKernelGGL(mixedProducers, dim3(NUM_BLOCKS),
|
||||
dim3(THREADS_PER_BLOCK), 0, 0, stream);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
if (!checkMixedProducers(stream)) {
|
||||
test_failed(__func__);
|
||||
success = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostFree(buffer));
|
||||
return success;
|
||||
}
|
||||
|
||||
#define STR27 "In et consectetur mi metus."
|
||||
#define STR64 "Praesent tempus arcu id ligula blandit, eget congue justo metus."
|
||||
#define STR40 "Sed at dolor ipsum. Curabitur cras amet."
|
||||
|
||||
__global__ void
|
||||
splitMessage(__ockl_as_stream_t *stream)
|
||||
{
|
||||
const char *str27 = STR27;
|
||||
const char *str64 = STR64;
|
||||
const char *str40 = STR40;
|
||||
const int numStr = 3;
|
||||
const char *strArray[] = {str27, str64, str40};
|
||||
int strLengths[] = {27, 64, 40};
|
||||
|
||||
const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
int service = tid % 3;
|
||||
int first = tid % 3;
|
||||
int second = (tid + 1) % 3;
|
||||
int third = (tid + 2) % 3;
|
||||
|
||||
uint64_t connection_id;
|
||||
__hip_as_write_block(stream, service, &connection_id,
|
||||
(unsigned const char *)strArray[first],
|
||||
strLengths[first], __OCKL_AS_CONNECTION_BEGIN);
|
||||
__hip_as_write_block(stream, service, &connection_id,
|
||||
(unsigned const char *)strArray[second],
|
||||
strLengths[second], 0);
|
||||
__hip_as_write_block(stream, service, &connection_id,
|
||||
(unsigned const char *)strArray[third],
|
||||
strLengths[third], __OCKL_AS_CONNECTION_END);
|
||||
}
|
||||
|
||||
bool
|
||||
checkSplitMessage(__ockl_as_stream_t *stream)
|
||||
{
|
||||
typedef std::unordered_map<uint64_t, std::string> stream_buffer_map_t;
|
||||
stream_buffer_map_t buffers;
|
||||
std::unordered_map<std::string, int> strRecd;
|
||||
|
||||
static const int numExpected = 3;
|
||||
const char *strExpected[numExpected] = {STR27 STR64 STR40,
|
||||
STR64 STR40 STR27,
|
||||
STR40 STR27 STR64};
|
||||
|
||||
for (unsigned long read_index = 0; read_index != stream->write_index;
|
||||
++read_index) {
|
||||
__ockl_as_packet_t *packet = stream->base_address + read_index;
|
||||
|
||||
uint header = packet->header;
|
||||
if (get_packet_type(header) != __OCKL_AS_PACKET_READY)
|
||||
return false;
|
||||
|
||||
uint bytes = get_packet_bytes(header);
|
||||
unsigned char flags = get_packet_flags(header);
|
||||
unsigned long connection_id = packet->connection_id;
|
||||
unsigned char *payload = packet->payload;
|
||||
uint8_t service = get_packet_service(header);
|
||||
|
||||
if ((flags & __OCKL_AS_CONNECTION_BEGIN) !=
|
||||
(buffers.count(connection_id) == 0))
|
||||
return false;
|
||||
|
||||
std::string &buf = buffers[connection_id];
|
||||
buf.insert(buf.end(), payload, payload + bytes);
|
||||
|
||||
if (flags & __OCKL_AS_CONNECTION_END) {
|
||||
if (buf != strExpected[service])
|
||||
return false;
|
||||
strRecd[buf] += 1;
|
||||
buffers.erase(connection_id);
|
||||
}
|
||||
}
|
||||
|
||||
int expected_counts[numExpected];
|
||||
for (int ii = 0; ii != numExpected; ++ii) {
|
||||
expected_counts[ii] = NUM_THREADS / numExpected;
|
||||
if (ii < (NUM_THREADS % numExpected)) {
|
||||
++expected_counts[ii];
|
||||
}
|
||||
}
|
||||
|
||||
if (strRecd.size() != numExpected)
|
||||
return false;
|
||||
|
||||
for (int ii = 0; ii != numExpected; ++ii) {
|
||||
std::string mystr(strExpected[ii]);
|
||||
if (strRecd[mystr] != expected_counts[ii])
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool
|
||||
ocklAsSplitMessage()
|
||||
{
|
||||
bool success = true;
|
||||
|
||||
unsigned int numPackets = NUM_PACKETS_LARGE;
|
||||
unsigned int bufferSize =
|
||||
sizeof(__ockl_as_stream_t) + numPackets * __OCKL_AS_PACKET_SIZE;
|
||||
|
||||
void *buffer;
|
||||
HIPCHECK(hipHostMalloc(&buffer, bufferSize));
|
||||
|
||||
__ockl_as_stream_t *stream = createStream(buffer, bufferSize, numPackets);
|
||||
|
||||
hipLaunchKernelGGL(splitMessage, dim3(NUM_BLOCKS), dim3(THREADS_PER_BLOCK),
|
||||
0, 0, stream);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
if (!checkSplitMessage(stream)) {
|
||||
test_failed(__func__);
|
||||
success = false;
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostFree(buffer));
|
||||
return success;
|
||||
}
|
||||
|
||||
#define TESTNAME "hipAsynchronousStreams"
|
||||
int
|
||||
main(int argc, char **argv)
|
||||
{
|
||||
bool success = true;
|
||||
|
||||
success &= ocklAsSinglePacketSingleProducer();
|
||||
success &= ocklAsMultipleProducers();
|
||||
success &= ocklAsDropPackets();
|
||||
success &= ocklAsMixedProducers();
|
||||
success &= ocklAsSplitMessage();
|
||||
|
||||
hipDeviceReset();
|
||||
|
||||
if (success) {
|
||||
test_passed(TESTNAME);
|
||||
return 0;
|
||||
}
|
||||
|
||||
failed(TESTNAME);
|
||||
}
|
||||
Ссылка в новой задаче
Block a user