AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Shared Library Constructor (rocprofv3 deadlock fix) (#599)

Browse Source 
* Moved tests/apps to tests/bin

* Renamed cmake project in tests/bin

* Update samples

- Use ROCPROFILER_DEFAULT_FAIL_REGEX
- tweaks to stdout messages

* Update tests

- Use ROCPROFILER_DEFAULT_FAIL_REGEX

* Add tests/lib

- libraries with HIP code

* Update PTL submodule

- remove atexit delete of thread_id_map

* Update cmake/rocprofiler_options.cmake

- Set ROCPROFILER_DEFAULT_FAIL_REGEX

* Update common lib: env + logging

- improved customization of logging settings
- default to disabling logging to files
- install failure handler for rocprofv3
- set_env support in environment.*

* Add lib/rocprofiler-sdk/shared_library.cpp

- shared library constructor

* Update lib/rocprofiler-sdk-tool/tool.cpp

- destructor thread safety
- convert callback_name_info and buffered_name_info to pointers
- install failure handler for logging

* Add tests/bin/hip-in-libraries

- hip-in-libraries is an exe which uses two shared libraries where each shared library contains HIP kernels
  - used for testing deadlocking within __hipRegisterFatBinary

* Update bin/rocprofv3

- reorganized the env variables
- use exec to launch command
- set ROCPROFILER_LIBRARY_CTOR=1

* Add tests/rocprofv3/tracing-hip-in-libraries

- uses hip-in-libraries exe for exe which uses shared libraries to launch HIP kernels

* Update bin/rocprofv3

- fix counter collection (no exec)

* Update lib/rocprofiler-sdk-tool/tool.cpp

- replace "Kernel-Name" with "Kernel_Name"

* Update lib/rocprofiler-sdk/registration.cpp

Use RTLD_LOCAL instead of RTLD_GLOBAL for env libraries

* Update tests/rocprofv3

- replace "Kernel-Name" with "Kernel_Name"

* Update tests

- vector-ops (bin) stream syncs + runs with 4 queues per device
- improve counter-collection/input1 validation
- rocprofv3/tracing-hip-in-libraries does not do sys-trace
- improved validation script for tracing-hip-in-libraries
- updated dispatch_callback in json-tool.cpp following reworking of prototypes for counter collection

* Update samples/counter_collection

- updated dispatch_callback(s) and record_callback(s) following reworking of prototypes

* Update bin/rocprofv3

- reorganized help menu
- added options for sub-HSA tables
- added --hip-runtime-trace
- changed --hip-trace to include --hip-compiler-trace

* Update lib/rocprofiler-sdk-tool

- improved kernel filtering
- removed arch_vgpr, accum_vgpr, sgpr code (in rocprofiler-sdk)
- fixed issue with counter-collection w/o tracing
- added support for fine grained HSA API tracing
- removed directly linking to HSA-runtime

* Update lib/rocprofiler-sdk/agent.cpp

- rocp_agents != hsa_agents is non-fatal when ROCPROFILER_BUILD_CI=OFF (CMake option)

* GPR (vector and scalar) info in kernel symbol data

- rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t contains general purpose register info

* Header include order fix

- Include repo headers first
- Third party library headers next
- standard library headers last

* Update dispatch profiling public API

- introduce rocprofiler_profile_counting_dispatch_data_t
- change signature of rocprofiler_profile_counting_dispatch_callback_t and rocprofiler_profile_counting_record_callback_t
- provide rocprofiler_user_data_t pointer in dispatch callback
- provide rocprofiler_user_data_t value (from dispatch cb) in record callback

* Update tests/bin/CMakeLists.txt

- fix add_subdirectory(hip-in-libraries) order

* Update VERSION

- bump to 0.2.0 in prep for AFAR

[ROCm/rocprofiler-sdk commit: 7b6d3c70bd]
This commit is contained in:
Jonathan R. Madsen
2024-03-07 22:21:26 -06:00
committed by GitHub
parent 1295c42022
commit 407fc57ede
85 changed files with 2497 additions and 856 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/tests/lib/CMakeLists.txt
+18
View File
@@ -0,0 +1,18 @@
#
# Integration test application libraries
#
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
project(rocprofiler-tests-lib LANGUAGES C CXX)
set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib")
# libraries used by integration test apps which DO NOT link to rocprofiler-sdk-roctx
add_subdirectory(vector-operations)
set(CMAKE_BUILD_RPATH
"\$ORIGIN:\$ORIGIN/../lib:$<TARGET_FILE_DIR:rocprofiler-sdk-roctx::rocprofiler-sdk-roctx-shared-library>"
)
# libraries used by integration test apps which DO link to rocprofiler-sdk-roctx
add_subdirectory(transpose)
projects/rocprofiler-sdk/tests/lib/transpose/CMakeLists.txt
+61
View File
@@ -0,0 +1,61 @@
#
#
#
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
if(NOT CMAKE_HIP_COMPILER)
find_program(
amdclangpp_EXECUTABLE
NAMES amdclang++
HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
PATH_SUFFIXES bin llvm/bin NO_CACHE)
mark_as_advanced(amdclangpp_EXECUTABLE)
if(amdclangpp_EXECUTABLE)
set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
endif()
endif()
project(rocprofiler-tests-lib-transpose-shared-library LANGUAGES CXX HIP)
foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
endif()
endforeach()
option(TRANSPOSE_USE_MPI "Enable MPI support in transpose-shared-library exe" OFF)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_EXTENSIONS OFF)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_HIP_STANDARD 17)
set(CMAKE_HIP_EXTENSIONS OFF)
set(CMAKE_HIP_STANDARD_REQUIRED ON)
set_source_files_properties(transpose.cpp PROPERTIES LANGUAGE HIP)
add_library(transpose-shared-library SHARED)
target_sources(transpose-shared-library PRIVATE transpose.cpp)
target_compile_options(transpose-shared-library PRIVATE -W -Wall -Wextra -Wpedantic
-Wshadow -Werror)
target_include_directories(transpose-shared-library PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
set_target_properties(transpose-shared-library PROPERTIES OUTPUT_NAME transpose)
find_package(Threads REQUIRED)
target_link_libraries(transpose-shared-library PRIVATE Threads::Threads)
find_package(rocprofiler-sdk-roctx REQUIRED)
target_link_libraries(transpose-shared-library
PRIVATE rocprofiler-sdk-roctx::rocprofiler-sdk-roctx)
if(TRANSPOSE_USE_MPI)
find_package(MPI REQUIRED)
target_compile_definitions(transpose-shared-library PRIVATE USE_MPI)
target_link_libraries(transpose-shared-library PRIVATE MPI::MPI_C)
endif()
install(
TARGETS transpose-shared-library
DESTINATION lib
COMPONENT tests)
projects/rocprofiler-sdk/tests/lib/transpose/transpose.cpp
+260
View File
@@ -0,0 +1,260 @@
// 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"
#include "rocprofiler-sdk-roctx/roctx.h"
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <iostream>
#include <mutex>
#include <random>
#include <stdexcept>
#if defined(USE_MPI)
# include <mpi.h>
#endif
#define HIP_API_CALL(CALL) \
{ \
hipError_t error_ = (CALL); \
if(error_ != hipSuccess) \
{ \
auto _hip_api_print_lk = auto_lock_t{print_lock}; \
fprintf(stderr, \
"%s:%d :: HIP error : %s\n", \
__FILE__, \
__LINE__, \
hipGetErrorString(error_)); \
throw std::runtime_error("hip_api_call"); \
} \
}
namespace
{
using auto_lock_t = std::unique_lock<std::mutex>;
auto print_lock = std::mutex{};
constexpr unsigned shared_mem_tile_dim = 32;
void
check_hip_error(void);
void
verify(int* in, int* out, int M, int N);
__global__ void
transpose(const int* in, int* out, int M, int N);
void
run_transpose_impl(int rank, int tid, int ndevice, size_t nitr, size_t nsync);
__global__ void
transpose(const int* in, int* out, int M, int N)
{
__shared__ int tile[shared_mem_tile_dim][shared_mem_tile_dim];
int idx = (blockIdx.y * blockDim.y + threadIdx.y) * M + blockIdx.x * blockDim.x + threadIdx.x;
tile[threadIdx.y][threadIdx.x] = in[idx];
__syncthreads();
idx = (blockIdx.x * blockDim.x + threadIdx.y) * N + blockIdx.y * blockDim.y + threadIdx.x;
out[idx] = tile[threadIdx.x][threadIdx.y];
}
void
run_transpose_impl(int rank, int tid, int devid, size_t nitr, size_t nsync)
{
roctxRangePush("run_transpose_impl");
constexpr unsigned int M = 4960 * 2;
constexpr unsigned int N = 4960 * 2;
hipStream_t stream = {};
printf("[transpose] Rank %i, thread %i assigned to device %i\n", rank, tid, devid);
HIP_API_CALL(hipSetDevice(devid));
HIP_API_CALL(hipStreamCreate(&stream));
auto_lock_t _lk{print_lock};
std::cout << "[transpose][" << rank << "][" << tid << "] M: " << M << " N: " << N << std::endl;
_lk.unlock();
std::default_random_engine _engine{std::random_device{}() * (rank + 1) * (tid + 1)};
std::uniform_int_distribution<int> _dist{0, 1000};
size_t size = sizeof(int) * M * N;
int* inp_matrix = new int[size];
int* out_matrix = new int[size];
for(size_t i = 0; i < M * N; i++)
{
inp_matrix[i] = _dist(_engine);
out_matrix[i] = 0;
}
int* in = nullptr;
int* out = nullptr;
HIP_API_CALL(hipMalloc(&in, size));
HIP_API_CALL(hipMalloc(&out, size));
HIP_API_CALL(hipMemsetAsync(in, 0, size, stream));
HIP_API_CALL(hipMemsetAsync(out, 0, size, stream));
HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, size, hipMemcpyHostToDevice, stream));
HIP_API_CALL(hipStreamSynchronize(stream));
dim3 grid(M / 32, N / 32, 1);
dim3 block(32, 32, 1); // transpose
auto t1 = std::chrono::high_resolution_clock::now();
for(size_t i = 0; i < nitr; ++i)
{
transpose<<<grid, block, 0, stream>>>(in, out, M, N);
check_hip_error();
if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream));
}
auto t2 = std::chrono::high_resolution_clock::now();
HIP_API_CALL(hipStreamSynchronize(stream));
HIP_API_CALL(hipMemcpyAsync(out_matrix, out, size, hipMemcpyDeviceToHost, stream));
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
float GB = (float) size * nitr * 2 / (1 << 30);
print_lock.lock();
std::cout << "[transpose][" << rank << "][" << tid << "] Runtime of transpose is " << time
<< " sec\n";
std::cout << "[transpose][" << rank << "][" << tid
<< "] The average performance of transpose is " << GB / time << " GBytes/sec"
<< std::endl;
print_lock.unlock();
HIP_API_CALL(hipStreamSynchronize(stream));
HIP_API_CALL(hipStreamDestroy(stream));
// cpu_transpose(matrix, out_matrix, M, N);
verify(inp_matrix, out_matrix, M, N);
HIP_API_CALL(hipFree(in));
HIP_API_CALL(hipFree(out));
delete[] inp_matrix;
delete[] out_matrix;
roctxRangePop();
}
void
check_hip_error(void)
{
hipError_t err = hipGetLastError();
if(err != hipSuccess)
{
auto_lock_t _lk{print_lock};
std::cerr << "Error: " << hipGetErrorString(err) << std::endl;
throw std::runtime_error("hip_api_call");
}
}
void
verify(int* in, int* out, int M, int N)
{
for(int i = 0; i < 10; i++)
{
int row = rand() % M;
int col = rand() % N;
if(in[row * N + col] != out[col * M + row])
{
auto_lock_t _lk{print_lock};
std::cout << "mismatch: " << row << ", " << col << " : " << in[row * N + col] << " | "
<< out[col * M + row] << "\n";
}
}
}
} // namespace
void
run_transpose(size_t nthreads, size_t nitr, size_t nsync)
{
auto range_id = roctxRangeStart("run_transpose");
int rank = 0;
int size = 1;
printf("[transpose] Number of threads: %zu\n", nthreads);
printf("[transpose] Number of iterations: %zu\n", nitr);
printf("[transpose] Syncing every %zu iterations\n", nsync);
#if defined(USE_MPI)
MPI_Comm_rank(MPI_COMM_WORLD, &rank);
MPI_Comm_size(MPI_COMM_WORLD, &size);
#else
(void) size;
#endif
// this is a temporary workaround in omnitrace when HIP + MPI is enabled
int ndevice = 0;
HIP_API_CALL(hipGetDeviceCount(&ndevice));
printf("[transpose] Number of devices found: %i\n", ndevice);
auto devids = std::vector<int>{};
devids.resize(size * nthreads, 0);
int devid = 0;
for(size_t i = 0; i < nthreads; ++i)
{
for(int j = 0; j < size; ++j)
{
auto idx = (j * nthreads) + i;
devids.at(idx) = devid++ % ndevice;
}
}
auto devid_offset = (rank * nthreads);
auto _threads = std::vector<std::thread>{};
for(size_t i = 1; i < nthreads; ++i)
_threads.emplace_back(
run_transpose_impl, rank, i, devids.at(devid_offset + i), nitr, nsync);
run_transpose_impl(rank, 0, devids.at(devid_offset + 0), nitr, nsync);
for(auto& itr : _threads)
itr.join();
#if defined(USE_MPI)
MPI_Barrier(MPI_COMM_WORLD);
#endif
// for(int i = 0; i < ndevice; ++i)
// {
// HIP_API_CALL(hipSetDevice(i));
// HIP_API_CALL(hipDeviceSynchronize());
// }
// #if defined(USE_MPI)
// MPI_Barrier(MPI_COMM_WORLD);
// #endif
// if(rank == 0)
// {
// for(int i = 0; i < ndevice; ++i)
// {
// HIP_API_CALL(hipSetDevice(i));
// HIP_API_CALL(hipDeviceReset());
// }
// }
// #if defined(USE_MPI)
// MPI_Barrier(MPI_COMM_WORLD);
// #endif
roctxRangeStop(range_id);
}
projects/rocprofiler-sdk/tests/lib/transpose/transpose.hpp
+28
View File
@@ -0,0 +1,28 @@
// 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.
#pragma once
#include <cstddef>
void
run_transpose(size_t nthreads, size_t nitr, size_t nsync);
projects/rocprofiler-sdk/tests/lib/vector-operations/CMakeLists.txt
+49
View File
@@ -0,0 +1,49 @@
#
#
#
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
if(NOT CMAKE_HIP_COMPILER)
find_program(
amdclangpp_EXECUTABLE
NAMES amdclang++
HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
PATH_SUFFIXES bin llvm/bin NO_CACHE)
mark_as_advanced(amdclangpp_EXECUTABLE)
if(amdclangpp_EXECUTABLE)
set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
endif()
endif()
project(rocprofiler-tests-lib-vector-operations LANGUAGES CXX HIP)
foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
endif()
endforeach()
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_EXTENSIONS OFF)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_HIP_STANDARD 17)
set(CMAKE_HIP_EXTENSIONS OFF)
set(CMAKE_HIP_STANDARD_REQUIRED ON)
set_source_files_properties(vector-ops.cpp PROPERTIES LANGUAGE HIP)
add_library(vector-ops-shared-library SHARED)
target_sources(vector-ops-shared-library PRIVATE vector-ops.cpp)
target_compile_options(vector-ops-shared-library PRIVATE -W -Wall -Wextra -Wpedantic
-Wshadow -Werror)
target_include_directories(vector-ops-shared-library PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
set_target_properties(vector-ops-shared-library PROPERTIES OUTPUT_NAME vector-ops)
find_package(Threads REQUIRED)
target_link_libraries(vector-ops-shared-library PRIVATE Threads::Threads)
install(
TARGETS vector-ops-shared-library
DESTINATION lib
COMPONENT tests)
projects/rocprofiler-sdk/tests/lib/vector-operations/vector-ops.cpp
+291
View File
@@ -0,0 +1,291 @@
// 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 <assert.h>
#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <algorithm>
#include <iostream>
#include <mutex>
#include <vector>
#define HIP_API_CALL(CALL) \
{ \
hipError_t error_ = (CALL); \
if(error_ != hipSuccess) \
{ \
auto _hip_api_print_lk = auto_lock_t{print_lock}; \
fprintf(stderr, \
"%s:%d :: HIP error : %s\n", \
__FILE__, \
__LINE__, \
hipGetErrorString(error_)); \
throw std::runtime_error("hip_api_call"); \
} \
}
namespace
{
using auto_lock_t = std::unique_lock<std::mutex>;
auto print_lock = std::mutex{};
constexpr auto WIDTH = (1 << 12); // 4096
constexpr auto HEIGHT = (1 << 11); // 2048
constexpr auto DEPTH = (1 << 0); // 1
constexpr auto NUM = (WIDTH * HEIGHT * DEPTH);
struct dimensions
{
int x = 1;
int y = 1;
int z = 1;
};
constexpr auto threads_per_block = dimensions{64, 1, 1};
// Computes vectorAdd with matrix-multiply
template <typename Tp>
__global__ void
addition_kernel(Tp* __restrict__ a,
const Tp* __restrict__ b,
const Tp* __restrict__ c,
int width,
int /*height*/)
{
// printf("addition kernel\n");
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
if(x >= WIDTH || y >= HEIGHT) return;
int index = y * width + x;
a[index] = b[index] + c[index];
}
template <typename Tp>
__global__ void
subtract_kernel(Tp* __restrict__ a,
const Tp* __restrict__ b,
const Tp* __restrict__ c,
int width,
int /*height*/)
{
// printf("subtract kernel\n");
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
if(x >= WIDTH || y >= HEIGHT) return;
int index = y * width + x;
a[index] = abs(b[index] - c[index]);
}
template <typename Tp>
__global__ void
multiply_kernel(Tp* __restrict__ a,
const Tp* __restrict__ b,
const Tp* __restrict__ c,
int width,
int /*height*/)
{
// printf("multiply kernel\n");
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
if(x >= WIDTH || y >= HEIGHT) return;
int index = y * width + x;
a[index] = (b[index] - 1) * (c[index] - 1) + 1;
}
template <typename Tp>
__global__ void
divide_kernel(Tp* __restrict__ a,
const Tp* __restrict__ b,
const Tp* __restrict__ c,
int width,
int /*height*/)
{
// printf("divide kernel\n");
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
if(x >= WIDTH || y >= HEIGHT) return;
int index = y * width + x;
a[index] = (b[index] - c[index]) / abs(c[index] + b[index]) + 1;
}
void
run_vector_ops_impl(int num_queue, int device_id)
{
auto t1 = std::chrono::high_resolution_clock::now();
HIP_API_CALL(hipSetDevice(device_id));
std::vector<float*> hostA(num_queue);
std::vector<float*> hostB(num_queue);
std::vector<float*> hostC(num_queue);
std::vector<float*> deviceA(num_queue);
std::vector<float*> deviceB(num_queue);
std::vector<float*> deviceC(num_queue);
std::vector<hipStream_t> streams(num_queue);
auto sync_stream = [num_queue, streams](int q) {
if(q < 0 || q >= num_queue)
throw std::runtime_error{std::string{"invalid stream id: "} + std::to_string(q)};
HIP_API_CALL(hipStreamSynchronize(streams.at(q)));
};
auto sync_streams = [num_queue, sync_stream]() {
for(int i = 0; i < num_queue; ++i)
sync_stream(i);
};
for(int q = 0; q < num_queue; q++)
{
HIP_API_CALL(hipStreamCreateWithFlags(&streams[q], hipStreamNonBlocking));
HIP_API_CALL(hipHostMalloc(&hostA[q], NUM * sizeof(float), 0));
HIP_API_CALL(hipHostMalloc(&hostB[q], NUM * sizeof(float), 0));
HIP_API_CALL(hipHostMalloc(&hostC[q], NUM * sizeof(float), 0));
// initialize the input data
for(int i = 0; i < NUM; i++)
{
hostB[q][i] = static_cast<float>(i);
hostC[q][i] = static_cast<float>(i * 100.0f);
}
HIP_API_CALL(hipMallocAsync(&deviceA[q], NUM * sizeof(float), streams[q]));
HIP_API_CALL(hipMallocAsync(&deviceB[q], NUM * sizeof(float), streams[q]));
HIP_API_CALL(hipMallocAsync(&deviceC[q], NUM * sizeof(float), streams[q]));
HIP_API_CALL(hipMemcpyAsync(
deviceB[q], hostB[q], NUM * sizeof(float), hipMemcpyHostToDevice, streams[q]));
HIP_API_CALL(hipMemcpyAsync(
deviceC[q], hostC[q], NUM * sizeof(float), hipMemcpyHostToDevice, streams[q]));
}
sync_streams();
for(int q = 0; q < num_queue; q++)
{
hipLaunchKernelGGL(addition_kernel,
dim3(WIDTH / threads_per_block.x, HEIGHT / threads_per_block.y),
dim3(threads_per_block.x, threads_per_block.y),
0,
streams[q],
deviceA[q],
deviceB[q],
deviceC[q],
WIDTH,
HEIGHT);
hipLaunchKernelGGL(subtract_kernel,
dim3(WIDTH / threads_per_block.x, HEIGHT / threads_per_block.y),
dim3(threads_per_block.x, threads_per_block.y),
0,
streams[q],
deviceA[q],
deviceB[q],
deviceC[q],
WIDTH,
HEIGHT);
hipLaunchKernelGGL(multiply_kernel,
dim3(WIDTH / threads_per_block.x, HEIGHT / threads_per_block.y),
dim3(threads_per_block.x, threads_per_block.y),
0,
streams[q],
deviceA[q],
deviceB[q],
deviceC[q],
WIDTH,
HEIGHT);
hipLaunchKernelGGL(divide_kernel,
dim3(WIDTH / threads_per_block.x, HEIGHT / threads_per_block.y),
dim3(threads_per_block.x, threads_per_block.y),
0,
streams[q],
deviceB[q],
deviceA[q],
deviceC[q],
WIDTH,
HEIGHT);
}
sync_streams();
for(int q = 0; q < num_queue; q++)
{
HIP_API_CALL(hipMemcpyAsync(
hostA[q], deviceA[q], NUM * sizeof(float), hipMemcpyDeviceToHost, streams[q]));
sync_stream(q);
HIP_API_CALL(hipFree(deviceA[q]));
HIP_API_CALL(hipFree(deviceB[q]));
HIP_API_CALL(hipFree(deviceC[q]));
HIP_API_CALL(hipHostFree(hostA[q]));
HIP_API_CALL(hipHostFree(hostB[q]));
HIP_API_CALL(hipHostFree(hostC[q]));
HIP_API_CALL(hipStreamDestroy(streams[q]));
}
auto t2 = std::chrono::high_resolution_clock::now();
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
print_lock.lock();
std::cout << "[vector-ops] Runtime of vector-ops is " << time << " sec\n";
print_lock.unlock();
}
} // namespace
void
run_vector_ops(int num_threads, int num_queue)
{
int device_count = 0;
HIP_API_CALL(hipGetDeviceCount(&device_count));
if(device_count == 0) throw std::runtime_error{"No HIP devices found"};
num_threads = std::max<int>(num_threads, 1);
num_queue = std::max<int>(num_queue, 1);
auto _threads = std::vector<std::thread>{};
_threads.reserve(num_threads);
for(int i = 0; i < num_threads; ++i)
_threads.emplace_back(run_vector_ops_impl, num_queue, i % device_count);
for(auto& itr : _threads)
itr.join();
}
projects/rocprofiler-sdk/tests/lib/vector-operations/vector-ops.hpp
+26
View File
@@ -0,0 +1,26 @@
// 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.
#pragma once
void
run_vector_ops(int num_threads, int num_queue);