Merge 'master' into 'amd-master'

Change-Id: I387d49269f0314b38db5e77eacc1be636280620d
2018-05-08 04:09:43 -05:00
commit b5fa30cfc6
@@ -102,9 +102,9 @@
 |   **CUDA**                                                |   **HIP**                     |
 |-----------------------------------------------------------|-------------------------------|
-| `cudaConfigureCall`                                       |                               |
+| `cudaConfigureCall`                                       | `hipConfigureCall`            |
-| `cudaLaunch`                                              |                               |
+| `cudaLaunch`                                              | `hipLaunchByPtr`              |
-| `cudaSetupArgument`                                       |                               |
+| `cudaSetupArgument`                                       | `hipSetupArgument`            |
 ## **9. Memory Management**
@@ -158,7 +158,7 @@
 | `cudaMemcpyToSymbolAsync`                                 | `hipMemcpyToSymbolAsync`      |
 | `cudaMemset`                                              | `hipMemset`                   |
 | `cudaMemset2D`                                            | `hipMemset2D`                 |
-| `cudaMemset2DAsync`                                       |                               |
+| `cudaMemset2DAsync`                                       | `hipMemset2DAsync`            |
 | `cudaMemset3D`                                            |                               |
 | `cudaMemset3DAsync`                                       |                               |
 | `cudaMemsetAsync`                                         | `hipMemsetAsync`              |
@@ -338,8 +338,8 @@
 |   **CUDA**                                                |   **HIP**                     |
 |-----------------------------------------------------------|-------------------------------|
-| `cudaCreateSurfaceObject`                                 |                               |
+| `cudaCreateSurfaceObject`                                 | `hipCreateSurfaceObject`      |
-| `cudaDestroySurfaceObject`                                |                               |
+| `cudaDestroySurfaceObject`                                | `hipDestroySurfaceObject`     |
 | `cudaGetSurfaceObjectResourceDesc`                        |                               |
 ## **27. Version Management**
@@ -675,10 +675,10 @@
 |            0 |*`cudaSharedMemBankSizeDefault`*               |*`hipSharedMemBankSizeDefault`*                       |
 |            1 |*`cudaSharedMemBankSizeFourByte`*              |*`hipSharedMemBankSizeFourByte`*                      |
 |            2 |*`cudaSharedMemBankSizeEightByte`*             |*`hipSharedMemBankSizeEightByte`*                     |
-| enum         |***`cudaSurfaceBoundaryMode`***                |                                                      |
+| enum         |***`cudaSurfaceBoundaryMode`***                |***`hipSurfaceBoundaryMode`***                        |
-|            0 |*`cudaBoundaryModeZero`*                       |                                                      |
+|            0 |*`cudaBoundaryModeZero`*                       |*`hipBoundaryModeZero`*                               |
-|            1 |*`cudaBoundaryModeClamp`*                      |                                                      |
+|            1 |*`cudaBoundaryModeClamp`*                      |*`hipBoundaryModeClamp`*                              |
-|            2 |*`cudaBoundaryModeTrap`*                       |                                                      |
+|            2 |*`cudaBoundaryModeTrap`*                       |*`hipBoundaryModeTrap`*                               |
 | enum         |***`cudaSurfaceFormatMode`***                  |                                                      |
 |            0 |*`cudaFormatModeForced`*                       |                                                      |
 |            1 |*`cudaFormatModeAuto`*                         |                                                      |
@@ -742,7 +742,7 @@
 | typedef      | `cudaOutputMode_t`                            |                                                      |
 | typedef      | `cudaStream_t`                                | `hipStream_t`                                        |
 | typedef      | `cudaStreamCallback_t`                        | `hipStreamCallback_t`                                |
-| typedef      | `cudaSurfaceObject_t`                         |                                                      |
+| typedef      | `cudaSurfaceObject_t`                         | `hipSurfaceObject_t`                                 |
 | typedef      | `cudaTextureObject_t`                         | `hipTextureObject_t`                                 |
 | typedef      | `CUuuid_stcudaUUID_t`                         |                                                      |
 | define       | `CUDA_IPC_HANDLE_SIZE`                        |                                                      |
@@ -17,41 +17,146 @@
 ### Dependencies
-`hipify-clang` requires clang+llvm of at least version 3.8.
+`hipify-clang` requires:
 1. LLVM+CLANG of at least version 3.8.0, latest stable release is 6.0.0.
 2. CUDA at least version 7.5, latest supported release is 9.0.
-In most cases, you can get a suitable version of clang+llvm with your package manager.
+| **LLVM release version** | **CUDA latest supported version** |
 |:------------------------:|:---------------------------------:|
 | 3.8.0                    | 7.5                               |
 | 3.8.1                    | 7.5                               |
 | 3.9.0                    | 7.5                               |
 | 3.9.1                    | 7.5                               |
 | 4.0.0                    | 8.0                               |
 | 4.0.1                    | 8.0                               |
 | 5.0.0                    | 8.0                               |
 | 5.0.1                    | 8.0                               |
 | 6.0.0                    | 9.0                               |
