Merge 'master' into 'amd-master'

Change-Id: I4ea252fc1aab6aeeb31c858ca5fcc2e36d989ba0

Jenkinsfile

@@ -295,52 +295,7 @@ String build_config = 'Release'
 String job_name = env.JOB_NAME.toLowerCase( )
 
 // The following launches 3 builds in parallel: rocm-head, rocm-1.9.x and cuda-9.x
-parallel rocm_1_9:
-{
-  node('hip-rocm')
-  {
-    String hcc_ver = 'rocm-1.9.x'
-    String from_image = 'ci_test_nodes/rocm-1.9.x/ubuntu-16.04:latest'
-    String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video'
-
-    // Checkout source code, dependencies and version files
-    String source_hip_rel = checkout_and_version( hcc_ver )
-
-    // Create/reuse a docker image that represents the hip build environment
-    def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image )
-
-    // Print system information for the log
-    hip_build_image.inside( inside_args )
-    {
-      sh  """#!/usr/bin/env bash
-          set -x
-          /opt/rocm/bin/rocm_agent_enumerator -t ALL
-          /opt/rocm/bin/hcc --version
-        """
-    }
-
-    // Conctruct a binary directory path based on build config
-    String build_hip_rel = build_directory_rel( build_config );
-
-    // Build hip inside of the build environment
-    docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel )
-
-    // Clean docker build image
-    docker_clean_images( 'hip', docker_build_image_name( ) )
-
-    // After a successful build, upload a docker image of the results
-    /*
-    String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel )
-    if( params.push_image_to_docker_hub )
-    {
-      docker_upload_dockerhub( job_name, hip_image_name, 'rocm' )
-      docker_clean_images( 'rocm', hip_image_name )
-    }
-    docker_clean_images( job_name, hip_image_name )
-    */
-  }
-},
-rocm_2_0:
+parallel rocm_2_0:
 {
   node('hip-rocm')
   {

hipify-clang/CMakeLists.txt

@@ -1,6 +1,11 @@
-cmake_minimum_required(VERSION 2.8.12)
+cmake_minimum_required(VERSION 3.12.3)
 project(hipify-clang)
 
+if (MSVC AND MSVC_VERSION VERSION_LESS "1900")
+    message(SEND_ERROR "hipify-clang could be built by Visual Studio 14 2015 or higher.")
+    return()
+endif()
+
 find_package(LLVM REQUIRED)
 message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}:")
 message(STATUS "   - CMake module path: ${LLVM_CMAKE_DIR}")

hipify-clang/README.md

@@ -34,7 +34,7 @@
 `hipify-clang` requires:
 1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release: 6.0.1 (linux and windows).
 
-2. CUDA at least version 7.5, latest supported release is 9.0.
+2. CUDA at least version 7.0, latest supported version is 9.0.
 
 | **LLVM release version** | **CUDA latest supported version** | **Comments** |
 |:------------------------:|:---------------------------------:|:------------:|
@@ -55,7 +55,7 @@
 In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
 
 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.10/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.1\dist`
+[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.1\dist`
 
 ## <a name="build-and-install"></a> Build and install
 
@@ -175,7 +175,7 @@ LLVM 5.0.0 - 6.0.1, CUDA 8.0, cudnn-8.0
 
 Build system for the above configurations:
 
-Python 2.7 (min), cmake 3.5.2 (min), GNU C/C++ 5.4.0 (min).
+Python 2.7 (min), cmake 3.12.3 (min), GNU C/C++ 5.4.0 (min).
 
 Here is an example of building `hipify-clang` with testing support on `Ubuntu 16.04`:
 
@@ -278,7 +278,7 @@ LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-8.0
 
 Build system for the above configurations:
 
-Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min).
+Python 3.6 (min), cmake 3.12.3 (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`:
 
@@ -321,9 +321,7 @@ For example:
 ./hipify-clang \
   square.cu \
   -- \
-  -x cuda \
   --cuda-path=/usr/local/cuda-8.0 \
-  --cuda-gpu-arch=sm_50 \
   -isystem /usr/local/cuda-8.0/samples/common/inc
 ```
 
@@ -339,5 +337,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-2018 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2014-2019 Advanced Micro Devices, Inc. All rights reserved.
 

hipify-clang/src/LLVMCompat.cpp

@@ -94,4 +94,13 @@ clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc) {
 #endif
 }
 
