diff --git a/Jenkinsfile b/Jenkinsfile
index 1257bfb0b4..812b68633a 100644
--- a/Jenkinsfile
+++ b/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')
{
diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt
index 56de7c7904..964e30d66a 100644
--- a/hipify-clang/CMakeLists.txt
+++ b/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}")
diff --git a/hipify-clang/README.md b/hipify-clang/README.md
index 2df3a73bff..8ef9883b55 100644
--- a/hipify-clang/README.md
+++ b/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`
## 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.
diff --git a/hipify-clang/src/LLVMCompat.cpp b/hipify-clang/src/LLVMCompat.cpp
index f96bd86bb0..22a0b56d64 100644
--- a/hipify-clang/src/LLVMCompat.cpp
+++ b/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 &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
diff --git a/hipify-clang/src/LLVMCompat.h b/hipify-clang/src/LLVMCompat.h
index 4d6edcc9f8..98ff66fb41 100644
--- a/hipify-clang/src/LLVMCompat.h
+++ b/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 &output,
+ bool expand_tilde = false);
+
} // namespace llcompat
diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp
index 06981a9d70..2a3b86f204 100644
--- a/hipify-clang/src/main.cpp
+++ b/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(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 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
diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp
index e678f25aa2..2fbda48629 100644
--- a/include/hip/hcc_detail/functional_grid_launch.hpp
+++ b/include/hip/hcc_detail/functional_grid_launch.hpp
@@ -33,6 +33,7 @@ THE SOFTWARE.
#include
#include
+#include
#include
#include
#include
@@ -56,7 +57,9 @@ template <
typename... Ts,
typename std::enable_if::type* = nullptr>
inline std::vector make_kernarg(
- std::vector kernarg, const std::tuple&) {
+ const std::tuple&,
+ const std::vector>&,
+ std::vector kernarg) {
return kernarg;
}
@@ -65,7 +68,9 @@ template <
typename... Ts,
typename std::enable_if::type* = nullptr>
inline std::vector make_kernarg(
- std::vector kernarg, const std::tuple& formals) {
+ const std::tuple& formals,
+ const std::vector>& size_align,
+ std::vector kernarg) {
using T = typename std::tuple_element>::type;
static_assert(
@@ -80,24 +85,48 @@ inline std::vector 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(formals)};
+ std::memcpy(
+ kernarg.data() + kernarg.size() - size_align[n].first,
+ &std::get(formals),
+ size_align[n].first);
- return make_kernarg(std::move(kernarg), formals);
+ return make_kernarg(formals, size_align, std::move(kernarg));
}
template
inline std::vector make_kernarg(
- void (*)(Formals...), std::tuple actuals) {
+ void (*kernel)(Formals...), std::tuple 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(kernel));
+ if (it == function_names().cend()) {
+ it =
+ function_names(true).find(reinterpret_cast(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 to_formals{std::move(actuals)};
std::vector 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,
diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 97d6320f7f..0d526ca25d 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/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));
/**
diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp
index bdb87b3509..da13c7c3db 100644
--- a/include/hip/hcc_detail/program_state.hpp
+++ b/include/hip/hcc_detail/program_state.hpp
@@ -99,6 +99,9 @@ const std::unordered_map& function_names(bool rebuild = false);
std::unordered_map& globals(bool rebuild = false);
+const std::unordered_map<
+ std::string, std::vector>>&
+ kernargs(bool rebuild = false);
hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
hsa_agent_t agent);
diff --git a/src/program_state.cpp b/src/program_state.cpp
index 88cdeeb404..bb906b0ad9 100644
--- a/src/program_state.cpp
+++ b/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>& 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>>& 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(desc), static_cast(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>>& fu
// created previously
function_names(rebuild);
+ kernargs(rebuild);
kernels(rebuild);
globals(rebuild);
}
@@ -501,6 +586,45 @@ unordered_map& globals(bool rebuild) {
return r;
}
+const unordered_map>>& kernargs(
+ bool rebuild) {
+ static unordered_map>> 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 lck{mtx};
+
+ tmp.insert(r.cbegin(), r.cend()); // Should use merge in C++17.
+ }
+
+ build_map(tmp);
+
+ lock_guard 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;
diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg
index d0328c06c2..76a811031a 100644
--- a/tests/hipify-clang/lit.cfg
+++ b/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"
diff --git a/tests/hipify-clang/lit.site.cfg.in b/tests/hipify-clang/lit.site.cfg.in
index f7302252a8..f93c206be4 100644
--- a/tests/hipify-clang/lit.site.cfg.in
+++ b/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@"
diff --git a/tests/hipify-clang/unit_tests/samples/cudaRegister.cu b/tests/hipify-clang/unit_tests/samples/cudaRegister.cu
index 43b4345337..11192c452e 100644
--- a/tests/hipify-clang/unit_tests/samples/cudaRegister.cu
+++ b/tests/hipify-clang/unit_tests/samples/cudaRegister.cu
@@ -22,12 +22,6 @@ THE SOFTWARE.
#include
#include
#include
-#ifdef _WIN32
-#include
-#define sleep(x) Sleep(x)
-#else
-#include
-#endif
#include
#include
@@ -90,7 +84,6 @@ int main(){
// CHECK: hipLaunchKernelGGL(Inc1, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd);
Inc1<<>>(Ad, Bd);
- sleep(3);
A[0] = -(ITER*1.0f);
std::cout<<"Same cache line before completion: \t"<< A[0]<>>(Ad, Bd);
- sleep(3);
A[0] = -(ITER*1.0f);
std::cout<<"Diff cache line before completion: \t"<>>(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();
}
diff --git a/tests/src/runtimeApi/memory/hipMemset3D.cpp b/tests/src/runtimeApi/memory/hipMemset3D.cpp
index 53e21c93bb..151a86a6c0 100644
--- a/tests/src/runtimeApi/memory/hipMemset3D.cpp
+++ b/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
*/