-Failing that, you can [download a release archive](http://releases.llvm.org/), extract it somewhere, and set
+In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
-[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.0/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it.
+
 Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set
 [CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.0/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.0\dist`
 ### Build
 Assuming this repository is at `./HIP`:
 ```shell
-mkdir build inst
+cd hipify-clang
-
+mkdir build dist
 cd build
 cmake \
- -DCMAKE_INSTALL_PREFIX=../inst \
+ -DCMAKE_INSTALL_PREFIX=../dist \
 -DCMAKE_BUILD_TYPE=Release \
- -DBUILD_HIPIFY_CLANG=ON \
+ ..
 ../HIP
 make -j install
 ```
 On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make`.
-The binary can then be found at `./inst/bin/hipify-clang`.
+Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested, `LLVM+CLANG` should be built in `Debug` mode as well.
 64 bit build mode `-Thost=x64` is supported as well, `LLVM+CLANG` should be built (installed) in 64bit mode as well.
 The binary can then be found at `./dist/bin/hipify-clang`.
 ### Test
 `hipify-clang` has unit tests using LLVM [`lit`](https://llvm.org/docs/CommandGuide/lit.html)/[`FileCheck`](https://llvm.org/docs/CommandGuide/FileCheck.html).
 **LLVM+CLANG should be built from sources, Pre-Built Binaries are not exhaustive for testing.**
 To run it:
-1. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
+1. Download [`LLVM`](http://releases.llvm.org/6.0.0/llvm-6.0.0.src.tar.xz)+[`CLANG`](http://releases.llvm.org/6.0.0/cfe-6.0.0.src.tar.xz) sources.
-2. Ensure `socat` is installed - your distro almost certainly has a package for this.
+2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html).
-3. Build with the `HIPIFY_CLANG_TESTS` option turned on.
+    For instance:
-4. `make test-hipify`
+    ```shell
   cd llvm
   mkdir build dist
   cd build
   cmake \
     -DCMAKE_INSTALL_PREFIX=../dist \
     -DLLVM_SOURCE_DIR=../llvm \
     -DCMAKE_BUILD_TYPE=Debug \
     -Thost=x64 \
     ../llvm
    make -j install
    ```
    On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`.
 3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.5 is installed.
    * Having multiple CUDA installations in order to choose a concrete version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
      `-DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
      * On Windows `CUDA_SDK_ROOT_DIR` option should be specified as well:
      `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`
 4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of version corresponding to CUDA's version is installed.
    * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:
      `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1`
 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
 6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
    * installing `lit` into `python` might be required:
      `python f:/LLVM/6.0.0/llvm/utils/lit/setup.py install`,
      where `f:/LLVM/6.0.0/llvm` is LLVM sources root directory.
    * Starting with LLVM 6.0.0 path to llvm-lit.py script should be specified by the `LLVM_EXTERNAL_LIT` option:
      `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py`,
      where `f:/LLVM/6.0.0/build/Debug` is LLVM build directory.
 7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1.
 8. `make test-hipify`
    On Windows after `cmake` the project `test-hipify` in the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make test-hipify`.
 ### Windows
 On Windows the following tested configuration is recommended:
 LLVM 6.0.0 (exact), CUDA 9.0 (exact), cudnn-9.0 (exact), Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min).
 Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
 ```shell
 cmake
 -G "Visual Studio 15 2017 Win64" \
 -DHIPIFY_CLANG_TESTS=1 \
 -DCMAKE_BUILD_TYPE=Debug \
 -DCMAKE_INSTALL_PREFIX=../dist \
 -DCMAKE_PREFIX_PATH=f:/LLVM/6.0.0/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.1 \
 -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py \
 ..
 ```
 A corresponding successful output:
 ```shell
 -- Found LLVM 6.0.0:
 --    - CMake module path: F:/LLVM/6.0.0/dist/lib/cmake/llvm
 --    - Include path     : F:/LLVM/6.0.0/dist/include
 --    - Binary path      : F:/LLVM/6.0.0/dist/bin
 -- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "2.7")
 -- Found lit: C:/Program Files/Python36/Scripts/lit.exe
 -- Found FileCheck: F:/LLVM/6.0.0/dist/bin/FileCheck.exe
 -- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0")
 -- Configuring done
 -- Generating done
 -- Build files have been written to: f:/HIP/hipify-clang/build
 ```
 ## Running and using hipify-clang
@@ -65,6 +170,7 @@ hipify-clang square.cu -- \
  --cuda-path=/opt/cuda \
  --cuda-gpu-arch=sm_30 \
  -isystem /opt/cuda/samples/common/inc
  -I /opt/cuda/cuDNN
 ```
 `hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
@@ -79,5 +185,5 @@ The information contained herein is for informational purposes only, and is subj
 AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies.
-Copyright (c) 2014-2017 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2014-2018 Advanced Micro Devices, Inc. All rights reserved.
@@ -261,12 +261,13 @@ const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
    // typedefs
    {"cudaTextureObject_t",  {"hipTextureObject_t",  CONV_TEX,     API_RUNTIME}},
    {"cudaSurfaceObject_t",  {"hipSurfaceObject_t",  CONV_SURFACE, API_RUNTIME}},
    // enums
    {"cudaResourceType",        {"hipResourceType",        CONV_TEX,     API_RUNTIME}},    // API_Driver ANALOGUE (CUresourcetype)
    {"cudaResourceViewFormat",  {"hipResourceViewFormat",  CONV_TEX,     API_RUNTIME}},    // API_Driver ANALOGUE (CUresourceViewFormat)
    {"cudaTextureAddressMode",  {"hipTextureAddressMode",  CONV_TEX,     API_RUNTIME}},
-    {"cudaSurfaceBoundaryMode", {"hipSurfaceBoundaryMode", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaSurfaceBoundaryMode", {"hipSurfaceBoundaryMode", CONV_SURFACE, API_RUNTIME}},
    {"cudaSurfaceFormatMode", {"hipSurfaceFormatMode", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
@@ -403,6 +404,9 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
    {"curand_precalc.h",          {"hiprand_kernel.h",         CONV_INCLUDE,             API_RAND}},
    {"curand_uniform.h",          {"hiprand_kernel.h",         CONV_INCLUDE,             API_RAND}},
    // CUDNN includes
    {"cudnn.h",                   {"hipDNN.h",                 CONV_INCLUDE_CUDA_MAIN_H, API_RAND}},
    // HIP includes
    // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h
    //    {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}},
@@ -1484,7 +1488,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
    {"cudaMemset",        {"hipMemset",        CONV_MEM, API_RUNTIME}},
    {"cudaMemsetAsync",   {"hipMemsetAsync",   CONV_MEM, API_RUNTIME}},
    {"cudaMemset2D",      {"hipMemset2D",      CONV_MEM, API_RUNTIME}},
-    {"cudaMemset2DAsync", {"hipMemset2DAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaMemset2DAsync", {"hipMemset2DAsync", CONV_MEM, API_RUNTIME}},
    {"cudaMemset3D",      {"hipMemset3D",      CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
    {"cudaMemset3DAsync", {"hipMemset3DAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
@@ -1745,9 +1749,9 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
    {"cudaSetDoubleForHost",       {"hipSetDoubleForHost",       CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
    // Execution Control [deprecated since 7.0]
-    {"cudaConfigureCall", {"hipConfigureCall", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaConfigureCall", {"hipConfigureCall", CONV_EXEC, API_RUNTIME}},
-    {"cudaLaunch",        {"hipLaunch",        CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaLaunch",        {"hipLaunchByPtr",   CONV_EXEC, API_RUNTIME}},
-    {"cudaSetupArgument", {"hipSetupArgument", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaSetupArgument", {"hipSetupArgument", CONV_EXEC, API_RUNTIME}},
    // Version Management
    {"cudaDriverGetVersion",  {"hipDriverGetVersion",  CONV_VERSION, API_RUNTIME}},
@@ -1889,17 +1893,17 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
    {"cudaGetSurfaceReference", {"hipGetSurfaceReference", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
    // enum cudaSurfaceBoundaryMode
-    {"cudaBoundaryModeZero",    {"hipBoundaryModeZero",    CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaBoundaryModeZero",    {"hipBoundaryModeZero",    CONV_SURFACE, API_RUNTIME}},
-    {"cudaBoundaryModeClamp",   {"hipBoundaryModeClamp",   CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaBoundaryModeClamp",   {"hipBoundaryModeClamp",   CONV_SURFACE, API_RUNTIME}},
-    {"cudaBoundaryModeTrap",    {"hipBoundaryModeTrap",    CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaBoundaryModeTrap",    {"hipBoundaryModeTrap",    CONV_SURFACE, API_RUNTIME}},
    // enum cudaSurfaceFormatMode
    {"cudaFormatModeForced",    {"hipFormatModeForced",    CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
    {"cudaFormatModeAuto",      {"hipFormatModeAuto",      CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
    // Surface Object Management
-    {"cudaCreateSurfaceObject",          {"hipCreateSurfaceObject",          CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaCreateSurfaceObject",          {"hipCreateSurfaceObject",          CONV_SURFACE, API_RUNTIME}},
-    {"cudaDestroySurfaceObject",         {"hipDestroySurfaceObject",         CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
+    {"cudaDestroySurfaceObject",         {"hipDestroySurfaceObject",         CONV_SURFACE, API_RUNTIME}},
    {"cudaGetSurfaceObjectResourceDesc", {"hipGetSurfaceObjectResourceDesc", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
    // Inter-Process Communications (IPC)
@@ -2880,6 +2884,83 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
    {"curand_poisson4",                               {"hiprand_poisson4",                               CONV_DEVICE_FUNC, API_RAND}},
    {"curand_Philox4x32_10",                          {"hiprand_Philox4x32_10",                          CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
    // unchanged function names: skipahead, skipahead_sequence, skipahead_subsequence
    ///////////////////////////// cuDNN /////////////////////////////
    {"cudnnContext",                                  {"hipdnnContext",                                  CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
    {"cudnnHandle_t",                                 {"hipdnnHandle_t",                                 CONV_TYPE, API_DNN}},
    {"cudnnStatus_t",                                 {"hipdnnStatus_t",                                 CONV_TYPE, API_DNN}},
    {"CUDNN_STATUS_SUCCESS",                          {"HIPDNN_STATUS_SUCCESS",                          CONV_NUMERIC_LITERAL, API_DNN}},    //  0
    {"CUDNN_STATUS_NOT_INITIALIZED",                  {"HIPDNN_STATUS_NOT_INITIALIZED",                  CONV_NUMERIC_LITERAL, API_DNN}},    //  1
    {"CUDNN_STATUS_ALLOC_FAILED",                     {"HIPDNN_STATUS_ALLOC_FAILED",                     CONV_NUMERIC_LITERAL, API_DNN}},    //  2
    {"CUDNN_STATUS_BAD_PARAM",                        {"HIPDNN_STATUS_BAD_PARAM",                        CONV_NUMERIC_LITERAL, API_DNN}},    //  3
    {"CUDNN_STATUS_INTERNAL_ERROR",                   {"HIPDNN_STATUS_INTERNAL_ERROR",                   CONV_NUMERIC_LITERAL, API_DNN}},    //  4
    {"CUDNN_STATUS_INVALID_VALUE",                    {"HIPDNN_STATUS_INVALID_VALUE",                    CONV_NUMERIC_LITERAL, API_DNN}},    //  5
    {"CUDNN_STATUS_ARCH_MISMATCH",                    {"HIPDNN_STATUS_ARCH_MISMATCH",                    CONV_NUMERIC_LITERAL, API_DNN}},    //  6
    {"CUDNN_STATUS_MAPPING_ERROR",                    {"HIPDNN_STATUS_MAPPING_ERROR",                    CONV_NUMERIC_LITERAL, API_DNN}},    //  7
    {"CUDNN_STATUS_EXECUTION_FAILED",                 {"HIPDNN_STATUS_EXECUTION_FAILED",                 CONV_NUMERIC_LITERAL, API_DNN}},    //  8
    {"CUDNN_STATUS_NOT_SUPPORTED",                    {"HIPDNN_STATUS_NOT_SUPPORTED",                    CONV_NUMERIC_LITERAL, API_DNN}},    //  9
    {"CUDNN_STATUS_LICENSE_ERROR",                    {"HIPDNN_STATUS_LICENSE_ERROR",                    CONV_NUMERIC_LITERAL, API_DNN}},    // 10
    {"CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING",     {"HIPDNN_STATUS_RUNTIME_PREREQUISITE_MISSING",     CONV_NUMERIC_LITERAL, API_DNN}},    // 11
    {"CUDNN_STATUS_RUNTIME_IN_PROGRESS",              {"HIPDNN_STATUS_RUNTIME_IN_PROGRESS",              CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},    // 12
    {"CUDNN_STATUS_RUNTIME_FP_OVERFLOW",              {"HIPDNN_STATUS_RUNTIME_FP_OVERFLOW",              CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},    // 13
    {"cudnnRuntimeTag_t",                             {"hipdnnRuntimeTag_t",                             CONV_TYPE, API_DNN, HIP_UNSUPPORTED}},
    {"cudnnTensorDescriptor_t",                       {"hipdnnTensorDescriptor_t",                       CONV_TYPE, API_DNN}},
    {"cudnnConvolutionDescriptor_t",                  {"hipdnnConvolutionDescriptor_t",                  CONV_TYPE, API_DNN}},
    {"cudnnConvolutionMode_t",                        {"hipdnnConvolutionMode_t",                        CONV_TYPE, API_DNN}},
    {"CUDNN_CONVOLUTION",                             {"HIPDNN_CONVOLUTION",                             CONV_NUMERIC_LITERAL, API_DNN}},    // 0
    {"CUDNN_CROSS_CORRELATION",                       {"HIPDNN_CROSS_CORRELATION",                       CONV_NUMERIC_LITERAL, API_DNN}},    // 1
    {"cudnnTensorFormat_t",                           {"hipdnnTensorFormat_t",                           CONV_TYPE, API_DNN}},
    {"CUDNN_TENSOR_NCHW",                             {"HIPDNN_TENSOR_NCHW",                             CONV_NUMERIC_LITERAL, API_DNN}},    // 0
    {"CUDNN_TENSOR_NHWC",                             {"HIPDNN_TENSOR_NHWC",                             CONV_NUMERIC_LITERAL, API_DNN}},    // 1
    {"CUDNN_TENSOR_NCHW_VECT_C",                      {"HIPDNN_TENSOR_NCHW_VECT_C",                      CONV_NUMERIC_LITERAL, API_DNN}},    // 2
    {"cudnnDataType_t",                               {"hipdnnDataType_t",                               CONV_TYPE, API_DNN}},
    {"CUDNN_DATA_FLOAT",                              {"HIPDNN_DATA_FLOAT",                              CONV_NUMERIC_LITERAL, API_DNN}},    // 0
    {"CUDNN_DATA_DOUBLE",                             {"HIPDNN_DATA_DOUBLE",                             CONV_NUMERIC_LITERAL, API_DNN}},    // 1
    {"CUDNN_DATA_HALF",                               {"HIPDNN_DATA_HALF",                               CONV_NUMERIC_LITERAL, API_DNN}},    // 2
    {"CUDNN_DATA_INT8",                               {"HIPDNN_DATA_INT8",                               CONV_NUMERIC_LITERAL, API_DNN}},    // 3
    {"CUDNN_DATA_INT32",                              {"HIPDNN_DATA_INT32",                              CONV_NUMERIC_LITERAL, API_DNN}},    // 4
    {"CUDNN_DATA_INT8x4",                             {"HIPDNN_DATA_INT8x4",                             CONV_NUMERIC_LITERAL, API_DNN}},    // 5
    {"CUDNN_DATA_UINT8",                              {"HIPDNN_DATA_UINT8",                              CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},  // 6
    {"CUDNN_DATA_UINT8x4",                            {"HIPDNN_DATA_UINT8x4",                            CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}},  // 7
    {"cudnnConvolutionFwdAlgo_t",                           {"hipdnnConvolutionFwdAlgo_t",                           CONV_TYPE, API_DNN}},
    {"CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM",            {"HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM",            CONV_NUMERIC_LITERAL, API_DNN}},    // 0
    {"CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM",    {"HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM",    CONV_NUMERIC_LITERAL, API_DNN}},    // 1
    {"CUDNN_CONVOLUTION_FWD_ALGO_GEMM",                     {"HIPDNN_CONVOLUTION_FWD_ALGO_GEMM",                     CONV_NUMERIC_LITERAL, API_DNN}},    // 2
    {"CUDNN_CONVOLUTION_FWD_ALGO_DIRECT",                   {"HIPDNN_CONVOLUTION_FWD_ALGO_DIRECT",                   CONV_NUMERIC_LITERAL, API_DNN}},    // 3
    {"CUDNN_CONVOLUTION_FWD_ALGO_FFT",                      {"HIPDNN_CONVOLUTION_FWD_ALGO_FFT",                      CONV_NUMERIC_LITERAL, API_DNN}},    // 4
    {"CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING",               {"HIPDNN_CONVOLUTION_FWD_ALGO_FFT_TILING",               CONV_NUMERIC_LITERAL, API_DNN}},    // 5
    {"CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD",                 {"HIPDNN_CONVOLUTION_FWD_ALGO_WINOGRAD",                 CONV_NUMERIC_LITERAL, API_DNN}},    // 6
    {"CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED",        {"HIPDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED",        CONV_NUMERIC_LITERAL, API_DNN}},    // 7
    {"CUDNN_CONVOLUTION_FWD_ALGO_COUNT",                    {"HIPDNN_CONVOLUTION_FWD_ALGO_COUNT",                    CONV_NUMERIC_LITERAL, API_DNN}},    // 8
    {"cudnnConvolutionFwdPreference_t",                     {"hipdnnConvolutionFwdPreference_t",                     CONV_TYPE, API_DNN}},
    {"CUDNN_CONVOLUTION_FWD_NO_WORKSPACE",                  {"HIPDNN_CONVOLUTION_FWD_NO_WORKSPACE",                  CONV_NUMERIC_LITERAL, API_DNN}},    // 0
    {"CUDNN_CONVOLUTION_FWD_PREFER_FASTEST",                {"HIPDNN_CONVOLUTION_FWD_PREFER_FASTEST",                CONV_NUMERIC_LITERAL, API_DNN}},    // 1
    {"CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT",       {"HIPDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT",       CONV_NUMERIC_LITERAL, API_DNN}},    // 2
    {"cudnnFilterDescriptor_t",                             {"hipdnnFilterDescriptor_t",                             CONV_TYPE, API_DNN}},
    {"cudnnGetVersion",                               {"hipdnnGetVersion",                               CONV_VERSION, API_DNN}},
    {"cudnnGetCudartVersion",                         {"hipdnnGetCudartVersion",                         CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
    {"cudnnGetErrorString",                           {"hipdnnGetErrorString",                           CONV_ERROR, API_DNN}},
    {"cudnnCreate",                                   {"hipdnnCreate",                                   CONV_MATH_FUNC, API_DNN}},
    {"cudnnCreateTensorDescriptor",                   {"hipdnnCreateTensorDescriptor",                   CONV_MATH_FUNC, API_DNN}},
    {"cudnnSetTensor4dDescriptor",                    {"hipdnnSetTensor4dDescriptor",                    CONV_MATH_FUNC, API_DNN}},
    {"cudnnSetConvolution2dDescriptor",               {"hipdnnSetConvolution2dDescriptor",               CONV_MATH_FUNC, API_DNN}},
    {"cudnnGetConvolution2dForwardOutputDim",         {"hipdnnGetConvolution2dForwardOutputDim",         CONV_MATH_FUNC, API_DNN}},
    {"cudnnCreateFilterDescriptor",                   {"hipdnnCreateFilterDescriptor",                   CONV_MATH_FUNC, API_DNN}},
    {"cudnnSetFilter4dDescriptor",                    {"hipdnnSetFilter4dDescriptor",                    CONV_MATH_FUNC, API_DNN}},
    {"cudnnCreateConvolutionDescriptor",              {"hipdnnCreateConvolutionDescriptor",              CONV_MATH_FUNC, API_DNN}},
    {"cudnnGetConvolutionForwardAlgorithm",           {"hipdnnGetConvolutionForwardAlgorithm",           CONV_MATH_FUNC, API_DNN}},
    {"cudnnConvolutionForward",                       {"hipdnnConvolutionForward",                       CONV_MATH_FUNC, API_DNN}},
    {"cudnnGetConvolutionForwardWorkspaceSize",       {"hipdnnGetConvolutionForwardWorkspaceSize",       CONV_MATH_FUNC, API_DNN}},
    {"cudnnDestroyTensorDescriptor",                  {"hipdnnDestroyTensorDescriptor",                  CONV_MATH_FUNC, API_DNN}},
    {"cudnnDestroyConvolutionDescriptor",             {"hipdnnDestroyConvolutionDescriptor",             CONV_MATH_FUNC, API_DNN}},
    {"cudnnDestroyFilterDescriptor",                  {"hipdnnDestroyFilterDescriptor",                  CONV_MATH_FUNC, API_DNN}},
    {"cudnnDestroyFilterDescriptor",                  {"hipdnnDestroyFilterDescriptor",                  CONV_MATH_FUNC, API_DNN}},
    {"cudnnDestroy",                                  {"hipdnnDestroy",                                  CONV_MATH_FUNC, API_DNN}},
 };
 const std::map<llvm::StringRef, hipCounter>& CUDA_RENAMES_MAP() {
@@ -160,6 +160,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
                        insertedRANDHeader = true;
                        return false;
                    }
                case API_DNN:
                  if (insertedDNNHeader) { return true; }
                    insertedDNNHeader = true;
                    return false;
                default:
                    return false;
            }
@@ -27,6 +27,7 @@ private:
    bool insertedBLASHeader = false;
    bool insertedRANDHeader = false;
    bool insertedRAND_kernelHeader = false;
    bool insertedDNNHeader = false;
    bool firstHeader = false;
    bool pragmaOnce = false;
    clang::SourceLocation firstHeaderLoc;
@@ -14,7 +14,7 @@ const char *counterNames[NUM_CONV_TYPES] = {
 };
 const char *apiNames[NUM_API_TYPES] = {
-    "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API"
+    "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API"
 };
 namespace {
@@ -56,6 +56,7 @@ enum ApiTypes {
    API_RUNTIME,
    API_BLAS,
    API_RAND,
    API_DNN,
    API_LAST
 };
 constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
@@ -0,0 +1,254 @@
 // RUN: %run_test hipify "%s" "%t" %cuda_args
 #include <iomanip>
 #include <iostream>
 #include <cstdlib>
 #include <vector>
 // CHECK: #include <hip/hip_runtime.h>
 #include <cuda.h>
 // CHECK: #include "hipDNN.h"
 #include "cudnn.h"
 // CHECK: hipError_t err = (f); \
 // CHECK: if (err != hipSuccess) { \
 #define CUDA_CALL(f) { \
  cudaError_t err = (f); \
  if (err != cudaSuccess) { \
    std::cout \
        << "    Error occurred: " << err << std::endl; \
    std::exit(1); \
  } \
 }
 // CHECK: hipdnnStatus_t err = (f); \
 // CHECK: if (err != HIPDNN_STATUS_SUCCESS) { \
 #define CUDNN_CALL(f) { \
    cudnnStatus_t err = (f); \
    if (err != CUDNN_STATUS_SUCCESS) { \
    std::cout \
        << "    Error occurred: " << err << std::endl; \
    std::exit(1); \
  } \
 }
 __global__ void dev_const(float *px, float k) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  px[tid] = k;
 }
 __global__ void dev_iota(float *px) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  px[tid] = tid;
 }
 void print(const float *data, int n, int c, int h, int w) {
  std::vector<float> buffer(1 << 20);
  // CHECK: CUDA_CALL(hipMemcpy(
  CUDA_CALL(cudaMemcpy(
        buffer.data(), data,
        n * c * h * w * sizeof(float),
        // CHECK: hipMemcpyDeviceToHost));
        cudaMemcpyDeviceToHost));
  int a = 0;
  for (int i = 0; i < n; ++i) {
    for (int j = 0; j < c; ++j) {
      std::cout << "n=" << i << ", c=" << j << ":" << std::endl;
      for (int k = 0; k < h; ++k) {
        for (int l = 0; l < w; ++l) {
          std::cout << std::setw(4) << std::right << buffer[a];
          ++a;
        }
        std::cout << std::endl;
      }
    }
  }
  std::cout << std::endl;
 }
 int main() {
  // CHECK: hipdnnHandle_t cudnn;
  cudnnHandle_t cudnn;
  // CHECK: CUDNN_CALL(hipdnnCreate(&cudnn));
  CUDNN_CALL(cudnnCreate(&cudnn));
  // input
  const int in_n = 1;
  const int in_c = 1;
  const int in_h = 5;
  const int in_w = 5;
  std::cout << "in_n: " << in_n << std::endl;
  std::cout << "in_c: " << in_c << std::endl;
  std::cout << "in_h: " << in_h << std::endl;
  std::cout << "in_w: " << in_w << std::endl;
  std::cout << std::endl;
  // CHECK: hipdnnTensorDescriptor_t in_desc;
  cudnnTensorDescriptor_t in_desc;
  // CHECK: CUDNN_CALL(hipdnnCreateTensorDescriptor(&in_desc));
  CUDNN_CALL(cudnnCreateTensorDescriptor(&in_desc));
  // CHECK: CUDNN_CALL(hipdnnSetTensor4dDescriptor(
  CUDNN_CALL(cudnnSetTensor4dDescriptor(
  // CHECK: in_desc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_FLOAT,
        in_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT,
        in_n, in_c, in_h, in_w));
  float *in_data;
  // CHECK: CUDA_CALL(hipMalloc(
  CUDA_CALL(cudaMalloc(
        &in_data, in_n * in_c * in_h * in_w * sizeof(float)));
  // filter
  const int filt_k = 1;
  const int filt_c = 1;
  const int filt_h = 2;
  const int filt_w = 2;
  std::cout << "filt_k: " << filt_k << std::endl;
  std::cout << "filt_c: " << filt_c << std::endl;
  std::cout << "filt_h: " << filt_h << std::endl;
  std::cout << "filt_w: " << filt_w << std::endl;
  std::cout << std::endl;
  // CHECK: hipdnnFilterDescriptor_t filt_desc;
  cudnnFilterDescriptor_t filt_desc;
  // CHECK: CUDNN_CALL(hipdnnCreateFilterDescriptor(&filt_desc));
  CUDNN_CALL(cudnnCreateFilterDescriptor(&filt_desc));
  // CHECK: CUDNN_CALL(hipdnnSetFilter4dDescriptor(
  CUDNN_CALL(cudnnSetFilter4dDescriptor(
        // CHECK: filt_desc, HIPDNN_DATA_FLOAT, HIPDNN_TENSOR_NCHW,
        filt_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW,
        filt_k, filt_c, filt_h, filt_w));
  float *filt_data;
  //  CUDA_CALL(hipMalloc(
  CUDA_CALL(cudaMalloc(
      &filt_data, filt_k * filt_c * filt_h * filt_w * sizeof(float)));
  // convolution
  const int pad_h = 1;
  const int pad_w = 1;
  const int str_h = 1;
  const int str_w = 1;
  const int dil_h = 1;
  const int dil_w = 1;
  std::cout << "pad_h: " << pad_h << std::endl;
  std::cout << "pad_w: " << pad_w << std::endl;
  std::cout << "str_h: " << str_h << std::endl;
  std::cout << "str_w: " << str_w << std::endl;
  std::cout << "dil_h: " << dil_h << std::endl;
  std::cout << "dil_w: " << dil_w << std::endl;
  std::cout << std::endl;
  // CHECK: hipdnnConvolutionDescriptor_t conv_desc;
  cudnnConvolutionDescriptor_t conv_desc;
  // CUDNN_CALL(hipdnnCreateConvolutionDescriptor(&conv_desc));
  CUDNN_CALL(cudnnCreateConvolutionDescriptor(&conv_desc));
  // CHECK: CUDNN_CALL(hipdnnSetConvolution2dDescriptor(
  CUDNN_CALL(cudnnSetConvolution2dDescriptor(
        conv_desc,
        pad_h, pad_w, str_h, str_w, dil_h, dil_w,
        // CHECK: HIPDNN_CONVOLUTION, HIPDNN_DATA_FLOAT));
        CUDNN_CONVOLUTION, CUDNN_DATA_FLOAT));
  // output
  int out_n;
  int out_c;
  int out_h;
  int out_w;
  // CHECK: CUDNN_CALL(hipdnnGetConvolution2dForwardOutputDim(
  CUDNN_CALL(cudnnGetConvolution2dForwardOutputDim(
        conv_desc, in_desc, filt_desc,
        &out_n, &out_c, &out_h, &out_w));
  std::cout << "out_n: " << out_n << std::endl;
  std::cout << "out_c: " << out_c << std::endl;
  std::cout << "out_h: " << out_h << std::endl;
  std::cout << "out_w: " << out_w << std::endl;
  std::cout << std::endl;
  // CHECK: hipdnnTensorDescriptor_t out_desc;
  cudnnTensorDescriptor_t out_desc;
  // CHECK: CUDNN_CALL(hipdnnCreateTensorDescriptor(&out_desc));
  CUDNN_CALL(cudnnCreateTensorDescriptor(&out_desc));
  // CHECK: CUDNN_CALL(hipdnnSetTensor4dDescriptor(
  CUDNN_CALL(cudnnSetTensor4dDescriptor(
        // CHECK: out_desc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_FLOAT,
        out_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT,
        out_n, out_c, out_h, out_w));
  float *out_data;
  // CHECK: CUDA_CALL(hipMalloc(
  CUDA_CALL(cudaMalloc(
        &out_data, out_n * out_c * out_h * out_w * sizeof(float)));
  // algorithm
  // CHECK: hipdnnConvolutionFwdAlgo_t algo;
  cudnnConvolutionFwdAlgo_t algo;
  // CHECK: CUDNN_CALL(hipdnnGetConvolutionForwardAlgorithm(
  CUDNN_CALL(cudnnGetConvolutionForwardAlgorithm(
        cudnn,
        in_desc, filt_desc, conv_desc, out_desc,
        // CHECK: HIPDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo));
        CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo));
  std::cout << "Convolution algorithm: " << algo << std::endl;
  std::cout << std::endl;
  // workspace
  size_t ws_size;
  // CHECK: CUDNN_CALL(hipdnnGetConvolutionForwardWorkspaceSize(
  CUDNN_CALL(cudnnGetConvolutionForwardWorkspaceSize(
        cudnn, in_desc, filt_desc, conv_desc, out_desc, algo, &ws_size));
  float *ws_data;
  // CHECK: CUDA_CALL(hipMalloc(&ws_data, ws_size));
  CUDA_CALL(cudaMalloc(&ws_data, ws_size));
  std::cout << "Workspace size: " << ws_size << std::endl;
  std::cout << std::endl;
  // perform
  float alpha = 1.f;
  float beta = 0.f;
  // CHECK: hipLaunchKernelGGL(dev_iota, dim3(in_w * in_h), dim3(in_n * in_c), 0, 0, in_data);
  // CHECK: hipLaunchKernelGGL(dev_const, dim3(filt_w * filt_h), dim3(filt_k * filt_c), 0, 0, filt_data, 1.f);
  dev_iota<<<in_w * in_h, in_n * in_c>>>(in_data);
  dev_const<<<filt_w * filt_h, filt_k * filt_c>>>(filt_data, 1.f);
  // CHECK: CUDNN_CALL(hipdnnConvolutionForward(
  CUDNN_CALL(cudnnConvolutionForward(
      cudnn,
      &alpha, in_desc, in_data, filt_desc, filt_data,
      conv_desc, algo, ws_data, ws_size,
      &beta, out_desc, out_data));
  // results
  std::cout << "in_data:" << std::endl;
  print(in_data, in_n, in_c, in_h, in_w);
  std::cout << "filt_data:" << std::endl;
  print(filt_data, filt_k, filt_c, filt_h, filt_w);
  std::cout << "out_data:" << std::endl;
  print(out_data, out_n, out_c, out_h, out_w);
  // finalizing
  // CHECK: CUDA_CALL(hipFree(ws_data));
  CUDA_CALL(cudaFree(ws_data));
  // CHECK: CUDA_CALL(hipFree(out_data));
  CUDA_CALL(cudaFree(out_data));
  // CHECK: CUDNN_CALL(hipdnnDestroyTensorDescriptor(out_desc));
  CUDNN_CALL(cudnnDestroyTensorDescriptor(out_desc));
  // CHECK: CUDNN_CALL(hipdnnDestroyConvolutionDescriptor(conv_desc));
  CUDNN_CALL(cudnnDestroyConvolutionDescriptor(conv_desc));
  // CHECK: CUDA_CALL(hipFree(filt_data));
  CUDA_CALL(cudaFree(filt_data));
  // CHECK: CUDNN_CALL(hipdnnDestroyFilterDescriptor(filt_desc));
  CUDNN_CALL(cudnnDestroyFilterDescriptor(filt_desc));
  // CHECK: CUDA_CALL(hipFree(in_data));
  CUDA_CALL(cudaFree(in_data));
  // CHECK: CUDNN_CALL(hipdnnDestroyTensorDescriptor(in_desc));
  CUDNN_CALL(cudnnDestroyTensorDescriptor(in_desc));
  // CHECK: CUDNN_CALL(hipdnnDestroy(cudnn));
  CUDNN_CALL(cudnnDestroy(cudnn));
  return 0;
 }
@@ -35,7 +35,7 @@
 // CHECK: #include <hiprand.h>
 #include <curand.h>
-// CHECK: if((x)!=hipSuccess) {
+// CHECK: if ((x) != hipSuccess) {
 #define CUDA_CALL(x)                                                                               \
    do {                                                                                           \
        if ((x) != cudaSuccess) {                                                                  \
@@ -43,7 +43,7 @@
            exit(EXIT_FAILURE);                                                                    \
        }                                                                                          \
    } while (0)
-// CHECK: if((x)!=HIPRAND_STATUS_SUCCESS) {
+// CHECK: if ((x) != HIPRAND_STATUS_SUCCESS) {
 #define CURAND_CALL(x)                                                                             \
    do {                                                                                           \
        if ((x) != CURAND_STATUS_SUCCESS) {                                                        \
@@ -59,9 +59,8 @@ const size_t DEFAULT_RAND_N = 1024 * 1024 * 128;
 // CHECK: typedef hiprandRngType_t rng_type_t;
 typedef curandRngType rng_type_t;
 // CHECK: using generate_func_type = std::function<hiprandStatus_t(hiprandGenerator_t, T *,
 // size_t)>;
 template <typename T>
 // CHECK: using generate_func_type = std::function<hiprandStatus_t(hiprandGenerator_t, T*, size_t)>;
 using generate_func_type = std::function<curandStatus_t(curandGenerator_t, T*, size_t)>;
 template <typename T>
@@ -71,7 +70,7 @@ void run_benchmark(const cli::Parser& parser, const rng_type_t rng_type,
    const size_t trials = parser.get<size_t>("trials");
    T* data;
-    // CHECK: CUDA_CALL(hipMalloc((void **)&data, size * sizeof(T)));
+    // CHECK: CUDA_CALL(hipMalloc((void**)&data, size * sizeof(T)));
    CUDA_CALL(cudaMalloc((void**)&data, size * sizeof(T)));
    // CHECK: hiprandGenerator_t generator;
@@ -80,8 +79,8 @@ void run_benchmark(const cli::Parser& parser, const rng_type_t rng_type,
    CURAND_CALL(curandCreateGenerator(&generator, rng_type));
    const size_t dimensions = parser.get<size_t>("dimensions");
-    // CHECK: hiprandStatus_t status = hiprandSetQuasiRandomGeneratorDimensions(generator,
+    // CHECK: hiprandStatus_t status = hiprandSetQuasiRandomGeneratorDimensions(generator, dimensions);
-    // dimensions); CHECK: if (status != HIPRAND_STATUS_TYPE_ERROR)
+    // CHECK: if (status != HIPRAND_STATUS_TYPE_ERROR)
    curandStatus_t status = curandSetQuasiRandomGeneratorDimensions(generator, dimensions);
    if (status != CURAND_STATUS_TYPE_ERROR)  // If the RNG is not quasi-random
    {
@@ -123,12 +122,12 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
                    const std::string& distribution) {
    if (distribution == "uniform-uint") {
        // CHECK: if (rng_type != HIPRAND_RNG_QUASI_SOBOL64 &&
-        // CHECK: rng_type != HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64)
+        // CHECK: rng_type != HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
        if (rng_type != CURAND_RNG_QUASI_SOBOL64 &&
            rng_type != CURAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
            run_benchmark<unsigned int>(
                parser, rng_type,
-                // CHECK: [](hiprandGenerator_t gen, unsigned int * data, size_t size) {
+                // CHECK: [](hiprandGenerator_t gen, unsigned int* data, size_t size) {
                // CHECK: return hiprandGenerate(gen, data, size);
                [](curandGenerator_t gen, unsigned int* data, size_t size) {
                    return curandGenerate(gen, data, size);
@@ -142,7 +141,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
            rng_type == CURAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
            run_benchmark<unsigned long long>(
                parser, rng_type,
-                // CHECK: [](hiprandGenerator_t gen, unsigned long long * data, size_t size) {
+                // CHECK: [](hiprandGenerator_t gen, unsigned long long* data, size_t size) {
                [](curandGenerator_t gen, unsigned long long* data, size_t size) {
                    // curandGenerateLongLong is yet unsupported by HIP
                    // CHECK-NOT: return hiprandGenerateLongLong(gen, data, size);
@@ -152,7 +151,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
    }
    if (distribution == "uniform-float") {
        run_benchmark<float>(parser, rng_type,
-                             // CHECK: [](hiprandGenerator_t gen, float * data, size_t size) {
+                             // CHECK: [](hiprandGenerator_t gen, float* data, size_t size) {
                             // CHECK: return hiprandGenerateUniform(gen, data, size);
                             [](curandGenerator_t gen, float* data, size_t size) {
                                 return curandGenerateUniform(gen, data, size);
@@ -160,7 +159,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
    }
    if (distribution == "uniform-double") {
        run_benchmark<double>(parser, rng_type,
-                              // CHECK: [](hiprandGenerator_t gen, double * data, size_t size) {
+                              // CHECK: [](hiprandGenerator_t gen, double* data, size_t size) {
                              // CHECK: return hiprandGenerateUniformDouble(gen, data, size);
                              [](curandGenerator_t gen, double* data, size_t size) {
                                  return curandGenerateUniformDouble(gen, data, size);
@@ -168,7 +167,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
    }
    if (distribution == "normal-float") {
        run_benchmark<float>(parser, rng_type,
-                             // CHECK: [](hiprandGenerator_t gen, float * data, size_t size) {
+                             // CHECK: [](hiprandGenerator_t gen, float* data, size_t size) {
                             // CHECK: return hiprandGenerateNormal(gen, data, size, 0.0f, 1.0f);
                             [](curandGenerator_t gen, float* data, size_t size) {
                                 return curandGenerateNormal(gen, data, size, 0.0f, 1.0f);
@@ -177,7 +176,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
    if (distribution == "normal-double") {
        run_benchmark<double>(
            parser, rng_type,
-            // CHECK: [](hiprandGenerator_t gen, double * data, size_t size) {
+            // CHECK: [](hiprandGenerator_t gen, double* data, size_t size) {
            // CHECK: return hiprandGenerateNormalDouble(gen, data, size, 0.0, 1.0);
            [](curandGenerator_t gen, double* data, size_t size) {
                return curandGenerateNormalDouble(gen, data, size, 0.0, 1.0);
@@ -185,7 +184,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
    }
    if (distribution == "log-normal-float") {
        run_benchmark<float>(parser, rng_type,
-                             // CHECK: [](hiprandGenerator_t gen, float * data, size_t size) {
+                             // CHECK: [](hiprandGenerator_t gen, float* data, size_t size) {
                             // CHECK: return hiprandGenerateLogNormal(gen, data, size, 0.0f, 1.0f);
                             [](curandGenerator_t gen, float* data, size_t size) {
                                 return curandGenerateLogNormal(gen, data, size, 0.0f, 1.0f);
@@ -194,7 +193,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
    if (distribution == "log-normal-double") {
        run_benchmark<double>(
            parser, rng_type,
-            // CHECK: [](hiprandGenerator_t gen, double * data, size_t size) {
+            // CHECK: [](hiprandGenerator_t gen, double* data, size_t size) {
            // CHECK: return hiprandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0);
            [](curandGenerator_t gen, double* data, size_t size) {
                return curandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0);
@@ -207,7 +206,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
                      << "lambda " << std::fixed << std::setprecision(1) << lambda << std::endl;
            run_benchmark<unsigned int>(
                parser, rng_type,
-                // CHECK: [lambda](hiprandGenerator_t gen, unsigned int * data, size_t size) {
+                // CHECK: [lambda](hiprandGenerator_t gen, unsigned int* data, size_t size) {
                // CHECK: return hiprandGeneratePoisson(gen, data, size, lambda);
                [lambda](curandGenerator_t gen, unsigned int* data, size_t size) {
                    return curandGeneratePoisson(gen, data, size, lambda);
@@ -42,16 +42,15 @@
 #include <curand_mtgp32_host.h>
 #include <curand_mtgp32dc_p_11213.h>
-// CHECK: hipError_t error = (x);
+// CHECK: if ((x) != hipSuccess) {
 // CHECK: if(error!=hipSuccess) {
 #define CUDA_CALL(x)                                                                               \
    do {                                                                                           \
-        cudaError_t error = (x);                                                                   \
+        if ((x) != cudaSuccess) {                                                                  \
-        if (error != cudaSuccess) {                                                                \
+            printf("Error at %s:%d\n", __FILE__, __LINE__);                                        \
            printf("Error %d at %s:%d\n", error, __FILE__, __LINE__);                              \
            exit(EXIT_FAILURE);                                                                    \
        }                                                                                          \
    } while (0)
 // CHECK: if ((x) != HIPRAND_STATUS_SUCCESS) {
 #define CURAND_CALL(x)                                                                             \
    do {                                                                                           \
        if ((x) != CURAND_STATUS_SUCCESS) {                                                        \
@@ -64,17 +63,22 @@
 const size_t DEFAULT_RAND_N = 1024 * 1024 * 128;
 #endif
-size_t next_power2(size_t x) {
+size_t next_power2(size_t x)
 {
    size_t power = 1;
-    while (power < x) {
+    while (power < x)
    {
        power *= 2;
    }
    return power;
 }
-template <typename GeneratorState>
+template<typename GeneratorState>
-__global__ void init_kernel(GeneratorState* states, const unsigned long long seed,
+__global__
-                            const unsigned long long offset) {
+void init_kernel(GeneratorState * states,
                 const unsigned long long seed,
                 const unsigned long long offset)
 {
    const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x;
    GeneratorState state;
    // CHECK: hiprand_init(seed, state_id, offset, &state);
@@ -82,32 +86,42 @@ __global__ void init_kernel(GeneratorState* states, const unsigned long long see
    states[state_id] = state;
 }
-template <typename GeneratorState, typename T, typename GenerateFunc, typename Extra>
+template<typename GeneratorState, typename T, typename GenerateFunc, typename Extra>
-__global__ void generate_kernel(GeneratorState* states, T* data, const size_t size,
+__global__
-                                const GenerateFunc& generate_func, const Extra extra) {
+void generate_kernel(GeneratorState * states,
                     T * data,
                     const size_t size,
                     const GenerateFunc& generate_func,
                     const Extra extra)
 {
    const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int stride = gridDim.x * blockDim.x;
    GeneratorState state = states[state_id];
    unsigned int index = state_id;
-    while (index < size) {
+    while(index < size)
    {
        data[index] = generate_func(&state, extra);
        index += stride;
    }
    states[state_id] = state;
 }
-template <typename GeneratorState>
+template<typename GeneratorState>
-struct runner {
+struct runner
-    GeneratorState* states;
+{
    GeneratorState * states;
-    runner(const size_t dimensions, const size_t blocks, const size_t threads,
+    runner(const size_t dimensions,
-           const unsigned long long seed, const unsigned long long offset) {
+           const size_t blocks,
           const size_t threads,
           const unsigned long long seed,
           const unsigned long long offset)
    {
        const size_t states_size = blocks * threads;
        // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(GeneratorState)));
-        CUDA_CALL(cudaMalloc((void**)&states, states_size * sizeof(GeneratorState)));
+        CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(GeneratorState)));
-        // CHECK: hipLaunchKernelGGL(init_kernel, dim3(blocks), dim3(threads), 0, 0, states, seed,
+        // CHECK: hipLaunchKernelGGL(init_kernel, dim3(blocks), dim3(threads), 0, 0, states, seed, offset);
        // offset);
        init_kernel<<<blocks, threads>>>(states, seed, offset);
        // CHECK: CUDA_CALL(hipPeekAtLastError());
        // CHECK: CUDA_CALL(hipDeviceSynchronize());
@@ -115,21 +129,33 @@ struct runner {
        CUDA_CALL(cudaDeviceSynchronize());
    }
-    ~runner() { CUDA_CALL(cudaFree(states)); }
+    ~runner()
    {
        CUDA_CALL(cudaFree(states));
    }
-    template <typename T, typename GenerateFunc, typename Extra>
+    template<typename T, typename GenerateFunc, typename Extra>
-    void generate(const size_t blocks, const size_t threads, T* data, const size_t size,
+    void generate(const size_t blocks,
-                  const GenerateFunc& generate_func, const Extra extra) {
+                  const size_t threads,
-        // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(blocks), dim3(threads), 0, 0, states,
+                  T * data,
-        // data, size, generate_func, extra);
+                  const size_t size,
                  const GenerateFunc& generate_func,
                  const Extra extra)
    {
        // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra);
        generate_kernel<<<blocks, threads>>>(states, data, size, generate_func, extra);
    }
 };
 // CHECK: void generate_kernel(hiprandStateMtgp32_t * states,
-template <typename T, typename GenerateFunc, typename Extra>
+template<typename T, typename GenerateFunc, typename Extra>
-__global__ void generate_kernel(curandStateMtgp32_t* states, T* data, const size_t size,
+__global__
-                                const GenerateFunc& generate_func, const Extra extra) {
+void generate_kernel(curandStateMtgp32_t * states,
                     T * data,
                     const size_t size,
                     const GenerateFunc& generate_func,
                     const Extra extra)
 {
    const unsigned int state_id = blockIdx.x;
    const unsigned int thread_id = threadIdx.x;
    unsigned int index = blockIdx.x * blockDim.x + threadIdx.x;
@@ -137,67 +163,80 @@ __global__ void generate_kernel(curandStateMtgp32_t* states, T* data, const size
    // CHECK: __shared__ hiprandStateMtgp32_t state;
    __shared__ curandStateMtgp32_t state;
-    if (thread_id == 0) state = states[state_id];
+    if (thread_id == 0)
        state = states[state_id];
    __syncthreads();
-    const size_t r = size % blockDim.x;
+    const size_t r = size%blockDim.x;
    const size_t size_rounded_up = r == 0 ? size : size + (blockDim.x - r);
-    while (index < size_rounded_up) {
+    while(index < size_rounded_up)
    {
        auto value = generate_func(&state, extra);
-        if (index < size) data[index] = value;
+        if(index < size)
            data[index] = value;
        index += stride;
    }
    __syncthreads();
-    if (thread_id == 0) states[state_id] = state;
+    if (thread_id == 0)
        states[state_id] = state;
 }
 // CHECK: struct runner<hiprandStateMtgp32_t>
-template <>
+template<>
-struct runner<curandStateMtgp32_t> {
+struct runner<curandStateMtgp32_t>
 {
    // CHECK: hiprandStateMtgp32_t * states;
-    curandStateMtgp32_t* states;
+    curandStateMtgp32_t * states;
-    mtgp32_kernel_params_t* d_param;
+    mtgp32_kernel_params_t * d_param;
-    runner(const size_t dimensions, const size_t blocks, const size_t threads,
+    runner(const size_t dimensions,
-           const unsigned long long seed, const unsigned long long offset) {
+           const size_t blocks,
           const size_t threads,
           const unsigned long long seed,
           const unsigned long long offset)
    {
        const size_t states_size = std::min((size_t)200, blocks);
-        // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size *
+        // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(hiprandStateMtgp32_t)));
-        // sizeof(hiprandStateMtgp32_t)));
+        CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateMtgp32_t)));
        CUDA_CALL(cudaMalloc((void**)&states, states_size * sizeof(curandStateMtgp32_t)));
        // CHECK: CUDA_CALL(hipMalloc((void **)&d_param, sizeof(mtgp32_kernel_params)));
-        CUDA_CALL(cudaMalloc((void**)&d_param, sizeof(mtgp32_kernel_params)));
+        CUDA_CALL(cudaMalloc((void **)&d_param, sizeof(mtgp32_kernel_params)));
        // curandMakeMTGP32Constants is yet unsupported by HIP
        // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param));
        CURAND_CALL(curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param));
        // curandMakeMTGP32KernelState is yet unsupported by HIP
-        // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213,
+        // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed));
-        // d_param, states_size, seed));
+        CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed));
        CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param,
                                                states_size, seed));
    }
-    ~runner() {
+    ~runner()
    {
        // CHECK: CUDA_CALL(hipFree(states));
        // CHECK: CUDA_CALL(hipFree(d_param));
        CUDA_CALL(cudaFree(states));
        CUDA_CALL(cudaFree(d_param));
    }
-    template <typename T, typename GenerateFunc, typename Extra>
+    template<typename T, typename GenerateFunc, typename Extra>
-    void generate(const size_t blocks, const size_t threads, T* data, const size_t size,
+    void generate(const size_t blocks,
-                  const GenerateFunc& generate_func, const Extra extra) {
+                  const size_t threads,
-        // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(std::min((size_t)200, blocks)),
+                  T * data,
-        // dim3(256), 0, 0, states, data, size, generate_func, extra);
+                  const size_t size,
-        generate_kernel<<<std::min((size_t)200, blocks), 256>>>(states, data, size, generate_func,
+                  const GenerateFunc& generate_func,
-                                                                extra);
+                  const Extra extra)
    {
        // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra);
        generate_kernel<<<std::min((size_t)200, blocks), 256>>>(states, data, size, generate_func, extra);
    }
 };
 // CHECK: void init_kernel(hiprandStateSobol32_t * states,
-template <typename Directions>
+template<typename Directions>
-__global__ void init_kernel(curandStateSobol32_t* states, const Directions directions,
+__global__
-                            const unsigned long long offset) {
+void init_kernel(curandStateSobol32_t * states,
                 const Directions directions,
                 const unsigned long long offset)
 {
    const unsigned int dimension = blockIdx.y;
    const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x;
    // CHECK: hiprandStateSobol32_t state;
@@ -208,9 +247,14 @@ __global__ void init_kernel(curandStateSobol32_t* states, const Directions direc
 }
 // CHECK: void generate_kernel(hiprandStateSobol32_t * states,
-template <typename T, typename GenerateFunc, typename Extra>
+template<typename T, typename GenerateFunc, typename Extra>
-__global__ void generate_kernel(curandStateSobol32_t* states, T* data, const size_t size,
+__global__
-                                const GenerateFunc& generate_func, const Extra extra) {
+void generate_kernel(curandStateSobol32_t * states,
                     T * data,
                     const size_t size,
                     const GenerateFunc& generate_func,
                     const Extra extra)
 {
    const unsigned int dimension = blockIdx.y;
    const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int stride = gridDim.x * blockDim.x;
@@ -218,7 +262,8 @@ __global__ void generate_kernel(curandStateSobol32_t* states, T* data, const siz
    curandStateSobol32_t state = states[gridDim.x * blockDim.x * dimension + state_id];
    const unsigned int offset = dimension * size;
    unsigned int index = state_id;
-    while (index < size) {
+    while(index < size)
    {
        data[offset + index] = generate_func(&state, extra);
        skipahead(stride - 1, &state);
        index += stride;
@@ -229,39 +274,39 @@ __global__ void generate_kernel(curandStateSobol32_t* states, T* data, const siz
 }
 // CHECK: struct runner<hiprandStateSobol32_t>
-template <>
+template<>
-struct runner<curandStateSobol32_t> {
+struct runner<curandStateSobol32_t>
 {
    // CHECK: hiprandStateSobol32_t * states;
-    curandStateSobol32_t* states;
+    curandStateSobol32_t * states;
    size_t dimensions;
-    runner(const size_t dimensions, const size_t blocks, const size_t threads,
+    runner(const size_t dimensions,
-           const unsigned long long seed, const unsigned long long offset) {
+           const size_t blocks,
           const size_t threads,
           const unsigned long long seed,
           const unsigned long long offset)
    {
        this->dimensions = dimensions;
-        // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size *
+        // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(hiprandStateSobol32_t)));
        // sizeof(hiprandStateSobol32_t)));
        const size_t states_size = blocks * threads * dimensions;
-        CUDA_CALL(cudaMalloc((void**)&states, states_size * sizeof(curandStateSobol32_t)));
+        CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateSobol32_t)));
        // CHECK: hiprandDirectionVectors32_t * directions;
-        curandDirectionVectors32_t* directions;
+        curandDirectionVectors32_t * directions;
        // CHECK: const size_t size = dimensions * sizeof(hiprandDirectionVectors32_t);
        const size_t size = dimensions * sizeof(curandDirectionVectors32_t);
        // CHECK: CUDA_CALL(hipMalloc((void **)&directions, size));
-        CUDA_CALL(cudaMalloc((void**)&directions, size));
+        CUDA_CALL(cudaMalloc((void **)&directions, size));
        // CHECK: hiprandDirectionVectors32_t * h_directions;
-        curandDirectionVectors32_t* h_directions;
+        curandDirectionVectors32_t * h_directions;
-        // hiprandGetDirectionVectors32 and HIPRAND_DIRECTION_VECTORS_32_JOEKUO6 (of
+        // hiprandGetDirectionVectors32 and HIPRAND_DIRECTION_VECTORS_32_JOEKUO6 (of hiprandDirectionVectorSet_t) are yet unsupported by HIP
-        // hiprandDirectionVectorSet_t) are yet unsupported by HIP CHECK-NOT:
+        // CHECK-NOT: CURAND_CALL(hiprandGetDirectionVectors32(&h_directions, HIPRAND_DIRECTION_VECTORS_32_JOEKUO6));
-        // CURAND_CALL(hiprandGetDirectionVectors32(&h_directions,
+        CURAND_CALL(curandGetDirectionVectors32(&h_directions, CURAND_DIRECTION_VECTORS_32_JOEKUO6));
        // HIPRAND_DIRECTION_VECTORS_32_JOEKUO6));
        CURAND_CALL(
            curandGetDirectionVectors32(&h_directions, CURAND_DIRECTION_VECTORS_32_JOEKUO6));
        // CHECK: CUDA_CALL(hipMemcpy(directions, h_directions, size, hipMemcpyHostToDevice));
        CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice));
        const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
-        // CHECK: hipLaunchKernelGGL(init_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads),
+        // CHECK: hipLaunchKernelGGL(init_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, directions, offset);
        // 0, 0, states, directions, offset);
        init_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, directions, offset);
        // CHECK: CUDA_CALL(hipPeekAtLastError());
        // CHECK: CUDA_CALL(hipDeviceSynchronize());
@@ -271,25 +316,31 @@ struct runner<curandStateSobol32_t> {
        CUDA_CALL(cudaFree(directions));
    }
-    ~runner() {
+    ~runner()
    {
        // CHECK: CUDA_CALL(hipFree(states));
        CUDA_CALL(cudaFree(states));
    }
-    template <typename T, typename GenerateFunc, typename Extra>
+    template<typename T, typename GenerateFunc, typename Extra>
-    void generate(const size_t blocks, const size_t threads, T* data, const size_t size,
+    void generate(const size_t blocks,
-                  const GenerateFunc& generate_func, const Extra extra) {
+                  const size_t threads,
                  T * data,
                  const size_t size,
                  const GenerateFunc& generate_func,
                  const Extra extra)
    {
        const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
-        // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(dim3(blocks_x, dimensions)),
+        // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
-        // dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
+        generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra);
        generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions,
                                                                 generate_func, extra);
    }
 };
-template <typename T, typename GeneratorState, typename GenerateFunc, typename Extra>
+template<typename T, typename GeneratorState, typename GenerateFunc, typename Extra>
-void run_benchmark(const cli::Parser& parser, const GenerateFunc& generate_func,
+void run_benchmark(const cli::Parser& parser,
-                   const Extra extra) {
+                   const GenerateFunc& generate_func,
                   const Extra extra)
 {
    const size_t size = parser.get<size_t>("size");
    const size_t dimensions = parser.get<size_t>("dimensions");
    const size_t trials = parser.get<size_t>("trials");
@@ -297,14 +348,15 @@ void run_benchmark(const cli::Parser& parser, const GenerateFunc& generate_func,
    const size_t blocks = parser.get<size_t>("blocks");
    const size_t threads = parser.get<size_t>("threads");
-    T* data;
+    T * data;
    // CHECK: CUDA_CALL(hipMalloc((void **)&data, size * sizeof(T)));
-    CUDA_CALL(cudaMalloc((void**)&data, size * sizeof(T)));
+    CUDA_CALL(cudaMalloc((void **)&data, size * sizeof(T)));
    runner<GeneratorState> r(dimensions, blocks, threads, 12345ULL, 6789ULL);
    // Warm-up
-    for (size_t i = 0; i < 5; i++) {
+    for (size_t i = 0; i < 5; i++)
    {
        r.generate(blocks, threads, data, size, generate_func, extra);
        // CHECK: CUDA_CALL(hipPeekAtLastError());
        // CHECK: CUDA_CALL(hipDeviceSynchronize());
@@ -316,7 +368,8 @@ void run_benchmark(const cli::Parser& parser, const GenerateFunc& generate_func,
    // Measurement
    auto start = std::chrono::high_resolution_clock::now();
-    for (size_t i = 0; i < trials; i++) {
+    for (size_t i = 0; i < trials; i++)
    {
        r.generate(blocks, threads, data, size, generate_func, extra);
    }
    // CHECK: CUDA_CALL(hipPeekAtLastError());
@@ -326,132 +379,147 @@ void run_benchmark(const cli::Parser& parser, const GenerateFunc& generate_func,
    auto end = std::chrono::high_resolution_clock::now();
    std::chrono::duration<double, std::milli> elapsed = end - start;
-    std::cout << std::fixed << std::setprecision(3) << "      "
+    std::cout << std::fixed << std::setprecision(3)
-              << "Throughput = " << std::setw(8)
+              << "      "
-              << (trials * size * sizeof(T)) / (elapsed.count() / 1e3 * (1 << 30))
+              << "Throughput = "
-              << " GB/s, Samples = " << std::setw(8)
+              << std::setw(8) << (trials * size * sizeof(T)) /
-              << (trials * size) / (elapsed.count() / 1e3 * (1 << 30))
+                    (elapsed.count() / 1e3 * (1 << 30))
-              << " GSample/s, AvgTime (1 trial) = " << std::setw(8) << elapsed.count() / trials
+              << " GB/s, Samples = "
-              << " ms, Time (all) = " << std::setw(8) << elapsed.count() << " ms, Size = " << size
+              << std::setw(8) << (trials * size) /
                    (elapsed.count() / 1e3 * (1 << 30))
              << " GSample/s, AvgTime (1 trial) = "
              << std::setw(8) << elapsed.count() / trials
              << " ms, Time (all) = "
              << std::setw(8) << elapsed.count()
              << " ms, Size = " << size
              << std::endl;
    // CHECK: CUDA_CALL(hipFree(data));
    CUDA_CALL(cudaFree(data));
 }
-template <typename GeneratorState>
+template<typename GeneratorState>
-void run_benchmarks(const cli::Parser& parser, const std::string& distribution) {
+void run_benchmarks(const cli::Parser& parser,
-    if (distribution == "uniform-uint") {
+                    const std::string& distribution)
 {
    if (distribution == "uniform-uint")
    {
        // curandStateSobol64_t and curandStateScrambledSobol64_t are yet unsupported by HIP
        // CHECK-NOT: if (!std::is_same<GeneratorState, hiprandStateSobol64_t>::value &&
        // CHECK-NOT: !std::is_same<GeneratorState, hiprandStateScrambledSobol64_t>::value)
        if (!std::is_same<GeneratorState, curandStateSobol64_t>::value &&
-            !std::is_same<GeneratorState, curandStateScrambledSobol64_t>::value) {
+            !std::is_same<GeneratorState, curandStateScrambledSobol64_t>::value)
        {
            run_benchmark<unsigned int, GeneratorState>(parser,
-                                                        [] __device__(GeneratorState * state, int) {
+                [] __device__ (GeneratorState * state, int) {
-                                                            // CHECK: return hiprand(state);
+                    // CHECK: return hiprand(state);
-                                                            return curand(state);
+                    return curand(state);
-                                                        },
+                }, 0
-                                                        0);
+            );
        }
    }
-    if (distribution == "uniform-long-long") {
+    if (distribution == "uniform-long-long")
    {
        // curandStateSobol64_t and curandStateScrambledSobol64_t are yet unsupported by HIP
        // CHECK-NOT: if (!std::is_same<GeneratorState, hiprandStateSobol64_t>::value &&
        // CHECK-NOT: !std::is_same<GeneratorState, hiprandStateScrambledSobol64_t>::value)
        if (std::is_same<GeneratorState, curandStateSobol64_t>::value ||
-            std::is_same<GeneratorState, curandStateScrambledSobol64_t>::value) {
+            std::is_same<GeneratorState, curandStateScrambledSobol64_t>::value)
-            run_benchmark<unsigned long long, GeneratorState>(
+        {
-                parser,
+            run_benchmark<unsigned long long, GeneratorState>(parser,
-                [] __device__(GeneratorState * state, int) {
+                [] __device__ (GeneratorState * state, int) {
                    // CHECK: return hiprand(state);
                    return curand(state);
-                },
+                }, 0
-                0);
+            );
        }
    }
-    if (distribution == "uniform-float") {
+    if (distribution == "uniform-float")
    {
        run_benchmark<float, GeneratorState>(parser,
-                                             [] __device__(GeneratorState * state, int) {
+            [] __device__ (GeneratorState * state, int) {
-                                                 // CHECK: return hiprand_uniform(state);
+                // CHECK: return hiprand_uniform(state);
-                                                 return curand_uniform(state);
+                return curand_uniform(state);
-                                             },
+            }, 0
-                                             0);
+        );
    }
-    if (distribution == "uniform-double") {
+    if (distribution == "uniform-double")
    {
        run_benchmark<double, GeneratorState>(parser,
-                                              [] __device__(GeneratorState * state, int) {
+            [] __device__ (GeneratorState * state, int) {
-                                                  // CHECK: return hiprand_uniform_double(state);
+                // CHECK: return hiprand_uniform_double(state);
-                                                  return curand_uniform_double(state);
+                return curand_uniform_double(state);
-                                              },
+            }, 0
-                                              0);
+        );
    }
-    if (distribution == "normal-float") {
+    if (distribution == "normal-float")
    {
        run_benchmark<float, GeneratorState>(parser,
-                                             [] __device__(GeneratorState * state, int) {
+            [] __device__ (GeneratorState * state, int) {
-                                                 // CHECK: return hiprand_normal(state);
+                // CHECK: return hiprand_normal(state);
-                                                 return curand_normal(state);
+                return curand_normal(state);
-                                             },
+            }, 0
-                                             0);
+        );
    }
-    if (distribution == "normal-double") {
+    if (distribution == "normal-double")
    {
        run_benchmark<double, GeneratorState>(parser,
-                                              [] __device__(GeneratorState * state, int) {
+            [] __device__ (GeneratorState * state, int) {
-                                                  // CHECK: return hiprand_normal_double(state);
+                // CHECK: return hiprand_normal_double(state);
-                                                  return curand_normal_double(state);
+                return curand_normal_double(state);
-                                              },
+            }, 0
-                                              0);
+        );
    }
-    if (distribution == "log-normal-float") {
+    if (distribution == "log-normal-float")
    {
        run_benchmark<float, GeneratorState>(parser,
-                                             [] __device__(GeneratorState * state, int) {
+            [] __device__ (GeneratorState * state, int) {
-                                                 // CHECK: return hiprand_log_normal(state,
+                // CHECK: return hiprand_log_normal(state, 0.0f, 1.0f);
-                                                 // 0.0f, 1.0f);
+                return curand_log_normal(state, 0.0f, 1.0f);
-                                                 return curand_log_normal(state, 0.0f, 1.0f);
+            }, 0
-                                             },
+        );
                                             0);
    }
-    if (distribution == "log-normal-double") {
+    if (distribution == "log-normal-double")
    {
        run_benchmark<double, GeneratorState>(parser,
-                                              [] __device__(GeneratorState * state, int) {
+            [] __device__ (GeneratorState * state, int) {
-                                                  // CHECK: return hiprand_log_normal_double(state,
+                // CHECK: return hiprand_log_normal_double(state, 0.0, 1.0);
-                                                  // 0.0, 1.0);
+                return curand_log_normal_double(state, 0.0, 1.0);
-                                                  return curand_log_normal_double(state, 0.0, 1.0);
+            }, 0
-                                              },
+        );
                                              0);
    }
-    if (distribution == "poisson") {
+    if (distribution == "poisson")
    {
        const auto lambdas = parser.get<std::vector<double>>("lambda");
-        for (double lambda : lambdas) {
+        for (double lambda : lambdas)
-            std::cout << "    "
+        {
-                      << "lambda " << std::fixed << std::setprecision(1) << lambda << std::endl;
+            std::cout << "    " << "lambda "
-            run_benchmark<unsigned int, GeneratorState>(
+                 << std::fixed << std::setprecision(1) << lambda << std::endl;
-                parser,
+            run_benchmark<unsigned int, GeneratorState>(parser,
-                [] __device__(GeneratorState * state, double lambda) {
+                [] __device__ (GeneratorState * state, double lambda) {
                    // CHECK: return hiprand_poisson(state, lambda);
                    return curand_poisson(state, lambda);
-                },
+                }, lambda
-                lambda);
+            );
        }
    }
-    if (distribution == "discrete-poisson") {
+    if (distribution == "discrete-poisson")
    {
        const auto lambdas = parser.get<std::vector<double>>("lambda");
-        for (double lambda : lambdas) {
+        for (double lambda : lambdas)
-            std::cout << "    "
+        {
-                      << "lambda " << std::fixed << std::setprecision(1) << lambda << std::endl;
+            std::cout << "    " << "lambda "
                 << std::fixed << std::setprecision(1) << lambda << std::endl;
            // CHECK: hiprandDiscreteDistribution_t discrete_distribution;
            curandDiscreteDistribution_t discrete_distribution;
            // CHECK: CURAND_CALL(hiprandCreatePoissonDistribution(lambda, &discrete_distribution));
            CURAND_CALL(curandCreatePoissonDistribution(lambda, &discrete_distribution));
-            run_benchmark<unsigned int, GeneratorState>(
+            run_benchmark<unsigned int, GeneratorState>(parser,
-                parser,
+                // CHECK: [] __device__ (GeneratorState * state, hiprandDiscreteDistribution_t discrete_distribution) {
-                // CHECK: [] __device__ (GeneratorState * state, hiprandDiscreteDistribution_t
+                [] __device__ (GeneratorState * state, curandDiscreteDistribution_t discrete_distribution) {
                // discrete_distribution) {
                [] __device__(GeneratorState * state,
                              curandDiscreteDistribution_t discrete_distribution) {
                    // CHECK: return hiprand_discrete(state, discrete_distribution);
                    return curand_discrete(state, discrete_distribution);
-                },
+                }, discrete_distribution
-                discrete_distribution);
+            );
            // CHECK: CURAND_CALL(hiprandDestroyDistribution(discrete_distribution));
            CURAND_CALL(curandDestroyDistribution(discrete_distribution));
        }
@@ -459,9 +527,12 @@ void run_benchmarks(const cli::Parser& parser, const std::string& distribution)
 }
 const std::vector<std::string> all_engines = {
-    "xorwow", "mrg32k3a", "mtgp32",
+    "xorwow",
    "mrg32k3a",
    "mtgp32",
    // "mt19937",
-    "philox", "sobol32",
+    "philox",
    "sobol32",
    // "scrambled_sobol32",
    // "sobol64",
    // "scrambled_sobol64",
@@ -480,42 +551,50 @@ const std::vector<std::string> all_distributions = {
    "discrete-poisson",
 };
-int main(int argc, char* argv[]) {
+int main(int argc, char *argv[])
 {
    cli::Parser parser(argc, argv);
    const std::string distribution_desc =
        "space-separated list of distributions:" +
        std::accumulate(all_distributions.begin(), all_distributions.end(), std::string(),
-                        [](std::string a, std::string b) { return a + "\n      " + b; }) +
+            [](std::string a, std::string b) {
                return a + "\n      " + b;
            }
        ) +
        "\n      or all";
    const std::string engine_desc =
        "space-separated list of random number engines:" +
        std::accumulate(all_engines.begin(), all_engines.end(), std::string(),
-                        [](std::string a, std::string b) { return a + "\n      " + b; }) +
+            [](std::string a, std::string b) {
                return a + "\n      " + b;
            }
        ) +
        "\n      or all";
    parser.set_optional<size_t>("size", "size", DEFAULT_RAND_N, "number of values");
-    parser.set_optional<size_t>("dimensions", "dimensions", 1,
+    parser.set_optional<size_t>("dimensions", "dimensions", 1, "number of dimensions of quasi-random values");
                                "number of dimensions of quasi-random values");
    parser.set_optional<size_t>("trials", "trials", 20, "number of trials");
    parser.set_optional<size_t>("blocks", "blocks", 256, "number of blocks");
    parser.set_optional<size_t>("threads", "threads", 256, "number of threads in each block");
-    parser.set_optional<std::vector<std::string>>("dis", "dis", {"uniform-uint"},
+    parser.set_optional<std::vector<std::string>>("dis", "dis", {"uniform-uint"}, distribution_desc.c_str());
-                                                  distribution_desc.c_str());
+    parser.set_optional<std::vector<std::string>>("engine", "engine", {"philox"}, engine_desc.c_str());
-    parser.set_optional<std::vector<std::string>>("engine", "engine", {"philox"},
+    parser.set_optional<std::vector<double>>("lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution");
                                                  engine_desc.c_str());
    parser.set_optional<std::vector<double>>(
        "lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution");
    parser.run_and_exit_if_error();
    std::vector<std::string> engines;
    {
        auto es = parser.get<std::vector<std::string>>("engine");
-        if (std::find(es.begin(), es.end(), "all") != es.end()) {
+        if (std::find(es.begin(), es.end(), "all") != es.end())
        {
            engines = all_engines;
-        } else {
+        }
-            for (auto e : all_engines) {
+        else
-                if (std::find(es.begin(), es.end(), e) != es.end()) engines.push_back(e);
+        {
            for (auto e : all_engines)
            {
                if (std::find(es.begin(), es.end(), e) != es.end())
                    engines.push_back(e);
            }
        }
    }
@@ -523,11 +602,16 @@ int main(int argc, char* argv[]) {
    std::vector<std::string> distributions;
    {
        auto ds = parser.get<std::vector<std::string>>("dis");
-        if (std::find(ds.begin(), ds.end(), "all") != ds.end()) {
+        if (std::find(ds.begin(), ds.end(), "all") != ds.end())
        {
            distributions = all_distributions;
-        } else {
+        }
-            for (auto d : all_distributions) {
+        else
-                if (std::find(ds.begin(), ds.end(), d) != ds.end()) distributions.push_back(d);
+        {
            for (auto d : all_distributions)
            {
                if (std::find(ds.begin(), ds.end(), d) != ds.end())
                    distributions.push_back(d);
            }
        }
    }
@@ -552,24 +636,35 @@ int main(int argc, char* argv[]) {
    std::cout << "Device: " << props.name;
    std::cout << std::endl << std::endl;
-    for (auto engine : engines) {
+    for (auto engine : engines)
    {
        std::cout << engine << ":" << std::endl;
-        for (auto distribution : distributions) {
+        for (auto distribution : distributions)
        {
            std::cout << "  " << distribution << ":" << std::endl;
            const std::string plot_name = engine + "-" + distribution;
-            if (engine == "xorwow") {
+            if (engine == "xorwow")
            {
                // CHECK: run_benchmarks<hiprandStateXORWOW_t>(parser, distribution);
                run_benchmarks<curandStateXORWOW_t>(parser, distribution);
-            } else if (engine == "mrg32k3a") {
+            }
            else if (engine == "mrg32k3a")
            {
                // CHECK: run_benchmarks<hiprandStateMRG32k3a_t>(parser, distribution);
                run_benchmarks<curandStateMRG32k3a_t>(parser, distribution);
-            } else if (engine == "philox") {
+            }
            else if (engine == "philox")
            {
                // CHECK: run_benchmarks<hiprandStatePhilox4_32_10_t>(parser, distribution);
                run_benchmarks<curandStatePhilox4_32_10_t>(parser, distribution);
-            } else if (engine == "sobol32") {
+            }
            else if (engine == "sobol32")
            {
                // CHECK: run_benchmarks<hiprandStateSobol32_t>(parser, distribution);
                run_benchmarks<curandStateSobol32_t>(parser, distribution);
-            } else if (engine == "mtgp32") {
+            }
            else if (engine == "mtgp32")
            {
                // CHECK: run_benchmarks<hiprandStateMtgp32_t>(parser, distribution);
                run_benchmarks<curandStateMtgp32_t>(parser, distribution);
            }
@@ -57,6 +57,8 @@ else:
    run_test_ext = ".sh"
    clang_args += " -isystem'%s'/samples/common/inc"
-config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root)))
+clang_args += " -I'%s'/include"
 config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root, config.cuda_dnn_root)))
 config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"'))
 config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"'))
@@ -4,6 +4,7 @@ import os
 config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
 config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
 config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
 config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@"
 if sys.platform in ['win32']:
    config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@"
    if not config.cuda_sdk_root or config.cuda_sdk_root == "CUDA_SDK_ROOT_DIR-NOTFOUND":