+std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
+                          bool expand_tilde) {
+#if LLVM_VERSION_MAJOR < 5
+  return sys::fs::make_absolute(path, output);
+#else
+  return sys::fs::real_path(path, output, expand_tilde);
+#endif
+}
+
 } // namespace llcompat

hipify-clang/src/LLVMCompat.h

@@ -78,4 +78,7 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
                                   size_t len,
                                   bool DisableMacroExpansion);
 
+std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
+                          bool expand_tilde = false);
+
 } // namespace llcompat

hipify-clang/src/main.cpp

@@ -46,7 +46,7 @@ std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC,
     return sDir;
   }
   SmallString<256> dirAbsPath;
-  EC = sys::fs::real_path(sDir, dirAbsPath, true);
+  EC = llcompat::real_path(sDir, dirAbsPath, true);
   if (!EC && sys::fs::is_regular_file(dirAbsPath)) {
     llvm::errs() << "\n" << sHipify << sError << sDir << " is not a directory\n";
     EC = std::error_code(static_cast<int>(std::errc::not_a_directory), std::generic_category());
@@ -58,7 +58,7 @@ std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC,
       llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
       return "";
     }
-    EC = sys::fs::real_path(sDir, dirAbsPath, true);
+    EC = llcompat::real_path(sDir, dirAbsPath, true);
     if (EC) {
       llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
       return "";
@@ -137,7 +137,7 @@ int main(int argc, const char **argv) {
     // Create a copy of the file to work on. When we're done, we'll move this onto the
     // output (which may mean overwriting the input, if we're in-place).
     // Should we fail for some reason, we'll just leak this file and not corrupt the input.
-    EC = sys::fs::real_path(src, sourceAbsPath, true);
+    EC = llcompat::real_path(src, sourceAbsPath, true);
     if (EC) {
       llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << src << "\n";
       Result = 1;
@@ -171,9 +171,20 @@ int main(int argc, const char **argv) {
     ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
     std::string sInclude = "-I" + sys::path::parent_path(sourceAbsPath).str();
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN));
+    // Includes for clang's CUDA wrappers for using by packaged hipify-clang
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN));
     // Ensure at least c++11 is used.
-    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-std=c++11", ct::ArgumentInsertPosition::BEGIN));
+    std::string stdCpp = "-std=c++11";
+#if defined(_MSC_VER)
+    stdCpp = "-std=c++14";
+#endif
+    Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN));
 #if defined(HIPIFY_CLANG_RES)
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
 #endif

include/hip/hcc_detail/functional_grid_launch.hpp

@@ -33,6 +33,7 @@ THE SOFTWARE.
 
 #include <cstddef>
 #include <cstdint>
+#include <cstring>
 #include <functional>
 #include <iostream>
 #include <mutex>
@@ -56,7 +57,9 @@ template <
     typename... Ts,
     typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
 inline std::vector<std::uint8_t> make_kernarg(
-    std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>&) {
+    const std::tuple<Ts...>&,
+    const std::vector<std::pair<std::size_t, std::size_t>>&,
+    std::vector<std::uint8_t> kernarg) {
     return kernarg;
 }
 
