Integration Testing (#211)
* Add external/cereal submodule
- used for integration testing
* Update lib/common/container/small_vector.hpp
- documentation notes
* Update tests/apps
- update transpose app (fix build)
- add reproducible-runtime app
* Update include/rocprofiler/fwd.h
- rocprofiler_service_callback_phase_t -> rocprofiler_callback_phase_t
* Update PTL submodule
- fix for task group: submitting tasks from different thread
* Update lib/rocprofiler/hsa/queue.cpp
- CHECK_NOTNULL(_buffer)
* Update lib/rocprofiler/hsa/hsa.cpp
- use buffer::get_buffer instead of manually looking for buffer
* Update lib/rocprofiler/internal_threading.cpp
- use buffer::get_buffer instead of manually looking for buffer
* Update lib/rocprofiler/buffer.cpp
- offset the buffer id
- properly handle rocprofiler_create_buffer reusing rocprofiler_buffer_id_t on a different context
* Update tests
- kernel tracing library for integration testing
* Add cereal submodule
* Update lib/rocprofiler/registration.*
- OnUnload
- Support ROCP_TOOL_LIBRARIES for python usage
- improve finalize function
- remove calling hsa_shut_down in finalize function
* Update lib/rocprofiler/buffer.*
- allocate_buffer sets the buffer id value
- expose (internally) is_valid_buffer_id
- update test
* Update tests/kernel-tracing
- installation
- better organization of JSON groups
- improved messaging
* Update lib/rocprofiler/registration.cpp
- add workaround for hsa-runtime supporting rocprofiler-register
* Update tests/kernel-tracing/kernel-tracing.cpp
- fix memory leaks
* cereal support for minimal JSON
- update cereal submodule to rocprofiler branch
- change REPO_BRANCH in rocprofiler_checkout_git_submodule for cereal
- update tests/kernel-tracing/kernel-tracing.cpp
- use minimal json
- slight tweak putting giving contexts name in storing name + context pointer pair in map
* Update tests/kernel-tracing/kernel-tracing.cpp
- support runtime selection of contexts via KERNEL_TRACING_CONTEXTS environment variable
* Update tests
- tests/CMakeLists.txt
- find_package(Python3 REQUIRED)
- tests/kernel-tracing
- pytest validation
* Update CI workflow
- install pytest
- add checks for test labels
* Update scripts/run-ci.py
- change --coverage options
- replace 'unittests' with 'tests'
- replace test label regex '-L unittests' with '-L tests'
* Update requirements.txt
- this is now an empty file since none of the packages are required for this repo
[ROCm/rocprofiler-sdk commit: cf5e4b4b1b]
This commit is contained in:
committed by
GitHub
parent
59b40edeab
commit
2918251d9c
@@ -1,24 +1,24 @@
|
||||
/*
|
||||
Copyright (c) 2015-2020 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in
|
||||
// all copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
// THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
@@ -66,7 +66,7 @@ verify(int* in, int* out, int M, int N);
|
||||
} // namespace
|
||||
|
||||
__global__ void
|
||||
transpose_a(int* in, int* out, int M, int N);
|
||||
transpose(int* in, int* out, int M, int N);
|
||||
|
||||
void
|
||||
run(int rank, int tid, hipStream_t stream, int argc, char** argv);
|
||||
@@ -148,7 +148,7 @@ main(int argc, char** argv)
|
||||
}
|
||||
|
||||
__global__ void
|
||||
transpose_a(int* in, int* out, int M, int N)
|
||||
transpose(int* in, int* out, int M, int N)
|
||||
{
|
||||
__shared__ int tile[shared_mem_tile_dim][shared_mem_tile_dim];
|
||||
|
||||
@@ -193,12 +193,12 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
dim3 grid(M / 32, N / 32, 1);
|
||||
dim3 block(32, 32, 1); // transpose_a
|
||||
dim3 block(32, 32, 1); // transpose
|
||||
|
||||
auto t1 = std::chrono::high_resolution_clock::now();
|
||||
for(size_t i = 0; i < nitr; ++i)
|
||||
{
|
||||
transpose_a<<<grid, block, 0, stream>>>(in, out, M, N);
|
||||
transpose<<<grid, block, 0, stream>>>(in, out, M, N);
|
||||
check_hip_error();
|
||||
if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user