AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Merge 'master' into 'amd-master'

فهرست منبع 
Change-Id: Ibb1dbe78aa52c8aeb77851ee1471e1edaec4c654
This commit is contained in:
Jenkins
2019-09-09 05:10:28 -04:00
والد 6c776daba1 29d09eb269
کامیت cbe6b65e2c
20فایلهای تغییر یافته به همراه1966 افزوده شده و 1107 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
CMakeLists.txt
+8
مشاهده پرونده
@@ -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})
bin/hipify-perl
+1506 -282
مشاهده پرونده
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است Diff را بارگزاری کن
docs/markdown/hip_porting_guide.md
+15 -16
مشاهده پرونده
@@ -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
hipify-clang/CMakeLists.txt
+1 -8
مشاهده پرونده
@@ -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)
hipify-clang/README.md
+72 -68
مشاهده پرونده
@@ -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")
hipify-clang/src/CUDA2HIP.h
+2
مشاهده پرونده
@@ -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.
hipify-clang/src/CUDA2HIP_Device_functions.cpp
+48
مشاهده پرونده
@@ -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}},
};
hipify-clang/src/HipifyAction.cpp
+35 -5
مشاهده پرونده
@@ -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;
}
hipify-clang/src/HipifyAction.h
+2
مشاهده پرونده
@@ -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);
};
hipify-clang/src/ReplacementsFrontendActionFactory.h
+6
مشاهده پرونده
@@ -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
};
hipify-clang/src/Statistics.cpp
+1
مشاهده پرونده
@@ -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
hipify-clang/src/Statistics.h
+1
مشاهده پرونده
@@ -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,
include/hip/hcc_detail/device_functions.h
+1 -1
مشاهده پرونده
@@ -970,7 +970,7 @@ static void __barrier(int n)
__device__
inline
__attribute__((noduplicate))
__attribute__((convergent))
void __syncthreads()
{
__barrier(__CLK_LOCAL_MEM_FENCE);
include/hip/hcc_detail/hip_runtime.h
+2 -5
مشاهده پرونده
@@ -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,
src/hip_hcc.cpp
+11 -2
مشاهده پرونده
@@ -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
};
}
src/hip_memory.cpp
-7
مشاهده پرونده
@@ -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);
tests/hipify-clang/unit_tests/device/math_functions.cu
+46
مشاهده پرونده
@@ -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;
}
tests/hipify-clang/unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp
+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);
}
tests/hipify-clang/unit_tests/samples/axpy.cu
+11
مشاهده پرونده
@@ -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;
tests/src/deviceLib/hipAsynchronousStreams.cpp
-713
مشاهده پرونده
@@ -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);
}