diff --git a/CMakeLists.txt b/CMakeLists.txt index b7a390c8..af6302d2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -98,6 +98,7 @@ add_caliper_option(BUILD_TESTING "Build continuous integration app and unit tes add_caliper_option(BUILD_DOCS "Build Caliper documentation" FALSE) add_caliper_option(RUN_MPI_TESTS "Run MPI tests (only applicable with BUILD_TESTING=On)" TRUE) +add_caliper_option(RUN_HIP_TESTS "Run HIP tests when BUILD_TESTING=On" TRUE) # allow disabling cmake/pkg-config files and headers for when caliper is a subproject add_caliper_option(INSTALL_CONFIG "Install cmake and pkg-config files" TRUE) diff --git a/test/ci_app_tests/CMakeLists.txt b/test/ci_app_tests/CMakeLists.txt index 865765be..bb06d1bf 100644 --- a/test/ci_app_tests/CMakeLists.txt +++ b/test/ci_app_tests/CMakeLists.txt @@ -98,6 +98,15 @@ if (CALIPER_HAVE_MPI) endif() endif() +if (CALIPER_HAVE_ROCPROFILER) + enable_language(HIP) + add_executable(vectoradd vectoradd.hip) + target_link_libraries(vectoradd PUBLIC rocprofiler-sdk-roctx::rocprofiler-sdk-roctx caliper) + if (RUN_HIP_TESTS) + list(APPEND PYTHON_SCRIPTS test_rocprofiler.py) + endif() +endif() + if (WITH_FORTRAN) foreach(app ${CALIPER_CI_Fortran_TEST_APPS}) add_executable(${app} ${app}.f) diff --git a/test/ci_app_tests/test_rocprofiler.py b/test/ci_app_tests/test_rocprofiler.py new file mode 100644 index 00000000..17fb3bbe --- /dev/null +++ b/test/ci_app_tests/test_rocprofiler.py @@ -0,0 +1,62 @@ +# HIP tests + +import io +import unittest + +import caliperreader +import calipertest as cat + +class CaliperRocmServicesTest(unittest.TestCase): + """ Caliper test class for linux-specific services """ + + def test_rocm_activity_profile(self): + target_cmd = [ './vectoradd', 'rocm-activity-profile,profile.roctx,rocm.counters=SQ_WAVES_sum,output=stdout' ] + env = { 'HIP_LAUNCH_BLOCKING': '1' } + + out,_ = cat.run_test(target_cmd, env) + snapshots,_ = caliperreader.read_caliper_contents(io.StringIO(out.decode())) + + self.assertTrue(len(snapshots) > 1) + + self.assertTrue(cat.has_snapshot_with_keys( + snapshots, { 'rocm.activity', + 'rocm.kernel.name', + 'scale#sum#rocm.activity.duration', + 'path', + 'rocm.marker' } + )) + self.assertTrue(cat.has_snapshot_with_attributes( + snapshots, { 'rocm.activity': 'KERNEL_DISPATCH_COMPLETE', + 'sum#sum#rocm.activity.count': '1', + 'path': ['main', 'vectoradd', 'hipLaunchKernel'] } + )) + self.assertTrue(cat.has_snapshot_with_attributes( + snapshots, { 'rocm.activity': 'MEMORY_COPY_DEVICE_TO_HOST', + 'sum#sum#rocm.activity.count': '1', + 'sum#sum#rocm.bytes': '4194304', + 'path': ['main', 'copy_d2h', 'hipMemcpy'] } + )) + + rec = cat.get_snapshot_with_keys(snapshots, ['path', 'sum#sum#rocm.SQ_WAVES_sum']) + self.assertIsNotNone(rec) + self.assertEqual(int(rec['sum#sum#rocm.SQ_WAVES_sum']), 16384) + self.assertEqual(rec['path'], ['main', 'vectoradd', 'hipLaunchKernel']) + + def test_rocm_opts(self): + target_cmd = [ './vectoradd', 'runtime-profile,profile.roctx,rocm.gputime,output=stdout' ] + env = { 'HIP_LAUNCH_BLOCKING': '1' } + + out,_ = cat.run_test(target_cmd, env) + snapshots,_ = caliperreader.read_caliper_contents(io.StringIO(out.decode())) + + self.assertTrue(len(snapshots) > 1) + + self.assertTrue(cat.has_snapshot_with_keys(snapshots, { 'iscale#t.gpu.r', 'path' })) + + rec = cat.get_snapshot_with_keys(snapshots, ['path', 'scale#t.gpu.r']) + self.assertIsNotNone(rec) + self.assertGreater(float(rec['scale#t.gpu.r']), 0.0) + + +if __name__ == "__main__": + unittest.main() diff --git a/test/ci_app_tests/vectoradd.hip b/test/ci_app_tests/vectoradd.hip new file mode 100644 index 00000000..8d254c52 --- /dev/null +++ b/test/ci_app_tests/vectoradd.hip @@ -0,0 +1,171 @@ +/* +Copyright (c) 2015-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 +#include +#include +#include +#include + +#include "hip/hip_runtime.h" + +#include + +#include +#include + +// Macro for checking GPU API return values +#define hipCheck(call) \ +do{ \ + hipError_t gpuErr = call; \ + if(hipSuccess != gpuErr){ \ + printf("GPU API Error - %s:%d: '%s'\n", __FILE__, __LINE__, hipGetErrorString(gpuErr)); \ + exit(1); \ + } \ +}while(0) + +#define WIDTH 1024 +#define HEIGHT 1024 + +#define NUM (WIDTH*HEIGHT) + +#define THREADS_PER_BLOCK_X 16 +#define THREADS_PER_BLOCK_Y 16 +#define THREADS_PER_BLOCK_Z 1 + +__global__ void +vectoradd_float(float* __restrict__ a, const float* __restrict__ b, const float* __restrict__ c, int width, int height) + + { + + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = b[i] + c[i]; + } + + } + +using namespace std; + +int main(int argc, char* argv[]) { + + float* hostA; + float* hostB; + float* hostC; + + float* deviceA; + float* deviceB; + float* deviceC; + + hipDeviceProp_t devProp; + hipCheck(hipGetDeviceProperties(&devProp, 0)); + cerr << " System minor " << devProp.minor << endl; + cerr << " System major " << devProp.major << endl; + cerr << " agent prop name " << devProp.name << endl; + cerr << "hip Device prop succeeded " << endl ; + + cali::ConfigManager mgr; + if (argc > 1) + mgr.add(argv[1]); + if (mgr.error()) { + cerr << "cali::ConfigManager: " << mgr.error_msg() << endl; + return 1; + } + + cali_init(); + mgr.start(); + + roctxRangePush("main"); + + int i; + int errors; + + hostA = (float*)malloc(NUM * sizeof(float)); + hostB = (float*)malloc(NUM * sizeof(float)); + hostC = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (float)i; + hostC[i] = (float)i*100.0f; + } + + roctxRangePush("copy_h2d"); + + hipCheck(hipMalloc((void**)&deviceA, NUM * sizeof(float))); + hipCheck(hipMalloc((void**)&deviceB, NUM * sizeof(float))); + hipCheck(hipMalloc((void**)&deviceC, NUM * sizeof(float))); + + hipCheck(hipMemcpy(deviceB, hostB, NUM*sizeof(float), hipMemcpyHostToDevice)); + hipCheck(hipMemcpy(deviceC, hostC, NUM*sizeof(float), hipMemcpyHostToDevice)); + + roctxRangePop(); + roctxRangePush("vectoradd"); + + //hipLaunchKernelGGL(vectoradd_float, + // dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + // dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + // 0, 0, + // deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + dim3 grid(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y); + dim3 block(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y); + vectoradd_float<<>>(deviceA, deviceB, deviceC, WIDTH, HEIGHT); + + roctxRangePop(); + roctxRangePush("copy_d2h"); + + hipCheck(hipMemcpy(hostA, deviceA, NUM*sizeof(float), hipMemcpyDeviceToHost)); + + roctxRangePop(); + roctxRangePush("verify"); + + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors!=0) { + cerr << "FAILED: " << errors << " errors" << endl; + } else { + cerr << "PASSED!" << endl; + } + + roctxRangePop(); + + hipCheck(hipFree(deviceA)); + hipCheck(hipFree(deviceB)); + hipCheck(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + //hipResetDefaultAccelerator(); + roctxRangePop(); // main + mgr.flush(); + + return errors; +}