@@ -65,7 +68,9 @@ template <
     typename... Ts,
     typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
 inline std::vector<std::uint8_t> make_kernarg(
-    std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>& formals) {
+    const std::tuple<Ts...>& formals,
+    const std::vector<std::pair<std::size_t, std::size_t>>& size_align,
+    std::vector<std::uint8_t> kernarg) {
     using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
 
     static_assert(
@@ -80,24 +85,48 @@ inline std::vector<std::uint8_t> make_kernarg(
     #endif
 
     kernarg.resize(round_up_to_next_multiple_nonnegative(
-        kernarg.size(), alignof(T)) + sizeof(T));
+        kernarg.size(), size_align[n].second) + size_align[n].first);
 
-    new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get<n>(formals)};
+    std::memcpy(
+        kernarg.data() + kernarg.size() - size_align[n].first,
+        &std::get<n>(formals),
+        size_align[n].first);
 
-    return make_kernarg<n + 1>(std::move(kernarg), formals);
+    return make_kernarg<n + 1>(formals, size_align, std::move(kernarg));
 }
 
 template <typename... Formals, typename... Actuals>
 inline std::vector<std::uint8_t> make_kernarg(
-    void (*)(Formals...), std::tuple<Actuals...> actuals) {
+    void (*kernel)(Formals...), std::tuple<Actuals...> actuals) {
     static_assert(sizeof...(Formals) == sizeof...(Actuals),
         "The count of formal arguments must match the count of actuals.");
 
+    if (sizeof...(Formals) == 0) return {};
+
+    auto it = function_names().find(reinterpret_cast<std::uintptr_t>(kernel));
+    if (it == function_names().cend()) {
+        it =
+            function_names(true).find(reinterpret_cast<std::uintptr_t>(kernel));
+        if (it == function_names().cend()) {
+            throw std::runtime_error{"Undefined __global__ function."};
+        }
+    }
+
+    auto it1 = kernargs().find(it->second);
+    if (it1 == kernargs().end()) {
+        it1 = kernargs(true).find(it->second);
+
+        if (it1 == kernargs().end()) {
+            throw std::runtime_error{
+                "Missing metadata for __global__ function: " + it->second};
+        }
+    }
+
     std::tuple<Formals...> to_formals{std::move(actuals)};
     std::vector<std::uint8_t> kernarg;
     kernarg.reserve(sizeof(to_formals));
 
-    return make_kernarg<0>(std::move(kernarg), to_formals);
+    return make_kernarg<0>(to_formals, it1->second, std::move(kernarg));
 }
 
 void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks,

include/hip/hcc_detail/hip_runtime_api.h

@@ -2573,7 +2573,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr);
  * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
  *
  */
-hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream);
+hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem __dparm(0), hipStream_t stream __dparm(0));
 
 
 /**

include/hip/hcc_detail/program_state.hpp

@@ -99,6 +99,9 @@ const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kern
 functions(bool rebuild = false);
 const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
 std::unordered_map<std::string, void*>& globals(bool rebuild = false);
+const std::unordered_map<
+    std::string, std::vector<std::pair<std::size_t, std::size_t>>>&
+    kernargs(bool rebuild = false);
 
 hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
                                  hsa_agent_t agent);

src/program_state.cpp

@@ -340,6 +340,90 @@ void load_code_object_and_freeze_executable(
         code_readers.push_back(move(tmp));
     }
 }
+
+size_t parse_args(
+    const string& metadata,
+    size_t f,
+    size_t l,
+    vector<pair<size_t, size_t>>& size_align) {
+    if (f == l) return f;
+    if (!size_align.empty()) return l;
+
+    do {
+        static constexpr size_t size_sz{5};
+        f = metadata.find("Size:", f) + size_sz;
+
+        if (l <= f) return f;
+
+        auto size = strtoul(&metadata[f], nullptr, 10);
+
+        static constexpr size_t align_sz{6};
+        f = metadata.find("Align:", f) + align_sz;
+
+        char* l{};
+        auto align = strtoul(&metadata[f], &l, 10);
+
+        f += (l - &metadata[f]) + 1;
+
+        size_align.emplace_back(size, align);
+    } while (true);
+}
+
+void read_kernarg_metadata(
+    elfio& reader,
+    unordered_map<string, vector<pair<size_t, size_t>>>& kernargs)
+{   // TODO: this is inefficient.
+    auto it = find_section_if(
+        reader, [](const section* x) { return x->get_type() == SHT_NOTE; });
+
+    if (!it) return;
+
+    const note_section_accessor acc{reader, it};
+    for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) {
+        ELFIO::Elf_Word type{};
+        string name{};
+        void* desc{};
+        Elf_Word desc_size{};
+
+        acc.get_note(i, type, name, desc, desc_size);
+
+        if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
+
+        string tmp{
+            static_cast<char*>(desc), static_cast<char*>(desc) + desc_size};
+
+        auto dx = tmp.find("Kernels:");
+
+        if (dx == string::npos) continue;
+
+        static constexpr decltype(tmp.size()) kernels_sz{8};
+        dx += kernels_sz;
+
+        do {
+            dx = tmp.find("Name:", dx);
+
+            if (dx == string::npos) break;
+
+            static constexpr decltype(tmp.size()) name_sz{5};
+            dx = tmp.find_first_not_of(" '", dx + name_sz);
+
+            auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx);
+            dx += fn.size();
+
+            auto dx1 = tmp.find("CodeProps", dx);
+            dx = tmp.find("Args:", dx);
+
+            if (dx1 < dx) {
+                dx = dx1;
+                continue;
+            }
+            if (dx == string::npos) break;
+
+            static constexpr decltype(tmp.size()) args_sz{5};
+            dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]);
+        } while (true);
+    }
+}
 }  // namespace
 
 namespace hip_impl {
@@ -454,6 +538,7 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
             // created previously
 
             function_names(rebuild);
+            kernargs(rebuild);
             kernels(rebuild);
             globals(rebuild);
         }
@@ -501,6 +586,45 @@ unordered_map<string, void*>& globals(bool rebuild) {
     return r;
 }
 
+const unordered_map<string, vector<pair<size_t, size_t>>>& kernargs(
+    bool rebuild) {
+    static unordered_map<string, vector<pair<size_t, size_t>>> r;
+    static once_flag f;
+
+    static const auto build_map = [](decltype(r)& x) {
+        for (auto&& isa_blobs : code_object_blobs()) {
+            for (auto&& blob : isa_blobs.second) {
+                stringstream tmp{std::string{blob.cbegin(), blob.cend()}};
+
+                elfio reader;
+                if (!reader.load(tmp)) continue;
+
+                read_kernarg_metadata(reader, x);
+            }
+        }
+    };
+    call_once(f, []() { r.reserve(function_names().size()); build_map(r); });
+
+    if (rebuild) {
+        static mutex mtx;
+        thread_local static decltype(r) tmp;
+
+        {
+            lock_guard<mutex> lck{mtx};
+
+            tmp.insert(r.cbegin(), r.cend()); // Should use merge in C++17.
+        }
+
+        build_map(tmp);
+
+        lock_guard<mutex> lck{mtx};
+
+        r.insert(tmp.cbegin(), tmp.cend());
+    }
+
+    return r;
+}
+
 hsa_executable_t load_executable(const string& file, hsa_executable_t executable,
                                  hsa_agent_t agent) {
     elfio reader;

tests/hipify-clang/lit.cfg

@@ -15,12 +15,16 @@ print("CUDA " + config.cuda_version + " will be used for testing.")
 
 config.excludes = ['cmdparser.hpp']
 
+if config.cuda_version_major == 7 and config.cuda_version_minor == 0:
+    config.excludes.append('headers_test_09.cu')
+if config.cuda_version_major < 8:
+    config.excludes.append('cuSPARSE_02.cu')
 if config.cuda_version_major < 9:
     config.excludes.append('cuSPARSE_04.cu')
     config.excludes.append('cuSPARSE_05.cu')
     config.excludes.append('cuSPARSE_06.cu')
     config.excludes.append('cuSPARSE_07.cu')
-
+    config.excludes.append('benchmark_curand_kernel.cpp')
 if config.cuda_version_major < 10:
     config.excludes.append('cuSPARSE_08.cu')
     config.excludes.append('cuSPARSE_09.cu')
@@ -60,12 +64,12 @@ if obj_root is not None:
     config.environment['PATH'] = path
 
 hipify_path = obj_root
-clang_args = "-x cuda -v --cuda-gpu-arch=sm_30 --cuda-path='%s'"
+clang_args = "-v --cuda-path='%s'"
 
 if sys.platform in ['win32']:
     run_test_ext = ".bat"
     hipify_path += "/" + config.build_type
-    clang_args += " -isystem'%s'/common/inc -std=c++14"
+    clang_args += " -isystem'%s'/common/inc"
 else:
     run_test_ext = ".sh"
     clang_args += " -isystem'%s'/samples/common/inc"

tests/hipify-clang/lit.site.cfg.in

@@ -6,6 +6,7 @@ config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
 config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
 config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@"
 config.cuda_version_major = int("@CUDA_VERSION_MAJOR@")
+config.cuda_version_minor = int("@CUDA_VERSION_MINOR@")
 config.cuda_version = "@CUDA_VERSION@"
 if sys.platform in ['win32']:
     config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@"

tests/hipify-clang/unit_tests/samples/cudaRegister.cu

@@ -22,12 +22,6 @@ THE SOFTWARE.
 #include<cuda.h>
 #include<cuda_runtime.h>
 #include<iostream>
-#ifdef _WIN32
-#include <windows.h>
-#define sleep(x) Sleep(x)
-#else
-#include <unistd.h>
-#endif
 #include<stdio.h>
 #include<malloc.h>
 
@@ -90,7 +84,6 @@ int main(){
 
   // CHECK: hipLaunchKernelGGL(Inc1, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd);
 	Inc1<<<dimGrid, dimBlock>>>(Ad, Bd);
-	sleep(3);
 	A[0] = -(ITER*1.0f);
 	std::cout<<"Same cache line before completion: \t"<< A[0]<<std::endl;
 
@@ -104,7 +97,6 @@ int main(){
 
   // CHECK: hipLaunchKernelGGL(Inc2, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd);
 	Inc2<<<dimGrid, dimBlock>>>(Ad, Bd);
-	sleep(3);
 	A[0] = -(ITER*1.0f);
 	std::cout<<"Diff cache line before completion: \t"<<A[0]<<std::endl;
 

tests/src/kernel/hipGridLaunch.cpp

@@ -74,12 +74,43 @@ int test_gl2(size_t N) {
     return 0;
 }
 
+#if __HIP__
+int test_triple_chevron(size_t N) {
+    size_t Nbytes = N * sizeof(int);
+
+    int *A_d, *B_d, *C_d;
+    int *A_h, *B_h, *C_h;
+
+    HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
+
+
+    unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+
+
+    // Full vadd in one large chunk, to get things started:
+    HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+    HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+
+    vectorADD2<<<dim3(blocks), dim3(threadsPerBlock)>>>(A_d, B_d, C_d, N);
+
+    HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+
+    HIPCHECK(hipDeviceSynchronize());
+
+    HipTest::checkVectorADD(A_h, B_h, C_h, N);
+
+    return 0;
+}
+#endif
 
 int main(int argc, char* argv[]) {
     HipTest::parseStandardArguments(argc, argv, true);
 
     test_gl2(N);
 
+#if __HIP__
+    test_triple_chevron(N);
+#endif
 
     passed();
 }

tests/src/runtimeApi/memory/hipMemset3D.cpp

@@ -24,8 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * //Small copy
- * RUN: %t -N 10    --memsetval 0x42
+ * RUN: %t
  * HIT_END
  */