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

[rocprofv3-benchmark] SDK and rocprofv3 Benchmarking Suite (#157)

Browse Source 
* Adding Benchmarking Stg1

* config fix

* reset

* add jpeg and decode traces in iteration

* address comments benchmark config files.

* address comments.

* address comments.

* address comments: revert cntrl ctx.

* address comments: revert csv output.

* resolve merge conflits.

* format.

* build fix.

* fix hip runtime api traces.

* loop cb services.

* format.

* bug fix.

* Fix operator>

- public C++ comparison operator

* Update configuration options

- support selected regions (--selected-regions)
- support writing output config json (--output-config)
- update serialization data

* rocprofv3 tool library misc updates

- lambda for starting context
- support for writing config json

* Tool library updates

- Finished support for all benchmarking modes
- Added build spec support to config json

* Fix ROCPROFILER_SOVERSION

- this value should not be multiplied by 10,000

* Minor tweak to rocprofv3

* Benchmarking scripts

* formatting

* Fix duplicate include

* Add reproducible-dispatch-count test app

- used in benchmarking

* registration logging

- report number of registered contexts and active contexts after client initialization

* Serialize environment in rocprofv3 output config

* ROCPROFILER_BUILD_BENCHMARK CMake option

* Update benchmark SQL schema

- hash_id is text
- add md5sum to benchmarked_app
- remove app_id from benchmarked_sdk
- add sdk_id to benchmark_config
- separate hip_trace into hip_runtime_trace and hip_compiler_trace
- use INT instead of INTEGER for MySQL compatibility
- add count column in benchmark_statistics
- allow std_dev to be NULL in benchmark_statistics

* Update rocprofv3-benchmark.py

- use md5 instead of python hash (which includes random seed)
- use args.mysql_database
- compute md5sum of executable
- fix insert_benchmark_config
  - marker trace fixes
  - memory allocation fixes
  - split hip_trace into hip_{runtime,compiler}_trace
- remove app_id from benchmarked_sdk
- support warmup runs
- count field in benchmark_statistics

* Support launcher and environment in YAML

* Update reproducible-dispatch-count.cpp

- support mode which doesn't use hip event timing

* Misc rocprofv3-benchmark.py updates

- fix some MySQL support
- remove some unnecessary logging

* support mysql db.

* Format.

* Updated SQL input files

- moved benchmark_schema.sql to benchmark_table.sql
- added benchmark_views.sql
  - uses {{metric}} syntax for variable substitution

* cmake formatting

* update rocprofv3-benchmark.py

- benchmark config labels
- overhead views

* Encode rocprofv3-benchmark PID in rocprofv3 and timem output files

* Minor tweak to benchmark_views.sql

- include count
- reorder fields for readability

* split statements and use IS if values is NONE.

* use backtick instead of double quotes and add IS before NOT NULL.:

* Adding Mandelbrot Benchmark App

* Adding Dockerfile example

* Update dockerfile

* Update dockerfile

* [SDK] rocprofiler_query_external_correlation_id_request_kind_name

* Execution-profile benchmark mode

* Execution profile SQL support

* Rename mandlebrot folder + misc clang-tidy

* [rocprofv3-benchmark] Execution profile support

* Update installation

* add work dir when setting git revision, useful when building outside src.

* Set FULL_VERSION_STRING and ROCPROFILER_SDK_GIT_REVISION

- when benchmark folder is top-level

* Remove unused python packages from requirements.txt

* Use ldd/pyelftools to include linked libs for md5sum

- also add --filter-benchmark and --filter-rocprofv3 options
- support labeling the rocprofv3 options
- use more argparse groups
- more generic application of filters
- support variable substitution in environment, e.g. PATH=/some/path:$PATH

* Environment improvements

- improve reproducibility when env set via input file vs. shell
- support "environment-ignore" to remove environment variables

* Misc formatting

* Misc. fix

* use backticks for defining new columns name

* Support shuffling the order of benchmark modes/rocprofv3 args

* Address review comments

* Update Dockerfile

- rename to Dockerfile
- reduce to one layer

* Support docker build arg BRANCH

---------

Co-authored-by: Ammar ELWazir <aelwazir@amd.com>
Co-authored-by: Kandula, Venkateshwar reddy <Venkateshwarreddy.Kandula@amd.com>
Co-authored-by: Venkateshwar Reddy Kandula <vkandula@amd.com>
Co-authored-by: Madsen, Jonathan <Jonathan.Madsen@amd.com>
Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
This commit is contained in:
Elwazir, Ammar
2025-05-13 16:18:23 -05:00
committed by GitHub
parent b297e72853
commit 6f17da7ade
43 changed files with 5206 additions and 280 deletions
Download Patch File Download Diff File Expand all files Collapse all files
benchmark/source/CMakeLists.txt
+7
View File
@@ -0,0 +1,7 @@
#
#
#
add_subdirectory(lib)
add_subdirectory(bin)
add_subdirectory(share)
benchmark/source/bin/CMakeLists.txt
+24
View File
@@ -0,0 +1,24 @@
#
#
#
# Adding main rocprofv3
configure_file(rocprofv3-benchmark.py
${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/rocprofv3-benchmark @ONLY)
install(
FILES ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/rocprofv3-benchmark
DESTINATION ${CMAKE_INSTALL_BINDIR}
PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ
WORLD_EXECUTE)
# downloaded timem
install(
FILES ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/timem
DESTINATION ${CMAKE_INSTALL_BINDIR}
PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ
WORLD_EXECUTE
OPTIONAL)
# Adding Benchmark Workloads
add_subdirectory(mandelbrot)
benchmark/source/bin/mandelbrot/CMakeLists.txt
+44
View File
@@ -0,0 +1,44 @@
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-sdk-benchmark-bin-mandelbrot 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(mandelbrot.cpp PROPERTIES LANGUAGE HIP)
set_source_files_properties(utils.cpp PROPERTIES LANGUAGE HIP)
add_executable(mandelbrot)
target_sources(mandelbrot PRIVATE mandelbrot.cpp utils.cpp)
target_compile_options(mandelbrot PRIVATE -W -Wall -Wextra -Wpedantic -Werror
-ffp-contract=fast)
target_include_directories(mandelbrot PRIVATE ${CMAKE_CURRENT_SOURCE_DIR})
install(
TARGETS mandelbrot
DESTINATION ${CMAKE_INSTALL_BINDIR}
COMPONENT benchmark)
benchmark/source/bin/mandelbrot/mandelbrot.cpp
+938
View File
@@ -0,0 +1,938 @@
/*
Copyright (c) 2015 - 2021 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.
*/
/* HIT_START
* BUILD: %t %s ../../src/test_common.cpp
* TEST: %t
* HIT_END
*/
#include "utils.hpp"
#include <hip/hip_vector_types.h>
#include <hip/math_functions.h>
#include <omp.h>
#include <chrono>
#include <iostream>
#include <map>
#include <string>
#include <vector>
typedef struct
{
double x;
double y;
double width;
} coordRec;
coordRec coords[] = {
{0.0, 0.0, 4.0}, // Whole set
{0.0, 0.0, 0.00001}, // All black
{-0.0180789661868, 0.6424294066162, 0.00003824140}, // Hit detail
};
static unsigned int numCoords = sizeof(coords) / sizeof(coordRec);
template <typename T>
__global__ void
float_mad_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T yStep, uint maxIter)
{
int tid = (blockIdx.x * blockDim.x + threadIdx.x);
int i = tid % width;
int j = tid / width;
float x0 = (float) (xPos + xStep * i);
float y0 = (float) (yPos + yStep * j);
float x = x0;
float y = y0;
uint iter = 0;
float tmp;
for(iter = 0; (x * x + y * y <= 4.0f) && (iter < maxIter); iter++)
{
tmp = x;
x = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * tmp, y, y0);
}
out[tid] = iter;
};
template <typename T>
__global__ void
float_mandel_unroll_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T yStep, uint maxIter)
{
int tid = (blockIdx.x * blockDim.x + threadIdx.x);
int i = tid % width;
int j = tid / width;
float x0 = (float) (xPos + xStep * (float) i);
float y0 = (float) (yPos + yStep * (float) j);
float x = x0;
float y = y0;
#define FAST
uint iter = 0;
float tmp;
int stay;
uint ccount = 0;
stay = (x * x + y * y) <= 4.0;
float savx = x;
float savy = y;
#ifdef FAST
for(iter = 0; (iter < maxIter); iter += 16)
{
#else
for(iter = 0; stay && (iter < maxIter); iter += 16)
{
#endif
x = savx;
y = savy;
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
stay = (x * x + y * y) <= 4.0;
savx = (stay ? x : savx);
savy = (stay ? y : savy);
ccount += stay * 16;
#ifdef FAST
if(!stay) break;
#endif
}
// Handle remainder
if(!stay)
{
iter = 16;
do
{
x = savx;
y = savy;
stay = ((x * x + y * y) <= 4.0) && (ccount < maxIter);
tmp = x;
x = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * tmp, y, y0);
ccount += stay;
iter--;
savx = (stay ? x : savx);
savy = (stay ? y : savy);
} while(stay && iter);
}
out[tid] = (uint) ccount;
};
template <typename T>
__global__ void
double_mad_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T yStep, uint maxIter)
{
int tid = (blockIdx.x * blockDim.x + threadIdx.x);
int i = tid % width;
int j = tid / width;
double x0 = (double) (xPos + xStep * i);
double y0 = (double) (yPos + yStep * j);
double x = x0;
double y = y0;
uint iter = 0;
double tmp;
for(iter = 0; (x * x + y * y <= 4.0f) && (iter < maxIter); iter++)
{
tmp = x;
x = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * tmp, y, y0);
}
out[tid] = iter;
};
template <typename T>
__global__ void
double_mandel_unroll_kernel(uint* out, uint width, T xPos, T yPos, T xStep, T yStep, uint maxIter)
{
int tid = (blockIdx.x * blockDim.x + threadIdx.x);
int i = tid % width;
int j = tid / width;
double x0 = (double) (xPos + xStep * (double) i);
double y0 = (double) (yPos + yStep * (double) j);
double x = x0;
double y = y0;
#define FAST
uint iter = 0;
double tmp;
int stay;
uint ccount = 0;
stay = (x * x + y * y) <= 4.0;
double savx = x;
double savy = y;
#ifdef FAST
for(iter = 0; (iter < maxIter); iter += 16)
#else
for(iter = 0; stay && (iter < maxIter); iter += 16)
#endif
{
x = savx;
y = savy;
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
// Two iterations
tmp = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * x, y, y0);
x = fma(-y, y, fma(tmp, tmp, x0));
y = fma(2.0f * tmp, y, y0);
stay = (x * x + y * y) <= 4.0;
savx = (stay ? x : savx);
savy = (stay ? y : savy);
ccount += stay * 16;
#ifdef FAST
if(!stay) break;
#endif
}
// Handle remainder
if(!stay)
{
iter = 16;
do
{
x = savx;
y = savy;
stay = ((x * x + y * y) <= 4.0) && (ccount < maxIter);
tmp = x;
x = fma(-y, y, fma(x, x, x0));
y = fma(2.0f * tmp, y, y0);
ccount += stay;
iter--;
savx = (stay ? x : savx);
savy = (stay ? y : savy);
} while(stay && iter);
}
out[tid] = (uint) ccount;
};
// Commenting it out as it is not used anywhere in the code
// static const unsigned int FMA_EXPECTEDVALUES_INDEX = 15;
// Expected results for each kernel run at each coord
unsigned long long expectedIters[] = {
203277748ull, 2147483648ull, 120254651ull, 203277748ull, 2147483648ull, 120254651ull,
203277748ull, 2147483648ull, 120254651ull, 203315114ull, 2147483648ull, 120042599ull,
203315114ull, 2147483648ull, 120042599ull, 203280620ull, 2147483648ull, 120485704ull,
203280620ull, 2147483648ull, 120485704ull, 203280620ull, 2147483648ull, 120485704ull,
203315114ull, 2147483648ull, 120042599ull, 203315114ull, 2147483648ull, 120042599ull};
class hipPerfMandelBrot
{
public:
hipPerfMandelBrot();
~hipPerfMandelBrot();
void setNumKernels(unsigned int num) { numKernels = num; }
unsigned int getNumKernels() const { return numKernels; }
void setNumStreams(unsigned int num) { numStreams = num; }
unsigned int getNumStreams() const { return numStreams; }
void open(int deviceID);
void run(unsigned int testCase, unsigned int deviceId);
void printResults(void);
// array of funtion pointers
typedef void (hipPerfMandelBrot::*funPtr)(uint* out,
uint width,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt);
// Wrappers
void float_mad(uint* out,
uint width,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt);
void float_mandel_unroll(uint* out,
uint width,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt);
void double_mad(uint* out,
uint width,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt);
void double_mandel_unroll(uint* out,
uint width,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt);
hipStream_t streams[2];
private:
void setData(void* ptr, unsigned int value);
void checkData(uint* ptr);
unsigned int numKernels;
unsigned int numStreams;
std::map<std::string, std::vector<double>> results;
unsigned int width_;
unsigned int bufSize;
unsigned int maxIter;
unsigned int coordIdx;
volatile unsigned long long totalIters = 0;
int numCUs;
static const unsigned int numLoops = 10;
};
hipPerfMandelBrot::hipPerfMandelBrot() = default;
hipPerfMandelBrot::~hipPerfMandelBrot() = default;
void
hipPerfMandelBrot::open(int deviceId)
{
int nGpu = 0;
HIPCHECK(hipGetDeviceCount(&nGpu));
if(nGpu < 1)
{
failed("No GPU!");
}
HIPCHECK(hipSetDevice(deviceId));
hipDeviceProp_t props = {};
HIPCHECK(hipGetDeviceProperties(&props, deviceId));
std::cout << "info: running on bus "
<< "0x" << props.pciBusID << " " << props.name << " with "
<< props.multiProcessorCount << " CUs"
<< " and device id: " << deviceId << std::endl;
numCUs = props.multiProcessorCount;
}
void
hipPerfMandelBrot::printResults()
{
// int numkernels = getNumKernels();
int numStreams = getNumStreams();
std::cout << "\n"
<< "Measured perf for kernels in GFLOPS on " << numStreams << " streams (s)"
<< std::endl;
std::map<std::string, std::vector<double>>::iterator itr;
for(itr = results.begin(); itr != results.end(); itr++)
{
std::cout << "\n" << std::setw(20) << itr->first << " ";
for(auto i : results[itr->first])
{
std::cout << std::setw(10) << i << " ";
}
}
results.clear();
std::cout << std::endl;
}
// Wrappers for the kernel launches
void
hipPerfMandelBrot::float_mad(uint* out,
uint /* width */,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt)
{
int streamCnt = getNumStreams();
hipLaunchKernelGGL(float_mad_kernel<float>,
dim3(blocks),
dim3(threads_per_block),
0,
streams[kernelCnt % streamCnt],
out,
width_,
xPos,
yPos,
xStep,
yStep,
maxIter);
}
void
hipPerfMandelBrot::float_mandel_unroll(uint* out,
uint /* width */,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt)
{
int streamCnt = getNumStreams();
hipLaunchKernelGGL(float_mandel_unroll_kernel<float>,
dim3(blocks),
dim3(threads_per_block),
0,
streams[kernelCnt % streamCnt],
out,
width_,
xPos,
yPos,
xStep,
yStep,
maxIter);
}
void
hipPerfMandelBrot::double_mad(uint* out,
uint /* width */,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt)
{
int streamCnt = getNumStreams();
hipLaunchKernelGGL(double_mad_kernel<double>,
dim3(blocks),
dim3(threads_per_block),
0,
streams[kernelCnt % streamCnt],
out,
width_,
xPos,
yPos,
xStep,
yStep,
maxIter);
}
void
hipPerfMandelBrot::double_mandel_unroll(uint* out,
uint /* width */,
float xPos,
float yPos,
float xStep,
float yStep,
uint maxIter,
hipStream_t* streams,
int blocks,
int threads_per_block,
int kernelCnt)
{
int streamCnt = getNumStreams();
hipLaunchKernelGGL(float_mandel_unroll_kernel<double>,
dim3(blocks),
dim3(threads_per_block),
0,
streams[kernelCnt % streamCnt],
out,
width_,
xPos,
yPos,
xStep,
yStep,
maxIter);
}
void
hipPerfMandelBrot::run(unsigned int testCase, unsigned int /* deviceId */)
{
unsigned int numStreams = getNumStreams();
coordIdx = testCase % numCoords;
funPtr p[] = {&hipPerfMandelBrot::float_mad,
&hipPerfMandelBrot::float_mandel_unroll,
&hipPerfMandelBrot::double_mad,
&hipPerfMandelBrot::double_mandel_unroll};
// Maximum iteration count
maxIter = 32768;
// Variable-length arrays (VLAs) are not part of the C++ standard, using std::vector instead
// uint * hPtr[numKernels];
// uint * dPtr[numKernels];
std::vector<uint*> hPtr(numKernels);
std::vector<uint*> dPtr(numKernels);
// Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once.
width_ = 256;
bufSize = width_ * width_ * sizeof(uint);
// Create streams for concurrency
for(uint i = 0; i < numStreams; i++)
{
HIPCHECK(hipStreamCreate(&streams[i]));
}
// Allocate memory on the host and device
for(uint i = 0; i < numKernels; i++)
{
HIPCHECK(hipHostMalloc((void**) &hPtr[i], bufSize, hipHostMallocDefault));
setData(hPtr[i], 0xdeadbeef);
HIPCHECK(hipMalloc((uint**) &dPtr[i], bufSize))
}
// Prepare kernel launch parameters
int threads = (bufSize / sizeof(uint));
int threads_per_block = 64;
int blocks = (threads / threads_per_block) + (threads % threads_per_block);
// float xStep = (float)(coords[coordIdx].width / (double)width_);
// float yStep = (float)(-coords[coordIdx].width / (double)width_);
// float xPos = (float)(coords[coordIdx].x - 0.5 * coords[coordIdx].width);
// float yPos = (float)(coords[coordIdx].y + 0.5 * coords[coordIdx].width);
// Copy memory asynchronously and concurrently from host to device
for(uint i = 0; i < numKernels; i++)
{
HIPCHECK(hipMemcpy(dPtr[i], hPtr[i], bufSize, hipMemcpyHostToDevice));
}
// Synchronize to make sure all the copies are completed
HIPCHECK(hipStreamSynchronize(nullptr));
int kernelIdx;
if(testCase == 0 || testCase == 5 || testCase == 10)
{
kernelIdx = 0;
}
else if(testCase == 1 || testCase == 6 || testCase == 11)
{
kernelIdx = 1;
}
else if(testCase == 2 || testCase == 7 || testCase == 12)
{
kernelIdx = 2;
}
else if(testCase == 3 || testCase == 8 || testCase == 13)
{
kernelIdx = 3;
}
double totalTime = 0.0;
for(unsigned int k = 0; k < numLoops; k++)
{
if((testCase == 0 || testCase == 1 || testCase == 2 || testCase == 5 || testCase == 6 ||
testCase == 7 || testCase == 10 || testCase == 11 || testCase == 12))
{
float xStep = (float) (coords[coordIdx].width / (double) width_);
float yStep = (float) (-coords[coordIdx].width / (double) width_);
float xPos = (float) (coords[coordIdx].x - 0.5 * coords[coordIdx].width);
float yPos = (float) (coords[coordIdx].y + 0.5 * coords[coordIdx].width);
// Time the kernel execution
auto all_start = std::chrono::steady_clock::now();
for(uint i = 0; i < numKernels; i++)
{
(this->*p[kernelIdx])(dPtr[i],
width_,
xPos,
yPos,
xStep,
yStep,
maxIter,
streams,
blocks,
threads_per_block,
i);
}
// Synchronize all the concurrent streams to have completed execution
HIPCHECK(hipStreamSynchronize(0));
auto all_end = std::chrono::steady_clock::now();
std::chrono::duration<double> all_kernel_time = all_end - all_start;
totalTime += all_kernel_time.count();
}
else
{
double xStep = coords[coordIdx].width / (double) width_;
double yStep = -coords[coordIdx].width / (double) width_;
double xPos = coords[coordIdx].x - 0.5 * coords[coordIdx].width;
double yPos = coords[coordIdx].y + 0.5 * coords[coordIdx].width;
// Time the kernel execution
auto all_start = std::chrono::steady_clock::now();
for(uint i = 0; i < numKernels; i++)
{
(this->*p[kernelIdx])(dPtr[i],
width_,
xPos,
yPos,
xStep,
yStep,
maxIter,
streams,
blocks,
threads_per_block,
i);
}
// Synchronize all the concurrent streams to have completed execution
HIPCHECK(hipStreamSynchronize(0));
auto all_end = std::chrono::steady_clock::now();
std::chrono::duration<double> all_kernel_time = all_end - all_start;
totalTime += all_kernel_time.count();
}
}
// Copy data back from device to the host
for(uint i = 0; i < numKernels; i++)
{
HIPCHECK(hipMemcpy(hPtr[i], dPtr[i], bufSize, hipMemcpyDeviceToHost));
}
for(uint i = 0; i < numKernels; i++)
{
checkData(hPtr[i]);
int j = 0;
while((totalIters != expectedIters[j] && totalIters > expectedIters[j]) && j < 30)
{
j++;
}
if(j == 30)
{
std::cout << "Incorrect iteration count detected. ";
}
}
// Compute GFLOPS. There are 7 FLOPs per iteration
double perf = ((double) (totalIters * numKernels) * 7 * (double) (1e-09)) /
(totalTime / (double) numLoops);
std::vector<std::string> kernelName = {"float", "float_unroll", "double", "double_unroll"};
// Print results except for Warm-up kernel
if(testCase != 100)
{
results[kernelName[testCase % 4]].push_back(perf);
}
for(uint i = 0; i < numStreams; i++)
{
HIPCHECK(hipStreamDestroy(streams[i]));
}
// Free host and device memory
for(uint i = 0; i < numKernels; i++)
{
HIPCHECK(hipHostFree(hPtr[i]));
HIPCHECK(hipFree(dPtr[i]));
}
}
void
hipPerfMandelBrot::setData(void* ptr, unsigned int value)
{
unsigned int* ptr2 = (unsigned int*) ptr;
for(unsigned int i = 0; i < width_ * width_; i++)
{
ptr2[i] = value;
}
}
void
hipPerfMandelBrot::checkData(uint* ptr)
{
totalIters = 0;
for(unsigned int i = 0; i < width_ * width_; i++)
{
totalIters += ptr[i];
}
}
int
main(int argc, char* argv[])
{
// Default values for kernels and streams
unsigned int numStreamsWarmup = 1, numKernelsWarmup = 1;
unsigned int numStreamsSync = 1, numKernelsSync = 1;
unsigned int numStreamsAsync = 2, numKernelsAsync = 2;
// Check for help arguments
if(argc > 1 && (std::string(argv[1]) == "-h" || std::string(argv[1]) == "--help" ||
std::string(argv[1]) == "help"))
{
std::cout << "Usage: " << argv[0]
<< " [--warmup <numStreams>:<numKernels>] [--sync <numStreams>:<numKernels>] "
"[--async <numStreams>:<numKernels>]"
<< std::endl;
std::cout << "Example: " << argv[0] << " --warmup 1:1 --sync 2:4 --async 3:6" << std::endl;
std::exit(EXIT_SUCCESS); // Exit with success status
}
// Parse command-line arguments
for(int i = 1; i < argc; i++)
{
std::string arg = argv[i];
if(arg == "--warmup" && i + 1 < argc)
{
std::string value = argv[++i];
std::stringstream ss(value);
char delimiter;
ss >> numStreamsWarmup >> delimiter >> numKernelsWarmup;
if(delimiter != ':' || ss.fail())
{
std::cerr << "Invalid format for --warmup. Expected <numStreams>:<numKernels>."
<< std::endl;
std::exit(EXIT_FAILURE);
}
}
else if(arg == "--sync" && i + 1 < argc)
{
std::string value = argv[++i];
std::stringstream ss(value);
char delimiter;
ss >> numStreamsSync >> delimiter >> numKernelsSync;
if(delimiter != ':' || ss.fail())
{
std::cerr << "Invalid format for --sync. Expected <numStreams>:<numKernels>."
<< std::endl;
std::exit(EXIT_FAILURE);
}
}
else if(arg == "--async" && i + 1 < argc)
{
std::string value = argv[++i];
std::stringstream ss(value);
char delimiter;
ss >> numStreamsAsync >> delimiter >> numKernelsAsync;
if(delimiter != ':' || ss.fail())
{
std::cerr << "Invalid format for --async. Expected <numStreams>:<numKernels>."
<< std::endl;
std::exit(EXIT_FAILURE);
}
}
else
{
std::cerr << "Unknown argument: " << arg << std::endl;
std::exit(EXIT_FAILURE);
}
}
int deviceCount = 0;
HIPCHECK(hipGetDeviceCount(&deviceCount)); // Get the number of devices
#pragma omp parallel for
for(int deviceId = 0; deviceId < deviceCount; deviceId++)
{
hipPerfMandelBrot mandelbrotCompute;
mandelbrotCompute.open(deviceId);
for(unsigned int testCase = 0; testCase < 3; testCase++)
{
switch(testCase)
{
case 0:
{
// Warmup-kernel - default stream executes serially
mandelbrotCompute.setNumStreams(numStreamsWarmup);
mandelbrotCompute.setNumKernels(numKernelsWarmup);
mandelbrotCompute.run(100 /*Random number*/, deviceId);
break;
}
case 1:
{
// run all - sync
int i = 0;
do
{
mandelbrotCompute.setNumStreams(numStreamsSync);
mandelbrotCompute.setNumKernels(numKernelsSync);
mandelbrotCompute.run(i, deviceId);
i++;
} while(i < 12);
mandelbrotCompute.printResults();
break;
}
case 2:
{
// run all - async
int i = 0;
do
{
mandelbrotCompute.setNumStreams(numStreamsAsync);
mandelbrotCompute.setNumKernels(numKernelsAsync);
mandelbrotCompute.run(i, deviceId);
i++;
} while(i < 12);
mandelbrotCompute.printResults();
break;
}
default:
{
break;
}
}
}
}
passed();
}
benchmark/source/bin/mandelbrot/utils.cpp
+324
View File
@@ -0,0 +1,324 @@
/*
Copyright (c) 2015 - 2021 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 "utils.hpp"
#include <thread>
#ifdef __linux__
# include <sys/sysinfo.h>
#elif defined(_WIN32)
# include <windows.h>
#endif
// standard global variables that can be set on command line
size_t N = 4 * 1024 * 1024;
char memsetval = 0x42;
int memsetD32val = 0xDEADBEEF;
short memsetD16val = 0xDEAD;
char memsetD8val = 0xDE;
int iterations = 1;
unsigned blocksPerCU = 6; // to hide latency
unsigned threadsPerBlock = 256;
int textureFilterMode = 0; // 0: hipFilterModePoint; 1: hipFilterModeLinear
int p_gpuDevice = 0;
unsigned p_verbose = 0;
int p_tests = -1; /*which tests to run. Interpretation is left to each test. default:all*/
int debug_test = 0;
#ifdef _WIN64
const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES=";
const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES=";
const char* PATH_SEPERATOR_STR = "\\";
const char* NULL_DEVICE = "NUL:";
#else
const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES";
const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES";
const char* PATH_SEPERATOR_STR = "/";
const char* NULL_DEVICE = "/dev/null";
#endif
#ifdef _WIN64
// Windows does not have rand_r, use srand and rand instead.
int
rand_r(unsigned int* s)
{
srand(*s);
return rand();
}
#endif
// Get Free Memory from the system
static size_t
getMemoryAmount()
{
#if __linux__
struct sysinfo info;
int _ = sysinfo(&info);
return info.freeram / (1024 * 1024); // MB
#elif defined(_WIN32)
MEMORYSTATUSEX statex;
statex.dwLength = sizeof(statex);
GlobalMemoryStatusEx(&statex);
return (statex.ullAvailPhys / (1024 * 1024)); // MB
#endif
}
size_t
getHostThreadCount(const size_t memPerThread, const size_t maxThreads)
{
if(memPerThread == 0) return 0;
auto memAmount = getMemoryAmount();
const auto processor_count = std::thread::hardware_concurrency();
if(processor_count == 0 || memAmount == 0) return 0;
size_t thread_count = 0;
if((processor_count * memPerThread) < memAmount)
thread_count = processor_count;
else
thread_count = reinterpret_cast<size_t>(memAmount / memPerThread);
if(maxThreads > 0)
{
return (thread_count > maxThreads) ? maxThreads : thread_count;
}
return thread_count;
}
// Function to determine if the device is of gfx11 architecture
bool
IsGfx11()
{
#if defined(__HIP_PLATFORM_NVIDIA__)
return false;
#elif defined(__HIP_PLATFORM_AMD__)
int device = -1;
hipDeviceProp_t props{};
HIPCHECK(hipGetDevice(&device));
HIPCHECK(hipGetDeviceProperties(&props, device));
// Get GCN Arch Name and compare to check if it is gfx11
std::string arch = std::string(props.gcnArchName);
auto pos = arch.find(":");
if(pos != std::string::npos) arch = arch.substr(0, pos);
if(arch.size() >= 5) arch = arch.substr(0, 5);
return (arch == std::string("gfx11")) ? true : false;
#else
std::cout << "Have to be either Nvidia or AMD platform, asserting" << std::endl;
assert(false);
#endif
}
namespace HipTest
{
double
elapsed_time(long long startTimeUs, long long stopTimeUs)
{
return ((double) (stopTimeUs - startTimeUs)) / ((double) (1000));
}
int
parseSize(const char* str, size_t* output)
{
char* next;
*output = strtoull(str, &next, 0);
int l = strlen(str);
if(l)
{
char c = str[l - 1]; // last char.
if((c == 'k') || (c == 'K'))
{
*output *= 1024;
}
if((c == 'm') || (c == 'M'))
{
*output *= (1024 * 1024);
}
if((c == 'g') || (c == 'G'))
{
*output *= (1024 * 1024 * 1024);
}
}
return 1;
}
int
parseUInt(const char* str, unsigned int* output)
{
char* next;
*output = strtoul(str, &next, 0);
return !strlen(next);
}
int
parseInt(const char* str, int* output)
{
char* next;
*output = strtol(str, &next, 0);
return !strlen(next);
}
int
parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg)
{
int extraArgs = 1;
for(int i = 1; i < argc; i++)
{
const char* arg = argv[i];
if(!strcmp(arg, " "))
{
// skip NULL args.
}
else if(!strcmp(arg, "--N") || (!strcmp(arg, "-N")))
{
if(++i >= argc || !HipTest::parseSize(argv[i], &N))
{
failed("Bad N size argument");
}
}
else if(!strcmp(arg, "--threadsPerBlock"))
{
if(++i >= argc || !HipTest::parseUInt(argv[i], &threadsPerBlock))
{
failed("Bad threadsPerBlock argument");
}
}
else if(!strcmp(arg, "--blocksPerCU"))
{
if(++i >= argc || !HipTest::parseUInt(argv[i], &blocksPerCU))
{
failed("Bad blocksPerCU argument");
}
}
else if(!strcmp(arg, "--memsetval"))
{
int ex;
if(++i >= argc || !HipTest::parseInt(argv[i], &ex))
{
failed("Bad memsetval argument");
}
memsetval = ex;
}
else if(!strcmp(arg, "--memsetD32val"))
{
int ex;
if(++i >= argc || !HipTest::parseInt(argv[i], &ex))
{
failed("Bad memsetD32val argument");
}
memsetD32val = ex;
}
else if(!strcmp(arg, "--memsetD16val"))
{
int ex;
if(++i >= argc || !HipTest::parseInt(argv[i], &ex))
{
failed("Bad memsetD16val argument");
}
memsetD16val = ex;
}
else if(!strcmp(arg, "--memsetD8val"))
{
int ex;
if(++i >= argc || !HipTest::parseInt(argv[i], &ex))
{
failed("Bad memsetD8val argument");
}
memsetD8val = ex;
}
else if(!strcmp(arg, "--textureFilterMode"))
{
int mode;
if(++i >= argc || !HipTest::parseInt(argv[i], &mode))
{
failed("Bad textureFilterMode argument");
}
textureFilterMode = mode;
}
else if(!strcmp(arg, "--iterations") || (!strcmp(arg, "-i")))
{
if(++i >= argc || !HipTest::parseInt(argv[i], &iterations))
{
failed("Bad iterations argument");
}
}
else if(!strcmp(arg, "--gpu") || (!strcmp(arg, "-gpuDevice")) || (!strcmp(arg, "-g")))
{
if(++i >= argc || !HipTest::parseInt(argv[i], &p_gpuDevice))
{
failed("Bad gpuDevice argument");
}
}
else if(!strcmp(arg, "--verbose") || (!strcmp(arg, "-v")))
{
if(++i >= argc || !HipTest::parseUInt(argv[i], &p_verbose))
{
failed("Bad verbose argument");
}
}
else if(!strcmp(arg, "--tests") || (!strcmp(arg, "-t")))
{
if(++i >= argc || !HipTest::parseInt(argv[i], &p_tests))
{
failed("Bad tests argument");
}
}
else if(!strcmp(arg, "--debug") || (!strcmp(arg, "-d")))
{
if(++i >= argc || !HipTest::parseInt(argv[i], &debug_test))
{
failed("Bad tests argument");
}
}
else
{
if(failOnUndefinedArg)
{
failed("Bad argument '%s'", arg);
}
else
{
argv[extraArgs++] = argv[i];
}
}
};
return extraArgs;
}
unsigned
setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N)
{
int device;
HIPCHECK(hipGetDevice(&device));
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device));
unsigned blocks = props.multiProcessorCount * blocksPerCU;
if(blocks * threadsPerBlock > N)
{
blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
}
return blocks;
}
} // namespace HipTest
benchmark/source/bin/mandelbrot/utils.hpp
+733
View File
@@ -0,0 +1,733 @@
/*
Copyright (c) 2015 - 2021 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.
*/
/*
* File is intended to C and CPP compliant hence any CPP specic changes
* should be added into CPP section
*
*/
#pragma once
#ifdef __cplusplus
# include <iomanip>
# include <iostream>
# if __CUDACC__
# include <sys/time.h>
# else
# include <chrono>
# endif
#endif
// ************************ GCC section **************************
#include <stddef.h>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#define HC __attribute__((hc))
#define KNRM "\x1B[0m"
#define KRED "\x1B[31m"
#define KGRN "\x1B[32m"
#define KYEL "\x1B[33m"
#define KBLU "\x1B[34m"
#define KMAG "\x1B[35m"
#define KCYN "\x1B[36m"
#define KWHT "\x1B[37m"
// HIP Skip Return code set at cmake
#define HIP_SKIP_RETURN_CODE 127
#define HIP_ENABLE_SKIP_TESTS 0
// Recommended thresholds for Tests
#define MAX_THREADS 100
inline bool
hip_skip_tests_enabled()
{
return HIP_ENABLE_SKIP_TESTS;
}
inline int
hip_skip_retcode()
{
// HIP Skip Return code set at cmake
return HIP_SKIP_RETURN_CODE;
}
// This must be called in the end of main() to indicate test passed with success.
// If it's called somewhere else, compiling issues or unexpected result will arise.
#define passed() \
printf("%sPASSED!%s\n", KGRN, KNRM); \
return 0;
// The real "assert" would have written to stderr. But it is
// sufficient to just fflush here without getting pedantic. This also
// ensures that we don't lose any earlier writes to stdout.
#define failed(...) \
printf("%serror: ", KRED); \
printf(__VA_ARGS__); \
printf("\n"); \
printf("error: TEST FAILED\n%s", KNRM); \
fflush(NULL); \
abort();
#define warn(...) \
printf("%swarn: ", KYEL); \
printf(__VA_ARGS__); \
printf("\n"); \
printf("warn: TEST WARNING\n%s", KNRM);
#define HIP_PRINT_STATUS(status) \
std::cout << hipGetErrorName(status) << " at line: " << __LINE__ << std::endl;
#define HIPCHECK(error) \
{ \
hipError_t localError = error; \
if((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) \
{ \
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
KRED, \
hipGetErrorString(localError), \
localError, \
#error, \
__FILE__, \
__LINE__, \
KNRM); \
failed("API returned error code."); \
} \
}
#define HIPASSERT(condition) \
if(!(condition)) \
{ \
failed("%sassertion %s at %s:%d%s \n", KRED, #condition, __FILE__, __LINE__, KNRM); \
}
#define HIPCHECK_API(API_CALL, EXPECTED_ERROR) \
{ \
hipError_t _e = (API_CALL); \
if(_e != (EXPECTED_ERROR)) \
{ \
failed("%sAPI '%s' returned %d(%s) but test expected %d(%s) at %s:%d%s \n", \
KRED, \
#API_CALL, \
_e, \
hipGetErrorName(_e), \
EXPECTED_ERROR, \
hipGetErrorName(EXPECTED_ERROR), \
__FILE__, \
__LINE__, \
KNRM); \
} \
}
#define HIPCHECK_RETURN_ONFAIL(func) \
do \
{ \
hipError_t herror = (func); \
if(herror != hipSuccess) \
{ \
return herror; \
} \
} while(0);
#ifdef _WIN64
# include <tchar.h>
# define aligned_alloc(x, y) _aligned_malloc(y, x)
# define aligned_free(x) _aligned_free(x)
# define popen(x, y) _popen(x, y)
# define pclose(x) _pclose(x)
# define setenv(x, y, z) _putenv_s(x, y)
# define unsetenv _putenv
# define fileno(x) _fileno(x)
# define dup(x) _dup(x)
# define dup2(x, y) _dup2(x, y)
# define pipe(x, y, z) _pipe(x, y, z)
# define sleep(x) _sleep(x)
#else
# define aligned_free(x) free(x)
#endif
// standard command-line variables:
extern size_t N;
extern char memsetval;
extern int memsetD32val;
extern short memsetD16val;
extern char memsetD8val;
extern int iterations;
extern unsigned blocksPerCU;
extern unsigned threadsPerBlock;
extern int textureFilterMode;
extern int p_gpuDevice;
extern unsigned p_verbose;
extern int p_tests;
extern int debug_test;
extern const char* HIP_VISIBLE_DEVICES_STR;
extern const char* CUDA_VISIBLE_DEVICES_STR;
extern const char* PATH_SEPERATOR_STR;
extern const char* NULL_DEVICE;
// ********************* CPP section *********************
#ifdef __cplusplus
# ifdef __HIP_PLATFORM_HCC
# define TYPENAME(T) typeid(T).name()
# else
# define TYPENAME(T) "?"
# endif
# ifdef _WIN64
int
rand_r(unsigned int* s);
# endif
// Get Optimal Thread count size
size_t
getHostThreadCount(const size_t memPerThread = 200 /* MB */, const size_t maxThreads = 0);
namespace HipTest
{
// Returns the current system time in microseconds
inline long long
get_time()
{
# if __CUDACC__
struct timeval tv;
gettimeofday(&tv, 0);
return (tv.tv_sec * 1000000) + tv.tv_usec;
# else
return std::chrono::high_resolution_clock::now().time_since_epoch() /
std::chrono::microseconds(1);
# endif
}
double
elapsed_time(long long startTimeUs, long long stopTimeUs);
int
parseSize(const char* str, size_t* output);
int
parseUInt(const char* str, unsigned int* output);
int
parseInt(const char* str, int* output);
int
parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg);
unsigned
setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N);
template <typename T> // pointer type
void
checkArray(T hData, T hOutputData, size_t width, size_t height, size_t depth)
{
for([[maybe_unused]] size_t i = 0; i < depth; i++)
{
for([[maybe_unused]] size_t j = 0; j < height; j++)
{
for([[maybe_unused]] size_t k = 0; k < width; k++)
{
int offset = i * width * height + j * width + k;
if(hData[offset] != hOutputData[offset])
{
std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----"
<< hOutputData[offset] << " ";
failed("mistmatch at:%d %d %d", i, j, k);
}
}
}
}
}
template <typename T>
void
checkArray(T input, T output, size_t height, size_t width)
{
for(size_t i = 0; i < height; i++)
{
for(size_t j = 0; j < width; j++)
{
int offset = i * width + j;
if(input[offset] != output[offset])
{
std::cerr << '[' << i << ',' << j << ',' << "]:" << input[offset] << "----"
<< output[offset] << " ";
failed("mistmatch at:%d %d", i, j);
}
}
}
}
template <typename T>
__global__ void
vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for(size_t i = offset; i < NELEM; i += stride)
{
C_d[i] = A_d[i] + B_d[i];
}
}
template <typename T>
__global__ void
vectorADDReverse(const T* A_d, const T* B_d, T* C_d, size_t NELEM)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for(int64_t i = NELEM - stride + offset; i >= 0; i -= stride)
{
C_d[i] = A_d[i] + B_d[i];
}
}
template <typename T>
__global__ void
addCount(const T* A_d, T* C_d, size_t NELEM, int count)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
// Deliberately do this in an inefficient way to increase kernel runtime
for(int i = 0; i < count; i++)
{
for(size_t i = offset; i < NELEM; i += stride)
{
C_d[i] = A_d[i] + (T) count;
}
}
}
template <typename T>
__global__ void
addCountReverse(const T* A_d, T* C_d, int64_t NELEM, int count)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
// Deliberately do this in an inefficient way to increase kernel runtime
for(int i = 0; i < count; i++)
{
for(int64_t i = NELEM - stride + offset; i >= 0; i -= stride)
{
C_d[i] = A_d[i] + (T) count;
}
}
}
template <typename T>
__global__ void
memsetReverse(T* C_d, T val, int64_t NELEM)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for(int64_t i = NELEM - stride + offset; i >= 0; i -= stride)
{
C_d[i] = val;
}
}
template <typename T>
void
setDefaultData(size_t numElements, T* A_h, T* B_h, T* C_h)
{
// Initialize the host data:
for(size_t i = 0; i < numElements; i++)
{
if(A_h) (A_h)[i] = 3.146f + i; // Pi
if(B_h) (B_h)[i] = 1.618f + i; // Phi
if(C_h) (C_h)[i] = 0.0f + i;
}
}
template <typename T>
void
initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost = false)
{
size_t Nbytes = N * sizeof(T);
if(usePinnedHost)
{
if(A_h)
{
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(A_h), Nbytes));
}
if(B_h)
{
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(B_h), Nbytes));
}
if(C_h)
{
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(C_h), Nbytes));
}
}
else
{
if(A_h)
{
*A_h = (T*) malloc(Nbytes);
HIPASSERT(*A_h != NULL);
}
if(B_h)
{
*B_h = (T*) malloc(Nbytes);
HIPASSERT(*B_h != NULL);
}
if(C_h)
{
*C_h = (T*) malloc(Nbytes);
HIPASSERT(*C_h != NULL);
}
}
setDefaultData(N, A_h ? *A_h : NULL, B_h ? *B_h : NULL, C_h ? *C_h : NULL);
}
template <typename T>
void
initArrays(T** A_d,
T** B_d,
T** C_d,
T** A_h,
T** B_h,
T** C_h,
size_t N,
bool usePinnedHost = false)
{
size_t Nbytes = N * sizeof(T);
if(A_d)
{
HIPCHECK(hipMalloc(A_d, Nbytes));
}
if(B_d)
{
HIPCHECK(hipMalloc(B_d, Nbytes));
}
if(C_d)
{
HIPCHECK(hipMalloc(C_d, Nbytes));
}
initArraysForHost(A_h, B_h, C_h, N, usePinnedHost);
}
template <typename T>
void
freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost)
{
if(usePinnedHost)
{
if(A_h)
{
HIPCHECK(hipHostFree(A_h));
}
if(B_h)
{
HIPCHECK(hipHostFree(B_h));
}
if(C_h)
{
HIPCHECK(hipHostFree(C_h));
}
}
else
{
if(A_h)
{
free(A_h);
}
if(B_h)
{
free(B_h);
}
if(C_h)
{
free(C_h);
}
}
}
template <typename T>
void
freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost)
{
if(A_d)
{
HIPCHECK(hipFree(A_d));
}
if(B_d)
{
HIPCHECK(hipFree(B_d));
}
if(C_d)
{
HIPCHECK(hipFree(C_d));
}
freeArraysForHost(A_h, B_h, C_h, usePinnedHost);
}
# if defined(__HIP_PLATFORM_AMD__)
template <typename T>
void
initArrays2DPitch(T** A_d,
T** B_d,
T** C_d,
size_t* pitch_A,
size_t* pitch_B,
size_t* pitch_C,
size_t numW,
size_t numH)
{
if(A_d)
{
HIPCHECK(hipMallocPitch((void**) A_d, pitch_A, numW * sizeof(T), numH));
}
if(B_d)
{
HIPCHECK(hipMallocPitch((void**) B_d, pitch_B, numW * sizeof(T), numH));
}
if(C_d)
{
HIPCHECK(hipMallocPitch((void**) C_d, pitch_C, numW * sizeof(T), numH));
}
HIPASSERT(*pitch_A == *pitch_B);
HIPASSERT(*pitch_A == *pitch_C)
}
inline void
initHIPArrays(hipArray** A_d,
hipArray** B_d,
hipArray** C_d,
const hipChannelFormatDesc* desc,
const size_t numW,
const size_t numH,
const unsigned int flags)
{
if(A_d)
{
HIPCHECK(hipMallocArray(A_d, desc, numW, numH, flags));
}
if(B_d)
{
HIPCHECK(hipMallocArray(B_d, desc, numW, numH, flags));
}
if(C_d)
{
HIPCHECK(hipMallocArray(C_d, desc, numW, numH, flags));
}
}
# endif
// Assumes C_h contains vector add of A_h + B_h
// Calls the test "failed" macro if a mismatch is detected.
template <typename T>
size_t
checkVectorADD(T* A_h,
T* B_h,
T* result_H,
size_t N,
bool expectMatch = true,
bool reportMismatch = true)
{
size_t mismatchCount = 0;
size_t firstMismatch = 0;
size_t mismatchesToPrint = 10;
for(size_t i = 0; i < N; i++)
{
T expected = A_h[i] + B_h[i];
if(result_H[i] != expected)
{
if(mismatchCount == 0)
{
firstMismatch = i;
}
mismatchCount++;
if((mismatchCount <= mismatchesToPrint) && expectMatch)
{
std::cout << std::fixed << std::setprecision(32);
std::cout << "At " << i << std::endl;
std::cout << " Computed:" << result_H[i] << std::endl;
std::cout << " Expected:" << expected << std::endl;
}
}
}
if(reportMismatch)
{
if(expectMatch)
{
if(mismatchCount)
{
failed("%zu mismatches ; first at index:%zu\n", mismatchCount, firstMismatch);
}
}
else
{
if(mismatchCount == 0)
{
failed("expected mismatches but did not detect any!");
}
}
}
return mismatchCount;
}
// Assumes C_h contains vector add of A_h + B_h
// Calls the test "failed" macro if a mismatch is detected.
template <typename T>
void
checkTest(T* expected_H, T* result_H, size_t N, bool expectMatch = true)
{
size_t mismatchCount = 0;
size_t firstMismatch = 0;
size_t mismatchesToPrint = 10;
for(size_t i = 0; i < N; i++)
{
if(result_H[i] != expected_H[i])
{
if(mismatchCount == 0)
{
firstMismatch = i;
}
mismatchCount++;
if((mismatchCount <= mismatchesToPrint) && expectMatch)
{
std::cout << std::fixed << std::setprecision(32);
std::cout << "At " << i << std::endl;
std::cout << " Computed:" << result_H[i] << std::endl;
std::cout << " Expected:" << expected_H[i] << std::endl;
}
}
}
if(expectMatch)
{
if(mismatchCount)
{
fprintf(stderr, "%zu mismatches ; first at index:%zu\n", mismatchCount, firstMismatch);
// failed("%zu mismatches ; first at index:%zu\n", mismatchCount, firstMismatch);
}
}
else
{
if(mismatchCount == 0)
{
failed("expected mismatches but did not detect any!");
}
}
}
//---
struct Pinned
{
static const bool isPinned = true;
static const char* str() { return "Pinned"; };
static void* Alloc(size_t sizeBytes)
{
void* p;
HIPCHECK(hipHostMalloc((void**) &p, sizeBytes));
return p;
};
};
//---
struct Unpinned
{
static const bool isPinned = false;
static const char* str() { return "Unpinned"; };
static void* Alloc(size_t sizeBytes)
{
void* p = malloc(sizeBytes);
HIPASSERT(p);
return p;
};
};
struct Memcpy
{
static const char* str() { return "Memcpy"; };
};
struct MemcpyAsync
{
static const char* str() { return "MemcpyAsync"; };
};
template <typename C>
struct MemTraits;
template <>
struct MemTraits<Memcpy>
{
static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t)
{
HIPCHECK(hipMemcpy(dest, src, sizeBytes, kind));
}
};
template <>
struct MemTraits<MemcpyAsync>
{
static void Copy(void* dest,
const void* src,
size_t sizeBytes,
hipMemcpyKind kind,
hipStream_t stream)
{
HIPCHECK(hipMemcpyAsync(dest, src, sizeBytes, kind, stream));
}
};
inline bool
isImageSupported()
{
int imageSupport = 1;
# ifdef __HIP_PLATFORM_AMD__
HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, p_gpuDevice));
# endif
return imageSupport != 0;
}
}; // namespace HipTest
// This must be called in the beginning of image test app's main() to indicate whether image
// is supported.
# define checkImageSupport() \
if(!HipTest::isImageSupported()) \
{ \
printf("Texture is not support on the device. Skipped.\n"); \
passed(); \
}
#endif //__cplusplus
// Function to determine if the device is of gfx11 architecture
bool
IsGfx11();
benchmark/source/bin/rocprofv3-benchmark.py
Executable
+1487
View File
File diff suppressed because it is too large Load Diff
benchmark/source/lib/CMakeLists.txt
+11
View File
@@ -0,0 +1,11 @@
#
#
#
# install the downloaded timem libraries
install(
DIRECTORY ${PROJECT_BINARY_DIR}/lib64/
DESTINATION ${CMAKE_INSTALL_LIB64DIR}
FILES_MATCHING
PATTERN "*libtimem.*"
PATTERN "*/timemory/*")
benchmark/source/share/CMakeLists.txt
+5
View File
@@ -0,0 +1,5 @@
#
#
#
add_subdirectory(rocprofiler-sdk)
benchmark/source/share/rocprofiler-sdk/CMakeLists.txt
+12
View File
@@ -0,0 +1,12 @@
#
#
#
set(DATA_FILES benchmark_tables.sql benchmark_views.sql)
foreach(_FILE ${DATA_FILES})
configure_file(${_FILE} ${CMAKE_DATAROOT_OUTPUT_DIRECTORY}/rocprofiler-sdk/${_FILE}
COPYONLY)
endforeach()
install(FILES ${DATA_FILES} DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/rocprofiler-sdk)
benchmark/source/share/rocprofiler-sdk/benchmark_tables.sql
+137
View File
@@ -0,0 +1,137 @@
-- Application used for benchmarking
-- Columns such "hip_compiler_api", ..., "scratch_memory" are
-- the number of events in the given category, e.g. kernel_dispatch
-- represents the number of kernel dispatches in the app. These
-- can be approximate since for a given application, the exact
-- count may vary.
CREATE TABLE IF NOT EXISTS
`benchmarked_app` (
id INT PRIMARY KEY AUTO_INCREMENT UNIQUE,
hash_id TEXT NOT NULL,
md5sum TEXT NOT NULL,
revision INT DEFAULT 0,
command JSON NOT NULL,
compiler_id TEXT,
compiler_version TEXT,
library_arch TEXT,
system_name TEXT,
system_processor TEXT,
system_version TEXT,
threads INT,
hip_compiler_api INT,
hip_runtime_api INT,
hsa_api INT,
kernel_dispatch INT,
marker_api INT,
memory_allocation INT,
memory_copy INT,
ompt INT,
rccl_api INT,
rocdecode_api INT,
rocjpeg_api INT,
scratch_memory INT,
environment JSON DEFAULT ("{}")
);
-- rocprofiler-sdk used for benchmarking
CREATE TABLE IF NOT EXISTS
`benchmarked_sdk` (
id INT PRIMARY KEY AUTO_INCREMENT UNIQUE,
hash_id TEXT NOT NULL,
version_major INT NOT NULL,
version_minor INT NOT NULL,
version_patch INT NOT NULL,
soversion INT NOT NULL,
compiler_id TEXT NOT NULL,
compiler_version TEXT NOT NULL,
git_revision TEXT NOT NULL,
library_arch TEXT NOT NULL,
system_name TEXT NOT NULL,
system_processor TEXT NOT NULL,
system_version TEXT NOT NULL
);
-- rocprofiler-sdk used for benchmarking
CREATE TABLE IF NOT EXISTS
`benchmark_config` (
id INT PRIMARY KEY AUTO_INCREMENT UNIQUE,
hash_id TEXT NOT NULL,
sdk_id INT,
label TEXT, -- name identifier
benchmark_mode TEXT CHECK (
benchmark_mode IN (
"baseline",
"disabled-sdk-contexts",
"sdk-buffer-overhead",
"sdk-callback-overhead",
"tool-runtime-overhead"
)
) NOT NULL,
kernel_rename INT,
group_by_queue INT,
kernel_trace INT,
hsa_trace INT,
hip_runtime_trace INT,
hip_compiler_trace INT,
marker_trace INT,
memory_copy_trace INT,
memory_allocation_trace INT,
scratch_memory_trace INT,
dispatch_counter_collection INT,
rccl_trace INT,
rocdecode_trace INT,
rocjpeg_trace INT,
pmc_counters JSON DEFAULT ("[]"),
pc_sampling_host_trap INT,
pc_sampling_stocastic INT,
advanced_thread_trace INT,
--
-- Eventually, we will create tables for storing the subconfigurations for pc sampling and ATT
--
-- pc_sampling_host_trap_config_id INT,
-- pc_sampling_stocastic_config_id INT,
-- advanced_thread_trace_config_id INT,
-- FOREIGN KEY (pc_sampling_host_trap_config_id) REFERENCES benchmark_pc_sampling_host_trap_config (id) ON UPDATE CASCADE,
-- FOREIGN KEY (pc_sampling_stocastic_config_id) REFERENCES benchmark_pc_sampling_stocastic_config (id) ON UPDATE CASCADE,
-- FOREIGN KEY (advanced_thread_trace_config_id) REFERENCES benchmark_advanced_thread_trace_config (id) ON UPDATE CASCADE
FOREIGN KEY (sdk_id) REFERENCES benchmarked_sdk (id) ON UPDATE CASCADE
);
-- metrics for the benchmark
CREATE TABLE IF NOT EXISTS
`benchmark_metrics` (
id INT PRIMARY KEY AUTO_INCREMENT UNIQUE,
app_id INT NOT NULL,
cfg_id INT NOT NULL,
sdk_id INT,
executed_at TIMESTAMP NOT NULL,
wall_time DOUBLE NOT NULL,
cpu_time DOUBLE NOT NULL,
cpu_util DOUBLE NOT NULL,
peak_rss DOUBLE NOT NULL,
page_rss DOUBLE NOT NULL,
virtual_memory DOUBLE NOT NULL,
major_page_faults BIGINT NOT NULL,
minor_page_faults BIGINT NOT NULL,
priority_context_switches BIGINT NOT NULL,
voluntary_context_switches BIGINT NOT NULL,
FOREIGN KEY (app_id) REFERENCES benchmarked_app (id) ON UPDATE CASCADE,
FOREIGN KEY (cfg_id) REFERENCES benchmark_config (id) ON UPDATE CASCADE,
FOREIGN KEY (sdk_id) REFERENCES benchmarked_sdk (id) ON UPDATE CASCADE
);
CREATE TABLE IF NOT EXISTS
`benchmark_statistics` (
id INT PRIMARY KEY AUTO_INCREMENT UNIQUE,
app_id INT NOT NULL,
cfg_id INT NOT NULL,
sdk_id INT,
metric_name TEXT NOT NULL,
metric_unit TEXT NOT NULL,
count INT NOT NULL,
sum DOUBLE NOT NULL,
mean DOUBLE NOT NULL,
min DOUBLE NOT NULL,
max DOUBLE NOT NULL,
std_dev DOUBLE
);
benchmark/source/share/rocprofiler-sdk/benchmark_views.sql
+75
View File
@@ -0,0 +1,75 @@
-- Analysis views used for benchmarking
CREATE VIEW IF NOT EXISTS
`benchmark_analysis_{{metric}}` AS
WITH
baseline AS (
SELECT
*
FROM
benchmark_statistics BL
WHERE
BL.sdk_id IS NULL
AND BL.metric_name = "{{metric}}"
)
SELECT
ST.id,
ST.app_id,
ST.cfg_id,
ST.sdk_id,
BS.git_revision,
BA.command,
ST.metric_name,
ST.metric_unit,
ST.count,
ST.mean,
ST.std_dev AS `+/-`,
BL.mean AS baseline_mean,
BL.std_dev AS `+/- (baseline)`,
((ST.mean - BL.mean) / BL.mean) * 100 AS `overhead (%)`,
BC.benchmark_mode,
BC.label AS benchmark_label
FROM
benchmark_statistics ST
JOIN benchmark_config BC ON BC.id = ST.cfg_id
JOIN benchmarked_sdk BS ON BS.id = ST.sdk_id
JOIN benchmarked_app BA ON BA.id = ST.app_id
JOIN baseline BL ON (
BL.app_id = ST.app_id
AND BL.metric_name = ST.metric_name
)
WHERE
ST.metric_name = "{{metric}}"
AND ST.sdk_id IS NOT NULL
ORDER BY
`overhead (%)` DESC;
-- benchmarked_app without environment info
CREATE VIEW IF NOT EXISTS
`benchmarked_app_without_env` AS
SELECT
id,
hash_id,
md5sum,
revision,
command,
compiler_id,
compiler_version,
library_arch,
system_name,
system_processor,
system_version,
threads,
hip_compiler_api,
hip_runtime_api,
hsa_api,
kernel_dispatch,
marker_api,
memory_allocation,
memory_copy,
ompt,
rccl_api,
rocdecode_api,
rocjpeg_api,
scratch_memory
FROM
benchmarked_app;