From 97f8b7b1ec0acd1f78fc252004532ba27f13b2b3 Mon Sep 17 00:00:00 2001 From: ywang103-amd Date: Tue, 16 Sep 2025 10:17:23 -0400 Subject: [PATCH] change to single-kernel workload for pc_sampling tests (#955) --- .../sample/mat_mul_max.hip | 75 +++++++++++++++++++ .../rocprofiler-compute/tests/CMakeLists.txt | 6 ++ .../tests/test_profile_general.py | 10 +-- .../rocprofiler-compute/tests/test_utils.py | 2 +- 4 files changed, 87 insertions(+), 6 deletions(-) create mode 100644 projects/rocprofiler-compute/sample/mat_mul_max.hip diff --git a/projects/rocprofiler-compute/sample/mat_mul_max.hip b/projects/rocprofiler-compute/sample/mat_mul_max.hip new file mode 100644 index 0000000000..3f32e4776b --- /dev/null +++ b/projects/rocprofiler-compute/sample/mat_mul_max.hip @@ -0,0 +1,75 @@ +#include +#include +#include +#include + +#define TILE_SIZE 32 // Maximum block size: 32 x 32 = 1024 threads/block +#define N 4096 // Matrix size: 4096 x 4096 (~67M elements) + +__global__ void matMulKernel(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int width) { + __shared__ float tileA[TILE_SIZE][TILE_SIZE]; + __shared__ float tileB[TILE_SIZE][TILE_SIZE]; + + int row = blockIdx.y * TILE_SIZE + threadIdx.y; + int col = blockIdx.x * TILE_SIZE + threadIdx.x; + + float sum = 0.0f; + + for (int t = 0; t < width / TILE_SIZE; ++t) { + tileA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x]; + tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col]; + __syncthreads(); + + #pragma unroll + for (int i = 0; i < TILE_SIZE; ++i) + sum += tileA[threadIdx.y][i] * tileB[i][threadIdx.x]; + + __syncthreads(); + } + + C[row * width + col] = sum; +} + +int main() { + size_t size = N * N * sizeof(float); + + float *h_A = new float[N * N]; + float *h_B = new float[N * N]; + float *d_A, *d_B, *d_C; + + // Initialize matrices with dummy values + for (int i = 0; i < N * N; ++i) { + h_A[i] = static_cast(i % 100) * 0.01f; + h_B[i] = static_cast((i + 1) % 100) * 0.01f; + } + + hipMalloc(&d_A, size); + hipMalloc(&d_B, size); + hipMalloc(&d_C, size); + + hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice); + hipMemcpy(d_B, h_B, size, hipMemcpyHostToDevice); + + dim3 blockDim(TILE_SIZE, TILE_SIZE); // 32 x 32 = 1024 threads + dim3 gridDim(N / TILE_SIZE, N / TILE_SIZE); // 128 x 128 = 16,384 thread blocks + + std::cout << "Launching kernel with grid: (" << gridDim.x << ", " << gridDim.y + << ") and block: (" << blockDim.x << ", " << blockDim.y << ")\n"; + + auto start = std::chrono::high_resolution_clock::now(); + + matMulKernel<<>>(d_A, d_B, d_C, N); + hipDeviceSynchronize(); + + auto end = std::chrono::high_resolution_clock::now(); + std::chrono::duration elapsed = end - start; + std::cout << "Execution time: " << elapsed.count() << " seconds\n"; + + hipFree(d_A); + hipFree(d_B); + hipFree(d_C); + delete[] h_A; + delete[] h_B; + + return 0; +} \ No newline at end of file diff --git a/projects/rocprofiler-compute/tests/CMakeLists.txt b/projects/rocprofiler-compute/tests/CMakeLists.txt index 19978d8a99..a4b3b9d660 100644 --- a/projects/rocprofiler-compute/tests/CMakeLists.txt +++ b/projects/rocprofiler-compute/tests/CMakeLists.txt @@ -36,3 +36,9 @@ set_source_files_properties(${OCCUPANCY} PROPERTIES LANGUAGE HIP) add_executable(occupancy ${OCCUPANCY}) set_target_properties(occupancy PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR}/tests) + +set(MAT_MUL_MAX ../sample/mat_mul_max.hip) +set_source_files_properties(${MAT_MUL_MAX} PROPERTIES LANGUAGE HIP) +add_executable(mat_mul_max ${MAT_MUL_MAX}) +set_target_properties(mat_mul_max PROPERTIES RUNTIME_OUTPUT_DIRECTORY + ${CMAKE_SOURCE_DIR}/tests) diff --git a/projects/rocprofiler-compute/tests/test_profile_general.py b/projects/rocprofiler-compute/tests/test_profile_general.py index 96d50abccc..07c7e247e2 100644 --- a/projects/rocprofiler-compute/tests/test_profile_general.py +++ b/projects/rocprofiler-compute/tests/test_profile_general.py @@ -60,7 +60,7 @@ CHIP_IDS = { config = {} config["kernel_name_1"] = "vecCopy" config["app_1"] = ["./tests/vcopy", "-n", "1048576", "-b", "256", "-i", "3"] -config["app_occupancy"] = ["./tests/occupancy"] +config["app_mat_mul_max"] = ["./tests/mat_mul_max"] config["cleanup"] = True config["COUNTER_LOGGING"] = False config["METRIC_COMPARE"] = False @@ -1826,10 +1826,10 @@ def test_pc_sampling_host_trap(binary_handler_profile_rocprof_compute): options, check_success=True, roof=False, - app_name="app_occupancy", + app_name="app_mat_mul_max", ) - file_dict = test_utils.check_csv_files(workload_dir, num_devices, num_kernels) + file_dict = test_utils.check_csv_files(workload_dir, num_devices, 1) assert sorted(list(file_dict.keys())) == sorted(PC_SAMPLING_HOST_TRAP_FILES) validate(inspect.stack()[0][3], workload_dir, file_dict) @@ -1858,10 +1858,10 @@ def test_pc_sampling_stochastic(binary_handler_profile_rocprof_compute): options, check_success=True, roof=False, - app_name="app_occupancy", + app_name="app_mat_mul_max", ) - file_dict = test_utils.check_csv_files(workload_dir, num_devices, num_kernels) + file_dict = test_utils.check_csv_files(workload_dir, num_devices, 1) assert sorted(list(file_dict.keys())) == sorted(PC_SAMPLING_STOCHASTIC_FILES) validate(inspect.stack()[0][3], workload_dir, file_dict) diff --git a/projects/rocprofiler-compute/tests/test_utils.py b/projects/rocprofiler-compute/tests/test_utils.py index fb45e8e26b..2b950fe322 100644 --- a/projects/rocprofiler-compute/tests/test_utils.py +++ b/projects/rocprofiler-compute/tests/test_utils.py @@ -157,7 +157,7 @@ def check_csv_files(output_dir, num_devices, num_kernels): file_dict[file] = pd.read_csv(output_dir + "/" + file) if "roofline" in file: assert len(file_dict[file].index) >= num_devices - elif not "sysinfo" in file: + elif "sysinfo" not in file and "ps_file" not in file: assert len(file_dict[file].index) >= num_kernels elif file.endswith(".pdf"): file_dict[file] = "pdf"