diff --git a/tools/LaunchBench/Common.hpp b/tools/JitterBench/Common.hpp similarity index 98% rename from tools/LaunchBench/Common.hpp rename to tools/JitterBench/Common.hpp index 2fd8aa081e..8fad0ac3c1 100644 --- a/tools/LaunchBench/Common.hpp +++ b/tools/JitterBench/Common.hpp @@ -21,7 +21,6 @@ THE SOFTWARE. */ #include -#include #define HIP_CALL(cmd) \ do { \ diff --git a/tools/JitterBench/Compatibility.hpp b/tools/JitterBench/Compatibility.hpp new file mode 100644 index 0000000000..25efd81a09 --- /dev/null +++ b/tools/JitterBench/Compatibility.hpp @@ -0,0 +1,83 @@ +/* +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 + +#if defined(__NVCC__) + +#include + +// ROCm specific +#define wall_clock64 clock64 +#define gcnArchName name + +// Datatypes +#define hipDeviceProp_t cudaDeviceProp +#define hipError_t cudaError_t +#define hipEvent_t cudaEvent_t +#define hipStream_t cudaStream_t + +// Enumerations +#define hipDeviceAttributeClockRate cudaDevAttrClockRate +#define hipDeviceAttributeMaxSharedMemoryPerMultiprocessor cudaDevAttrMaxSharedMemoryPerMultiprocessor +#define hipDeviceAttributeMultiprocessorCount cudaDevAttrMultiProcessorCount +#define hipErrorPeerAccessAlreadyEnabled cudaErrorPeerAccessAlreadyEnabled +#define hipFuncCachePreferShared cudaFuncCachePreferShared +#define hipMemcpyDefault cudaMemcpyDefault +#define hipMemcpyDeviceToHost cudaMemcpyDeviceToHost +#define hipMemcpyHostToDevice cudaMemcpyHostToDevice +#define hipSuccess cudaSuccess + +// Functions +#define hipDeviceCanAccessPeer cudaDeviceCanAccessPeer +#define hipDeviceEnablePeerAccess cudaDeviceEnablePeerAccess +#define hipDeviceGetAttribute cudaDeviceGetAttribute +#define hipDeviceGetPCIBusId cudaDeviceGetPCIBusId +#define hipDeviceSetCacheConfig cudaDeviceSetCacheConfig +#define hipDeviceSynchronize cudaDeviceSynchronize +#define hipEventCreate cudaEventCreate +#define hipEventDestroy cudaEventDestroy +#define hipEventElapsedTime cudaEventElapsedTime +#define hipEventRecord cudaEventRecord +#define hipFree cudaFree +#define hipGetDeviceCount cudaGetDeviceCount +#define hipGetDeviceProperties cudaGetDeviceProperties +#define hipGetErrorString cudaGetErrorString +#define hipHostFree cudaFreeHost +#define hipHostMalloc cudaMallocHost +#define hipMalloc cudaMalloc +#define hipMemcpy cudaMemcpy +#define hipMemcpyAsync cudaMemcpyAsync +#define hipMemset cudaMemset +#define hipMemsetAsync cudaMemsetAsync +#define hipSetDevice cudaSetDevice +#define hipStreamCreate cudaStreamCreate +#define hipStreamDestroy cudaStreamDestroy +#define hipStreamSynchronize cudaStreamSynchronize + +#else + +#include +#include +#include + +#endif diff --git a/tools/JitterBench/GetClosestNumaNode.hpp b/tools/JitterBench/GetClosestNumaNode.hpp new file mode 100644 index 0000000000..bb184ea71c --- /dev/null +++ b/tools/JitterBench/GetClosestNumaNode.hpp @@ -0,0 +1,151 @@ +/* +Copyright (c) 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 + +// Helper macro for checking HSA calls +#define HSA_CHECK(cmd) \ + do { \ + hsa_status_t error = (cmd); \ + if (error != HSA_STATUS_SUCCESS) { \ + const char* errString = NULL; \ + hsa_status_string(error, &errString); \ + std::cerr << "Encountered HSA error (" << errString << ") at line " \ + << __LINE__ << " in file " << __FILE__ << "\n"; \ + exit(-1); \ + } \ + } while (0) + +// Structure to hold HSA agent information +#if !defined(__NVCC__) +struct AgentData +{ + bool isInitialized; + std::vector cpuAgents; + std::vector gpuAgents; + std::vector closestNumaNode; +}; + +// Simple callback function to return any memory pool for an agent +hsa_status_t MemPoolInfoCallback(hsa_amd_memory_pool_t pool, void *data) +{ + hsa_amd_memory_pool_t* poolData = reinterpret_cast(data); + + // Check memory pool flags + uint32_t poolFlags; + HSA_CHECK(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &poolFlags)); + + // Only consider coarse-grained pools + if (!(poolFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)) return HSA_STATUS_SUCCESS; + + *poolData = pool; + return HSA_STATUS_SUCCESS; +} + +// Callback function to gather HSA agent information +hsa_status_t AgentInfoCallback(hsa_agent_t agent, void* data) +{ + AgentData* agentData = reinterpret_cast(data); + + // Get the device type + hsa_device_type_t deviceType; + HSA_CHECK(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &deviceType)); + if (deviceType == HSA_DEVICE_TYPE_CPU) + agentData->cpuAgents.push_back(agent); + if (deviceType == HSA_DEVICE_TYPE_GPU) + { + agentData->gpuAgents.push_back(agent); + agentData->closestNumaNode.push_back(0); + } + + return HSA_STATUS_SUCCESS; +} + +AgentData& GetAgentData() +{ + static AgentData agentData = {}; + + if (!agentData.isInitialized) + { + agentData.isInitialized = true; + + // Add all detected agents to the list + HSA_CHECK(hsa_iterate_agents(AgentInfoCallback, &agentData)); + + // Loop over each GPU + for (uint32_t i = 0; i < agentData.gpuAgents.size(); i++) + { + // Collect memory pool + hsa_amd_memory_pool_t pool; + HSA_CHECK(hsa_amd_agent_iterate_memory_pools(agentData.gpuAgents[i], MemPoolInfoCallback, &pool)); + + // Loop over each CPU agent and check distance + int bestDistance = -1; + for (uint32_t j = 0; j < agentData.cpuAgents.size(); j++) + { + // Determine number of hops from GPU memory pool to CPU agent + uint32_t hops = 0; + HSA_CHECK(hsa_amd_agent_memory_pool_get_info(agentData.cpuAgents[j], + pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, + &hops)); + // Gather link info + hsa_amd_memory_pool_link_info_t* link_info = + (hsa_amd_memory_pool_link_info_t *)malloc(hops * sizeof(hsa_amd_memory_pool_link_info_t)); + HSA_CHECK(hsa_amd_agent_memory_pool_get_info(agentData.cpuAgents[j], + pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO, + link_info)); + int numaDist = 0; + for (int k = 0; k < hops; k++) + { + numaDist += link_info[k].numa_distance; + } + if (bestDistance == -1 || numaDist < bestDistance) + { + agentData.closestNumaNode[i] = j; + bestDistance = numaDist; + } + free(link_info); + } + } + } + return agentData; +} +#endif + +// Returns closest CPU NUMA node to provided GPU +// NOTE: This assumes HSA GPU indexing is similar to HIP GPU indexing +int GetClosestNumaNode(int gpuIdx) +{ +#if defined(__NVCC__) + return -1; +#else + AgentData& agentData = GetAgentData(); + if (gpuIdx < 0 || gpuIdx >= agentData.closestNumaNode.size()) + { + printf("[ERROR] GPU index out is out of bounds\n"); + exit(1); + } + return agentData.closestNumaNode[gpuIdx]; +#endif +} diff --git a/tools/JitterBench/JitterBench.cpp b/tools/JitterBench/JitterBench.cpp new file mode 100644 index 0000000000..5aa9306a24 --- /dev/null +++ b/tools/JitterBench/JitterBench.cpp @@ -0,0 +1,522 @@ +/* +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 +#include +#include +#include +#include +#include +#include +#include + +#include "Common.hpp" +#include "Compatibility.hpp" +#include "GetClosestNumaNode.hpp" +#include "Timeline.hpp" + +#ifdef MPI_SUPPORT +#include +#endif + +struct SyncData +{ + uint64_t cpuStart; + uint64_t cpuStop; + int32_t xccId; + uint64_t gpuStart; + uint64_t gpuStop; +}; + +enum +{ + HOST_START = 0, + HOST_RETURN = 1, + DEV_START = 2, + HOST_ABORT = 3, + DEV_STOP = 4, + HOST_STOP = 5, + KERNEL_CPUTIME = 6, + KERNEL_GPUTIME = 7, + KERNEL_TIMEDIFF = 8, + NUM_COLUMNS = 9 +} Columns; + +bool printCol[NUM_COLUMNS] = +{ + false, + false, + true, + false, + true, + false, + true, + true, + true +}; + +#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_ACQUIRE) +#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_RELEASE) + +__global__ void SyncKernel(uint64_t* cpuTime, uint32_t* abortFlag, SyncData* syncData) +{ + SyncData sd; + // Only first thread in threadblock participates + if (threadIdx.x != 0) return; + + // Collect timestamp upon kernel entry + sd.cpuStart = LOAD(cpuTime); + sd.gpuStart = wall_clock64(); + + // Wait for abort flag to be modified + while (!LOAD(abortFlag)); + + // Collect timestamps after abort flag + sd.cpuStop = LOAD(cpuTime); + sd.gpuStop = wall_clock64(); + + // Save timestamps + GetXccId(sd.xccId); + syncData[blockIdx.x] = sd; +} + +void SetNumaNode(int numaId) +{ + // Move CPU thread to targeted NUMA node + if (numa_run_on_node(numaId)) + { + printf("[ERROR] Unable to migrate to NUMA node %d\n", numaId); + exit(1); + } + + // Set memory to allocate on targeted NUMA node + numa_set_preferred(numaId); +} + +void UpdateCpuTime(int const useNuma, int const numaId, uint64_t* cpuTimestamp, bool* abortThread) +{ + if (useNuma) SetNumaNode(numaId); + while (!LOAD(abortThread)) + { + // Unroll to increase update vs abort check ratio + #pragma unroll + for (int i = 0; i < 64; i++) + STORE(cpuTimestamp, std::chrono::steady_clock::now().time_since_epoch().count()); + } +} + +void HostMalloc(void** pinnedHostPtr, size_t size) +{ +#if !defined(__NVCC__) + HIP_CALL(hipHostMalloc(pinnedHostPtr, size, hipHostMallocNumaUser)); +#else + HIP_CALL(hipHostMalloc(pinnedHostPtr, size)); +#endif + memset(*pinnedHostPtr, 0, size); +} + +int main(int argc, char **argv) +{ + // Initialize MPI (if supported) and check for NUMA support +#ifdef MPI_SUPPORT + MPI_Init(&argc, &argv); +#endif + if (numa_available() == -1) + { + printf("[ERROR] NUMA library not supported. Check to see if libnuma has been installed on this system\n"); + exit(1); + } + + int numAvailableGpus; + HIP_CALL(hipGetDeviceCount(&numAvailableGpus)); + + // Figure out how many GPUs total / which GPU this process is responsible for + int numUsedGpus, numTotalGpus, rank; +#ifdef MPI_SUPPORT + numUsedGpus = 1; + MPI_Comm_rank(MPI_COMM_WORLD, &rank); + MPI_Comm_size(MPI_COMM_WORLD, &numTotalGpus); + if (numTotalGpus > numAvailableGpus) + { + if (rank == 0) printf("[ERROR] Machine only has %d devices but %d ranks requested\n", numAvailableGpus, numTotalGpus); + MPI_Abort(MPI_COMM_WORLD, -1); + } + if (rank == 0) + { + printf("Running MPI version with %d ranks\n", numTotalGpus); + } +#else + numUsedGpus = numTotalGpus = numAvailableGpus; + printf("Running OMP version with %d ranks\n", numTotalGpus); + rank = 0; +#endif + + // Collect arguments from commandline or environment variable + #define GETARG(IDX, STR, DEFAULT) \ + (argc > IDX ? atoi(argv[IDX]) : (getenv(STR) ? atoi(getenv(STR)) : DEFAULT)) + + int numBlocks = GETARG(1, "NUM_BLOCKS", 4); + int blockSize = GETARG(2, "BLOCKSIZE", 64); + int numUpdateThreads = GETARG(3, "NUM_UPDATERS", 0); + int useNuma = GETARG(4, "USE_NUMA", 1); + int numIterations = GETARG(5, "NUM_ITERATIONS", 10); + int numWarmups = GETARG(6, "NUM_WARMUPS", 1000); + int numSleepUsec = GETARG(7, "SLEEP_USEC", 100); + int totalIterations = numWarmups + numIterations; + + int verbose = (getenv("VERBOSE" ) ? atoi(getenv("VERBOSE")) : 1); + int launchMode = (getenv("LAUNCH_MODE") ? atoi(getenv("LAUNCH_MODE")) : 1); + + if (numUpdateThreads == 0) numUpdateThreads = numUsedGpus; + + // Print off configuration and machine information + if (rank == 0) + { + printf("NUM_BLOCKS = %8d\n", numBlocks); + printf("BLOCKSIZE = %8d\n", blockSize); + printf("NUM_UPDATERS = %8d\n", numUpdateThreads); + printf("USE_NUMA = %8d\n", useNuma); + printf("NUM_ITERATIONS = %8d\n", numIterations); + printf("NUM_WARMUPS = %8d\n", numWarmups); + printf("SLEEP_USEC = %8d\n", numSleepUsec); + } + + char archName[100]; + std::vector uSecPerCycle(numUsedGpus); + for (int i = 0; i < numUsedGpus; i++) + { + hipDeviceProp_t prop; + HIP_CALL(hipGetDeviceProperties(&prop, i + rank)); + sscanf(prop.gcnArchName, "%[^:]", archName); + int wallClockMhz; + HIP_CALL(hipDeviceGetAttribute(&wallClockMhz, hipDeviceAttributeWallClockRate, i)); + uSecPerCycle[i] = 1000.0 / wallClockMhz; + if (verbose) printf("GPU %02d: %s: Closest NUMA: %d usecPerWallClockCycle %g\n", i + rank, archName, GetClosestNumaNode(i + rank), uSecPerCycle[i]); + } + + typedef typename std::ratio_multiply::type MicroSec; + //printf("std::chrono::steady_clock precision: %8.3f usec\n", static_cast(MicroSec::num)/MicroSec::den); + + // Allocate per-update-thread resources and start update threads + bool abortUpdateThreads = false; + std::vector cpuTimestamps(numUpdateThreads); + std::vector updateThreads; + for (int i = 0; i < numUpdateThreads; i++) + { + int numaId = GetClosestNumaNode(i + rank); + HIP_CALL(hipSetDevice(i + rank)); + if (useNuma) SetNumaNode(numaId); + + HostMalloc((void**)&cpuTimestamps[i], 256); // Allocate larger buffer to avoid multiple timestamps on same cacheline + + // Launch update thread + updateThreads.push_back(std::thread(UpdateCpuTime, useNuma, numaId, cpuTimestamps[i], &abortUpdateThreads)); + } + + // Allocate per-GPU resources + std::vector syncDataGpu(numUsedGpus); + std::vector abortFlags(numUsedGpus); + std::vector streams(numUsedGpus); + for (int i = 0; i < numUsedGpus; i++) + { + HIP_CALL(hipSetDevice(i + rank)); + if (useNuma) SetNumaNode(GetClosestNumaNode(i + rank)); + + HIP_CALL(hipMalloc((void**)&syncDataGpu[i], totalIterations * numBlocks * sizeof(SyncData))); + HostMalloc((void**)&abortFlags[i], 256); // Allocate larger buffer to avoid multiple abort flags on same cacheline + + HIP_CALL(hipStreamCreate(&streams[i])); + } + + // Allocate per-iteration resources + std::vector> hostStartTimes(numTotalGpus, std::vector(totalIterations)); + std::vector> hostReturnTimes(numTotalGpus, std::vector(totalIterations)); + std::vector> hostAbortTimes(numTotalGpus, std::vector(totalIterations)); + std::vector> hostStopTimes(numTotalGpus, std::vector(totalIterations)); + +#ifndef MPI_SUPPORT + #pragma omp parallel num_threads(numTotalGpus) +#endif + { +#ifdef MPI_SUPPORT + int deviceId = rank; + int localIdx = 0; +#else + int deviceId = omp_get_thread_num(); + int localIdx = deviceId; +#endif + HIP_CALL(hipSetDevice(deviceId)); + if (useNuma) SetNumaNode(GetClosestNumaNode(deviceId)); + + uint64_t* cpuTimestamp = cpuTimestamps[localIdx % numUpdateThreads]; + uint32_t* abortFlag = abortFlags[localIdx]; + + for (int iteration = 0; iteration < totalIterations; iteration++) + { + // Prepare for this iteration + // Clear abort flag + STORE(abortFlag, 0); + SyncData* syncData = syncDataGpu[localIdx] + (iteration * numBlocks); + + // Wait for all threads to arrive before launching all kernels +#ifdef MPI_SUPPORT + MPI_Barrier(MPI_COMM_WORLD); +#else + #pragma omp barrier +#endif + + // Launch kernel + uint64_t cpuStart = std::chrono::steady_clock::now().time_since_epoch().count(); + if (launchMode == 0) + { + SyncKernel<<>>(cpuTimestamp, abortFlag, syncData); + } + else + { + hipLaunchKernelGGL(SyncKernel, numBlocks, blockSize, 0, streams[localIdx], cpuTimestamp, abortFlag, syncData); + } + uint64_t cpuReturn = std::chrono::steady_clock::now().time_since_epoch().count(); + + // Busy wait performs more accurately than usleep / sleep_for + while (std::chrono::steady_clock::now().time_since_epoch().count() - cpuStart < numSleepUsec * 1000); + STORE(abortFlag, 1); + uint64_t cpuAbort = std::chrono::steady_clock::now().time_since_epoch().count(); + + // Wait for kernel to finish + HIP_CALL(hipStreamSynchronize(streams[localIdx])); + uint64_t cpuStop = std::chrono::steady_clock::now().time_since_epoch().count(); + + // Store values (after all timings to avoid false sharing) + hostStartTimes [deviceId][iteration] = cpuStart; + hostReturnTimes[deviceId][iteration] = cpuReturn; + hostAbortTimes [deviceId][iteration] = cpuAbort; + hostStopTimes [deviceId][iteration] = cpuStop; + +#ifdef MPI_SUPPORT + MPI_Barrier(MPI_COMM_WORLD); +#else + #pragma omp barrier +#endif + } + } + + // Stop all the update threads + STORE(&abortUpdateThreads, true); + for (auto& t : updateThreads) + t.join(); + + std::vector> syncDataCpu(numTotalGpus, std::vector(totalIterations * numBlocks)); + for (int i = 0; i < numUsedGpus; i++) + { + HIP_CALL(hipMemcpy(syncDataCpu[i+rank].data(), syncDataGpu[i], totalIterations * numBlocks * sizeof(SyncData), hipMemcpyDeviceToHost)); + } + + std::vector> singleMinDiff(numTotalGpus, std::vector(NUM_COLUMNS, std::numeric_limits::max())); + std::vector> singleSumDiff(numTotalGpus, std::vector(NUM_COLUMNS, 0)); + std::vector> singleMaxDiff(numTotalGpus, std::vector(NUM_COLUMNS, std::numeric_limits::min())); + std::vector multiMinDiff(NUM_COLUMNS, std::numeric_limits::max()); + std::vector multiSumDiff(NUM_COLUMNS, 0); + std::vector multiMaxDiff(NUM_COLUMNS, std::numeric_limits::min()); + std::vector timelineData; + char buff[1000]; + +#ifdef MPI_SUPPORT + // Collect results from every rank + if (rank == 0) + { + for (int deviceId = 1; deviceId < numTotalGpus; deviceId++) + { + MPI_Recv( hostStartTimes[deviceId].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, deviceId, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE); + MPI_Recv(hostReturnTimes[deviceId].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, deviceId, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE); + MPI_Recv( hostAbortTimes[deviceId].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, deviceId, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE); + MPI_Recv( hostStopTimes[deviceId].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, deviceId, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE); + MPI_Recv(syncDataCpu[deviceId].data(), totalIterations * numBlocks * sizeof(SyncData), MPI_BYTE, deviceId, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE); + } + } + else + { + MPI_Send( hostStartTimes[rank].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, 0, 0, MPI_COMM_WORLD); + MPI_Send(hostReturnTimes[rank].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, 0, 0, MPI_COMM_WORLD); + MPI_Send( hostAbortTimes[rank].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, 0, 0, MPI_COMM_WORLD); + MPI_Send( hostStopTimes[rank].data(), totalIterations * sizeof(uint64_t), MPI_BYTE, 0, 0, MPI_COMM_WORLD); + MPI_Send(syncDataCpu[rank].data(), totalIterations * numBlocks * sizeof(SyncData), MPI_BYTE, 0, 0, MPI_COMM_WORLD); + goto end; + } +#endif + + for (int iteration = 1; iteration <= numIterations; iteration++) + { + // Ignore warmup iterations + int iter = iteration + numWarmups - 1; + + // Figure out which timestamp is "earliest" to use as origin for this iteration + uint64_t origin = hostStartTimes[0][iter]; + for (int gpu = 1; gpu < numTotalGpus; gpu++) + origin = std::min(origin, hostStartTimes[gpu][iter]); + + if (verbose) + { + printf("Iteration %d: (All times in usec)\n", iteration); + printf("------------------------------------------------------------------------------------------------------------------------------------------\n"); + printf("| GPU | BLOCK | XCC | START(CPU) | RETURN(CPU)| START(GPU) | ABORT(CPU) | STOP (GPU) | STOP (CPU) | Kernel(CPU)| Kernel(GPU)| AbsDiff |\n"); + } + + std::vector multiMinTime(NUM_COLUMNS, std::numeric_limits::max()); + std::vector multiMaxTime(NUM_COLUMNS, std::numeric_limits::min()); + + for (int gpu = 0; gpu < numTotalGpus; gpu++) + { + std::vector times(NUM_COLUMNS); + times[HOST_START] = ( hostStartTimes[gpu][iter] - origin) / 1000.0; + times[HOST_RETURN] = (hostReturnTimes[gpu][iter] - origin) / 1000.0; + times[HOST_ABORT] = ( hostAbortTimes[gpu][iter] - origin) / 1000.0; + times[HOST_STOP] = ( hostStopTimes[gpu][iter] - origin) / 1000.0; + + TimelineData td; + sprintf(buff, "Iteration %d GPU %02d (CPU)", iteration, gpu); td.rowLabel = buff; + td.barLabel = "Launch ("; + sprintf(buff, "%.3f to %.3f", times[HOST_START], times[HOST_RETURN]); td.toolTip = buff; + td.startTime = times[HOST_START]; + td.stopTime = times[HOST_RETURN]; + timelineData.push_back(td); + + td.barLabel = "Pause"; + sprintf(buff, "%.3f to %.3f", times[HOST_RETURN], times[HOST_ABORT]); td.toolTip = buff; + td.startTime = times[HOST_RETURN]; + td.stopTime = times[HOST_ABORT]; + timelineData.push_back(td); + + td.barLabel = "Sync"; + sprintf(buff, "%.3f to %.3f", times[HOST_ABORT], times[HOST_STOP]); td.toolTip = buff; + td.startTime = times[HOST_ABORT]; + td.stopTime = times[HOST_STOP]; + timelineData.push_back(td); + + std::vector singleMinTime(NUM_COLUMNS, std::numeric_limits::max()); + std::vector singleMaxTime(NUM_COLUMNS, std::numeric_limits::min()); + for (int block = 0; block < numBlocks; block++) + { + int blockIdx = iter * numBlocks + block; + int xccId = syncDataCpu[gpu][blockIdx].xccId; + + times[DEV_START] = (syncDataCpu[gpu][blockIdx].cpuStart - origin) / 1000.0; + times[DEV_STOP] = (syncDataCpu[gpu][blockIdx].cpuStop - origin) / 1000.0; + times[KERNEL_CPUTIME] = times[DEV_STOP] - times[DEV_START]; + times[KERNEL_GPUTIME] = (syncDataCpu[gpu][blockIdx].gpuStop - syncDataCpu[gpu][blockIdx].gpuStart) * uSecPerCycle[gpu]; + times[KERNEL_TIMEDIFF] = fabs(times[KERNEL_CPUTIME] - times[KERNEL_GPUTIME]); + + for (int col = 0; col < NUM_COLUMNS; col++) + { + singleMinTime[col] = std::min(singleMinTime[col], times[col]); + singleMaxTime[col] = std::max(singleMaxTime[col], times[col]); + multiMinTime[col] = std::min( multiMinTime[col], times[col]); + multiMaxTime[col] = std::max( multiMaxTime[col], times[col]); + } + + if (verbose) + { + printf("| %3d | %3d | %3d |", gpu, block, xccId); + for (auto x : times) printf(" %10.3f |", x); + printf("\n"); + } + + sprintf(buff, "Iteration %d GPU %02d (GPU)", iteration, gpu); td.rowLabel = buff; + sprintf(buff, "Block %02d", block); td.barLabel = buff; + sprintf(buff, "%.3f to %.3f", times[DEV_START], times[DEV_STOP]); td.toolTip = buff; + td.startTime = times[DEV_START]; + td.stopTime = times[DEV_STOP]; + timelineData.push_back(td); + } + + for (int col = 0; col < NUM_COLUMNS; col++) + { + double const diff = singleMaxTime[col] - singleMinTime[col]; + singleMinDiff[gpu][col] = std::min(singleMinDiff[gpu][col], diff); + singleSumDiff[gpu][col] += diff; + singleMaxDiff[gpu][col] = std::max(singleMaxDiff[gpu][col], diff); + } + + if (verbose) + { + printf("| %3d | MAX ABS DIFF|", gpu); + for (int col = 0; col < NUM_COLUMNS; col++) + printCol[col] ? printf(" %10.3f |", singleMaxTime[col] - singleMinTime[col]) : printf(" |"); + printf("\n"); + } + } + for (int col = 0; col < NUM_COLUMNS; col++) + { + double const diff = multiMaxTime[col] - multiMinTime[col]; + multiMinDiff[col] = std::min(multiMinDiff[col], diff); + multiSumDiff[col] += diff; + multiMaxDiff[col] = std::max(multiMaxDiff[col], diff); + } + + if (verbose) + { + printf("------------------------------------------------------------------------------------------------------------------------------------------\n"); + printf("| ALL | MIN |"); for (auto x : multiMinTime) printf(" %10.3f |", x); printf("\n"); + printf("| ALL | MAX |"); for (auto x : multiMaxTime) printf(" %10.3f |", x); printf("\n"); + printf("| ALL | DIFF |"); for (int col = 0; col < NUM_COLUMNS; col++) printf(" %10.3f |", multiMaxTime[col] - multiMinTime[col]); printf("\n"); + } + } + + printf("==========================================================================================================================================\n"); + printf("| SUMMARY (All iter)| START(CPU) | RETURN(CPU)| START(GPU) | ABORT(CPU) | STOP (GPU) | STOP (CPU) | Kernel(CPU)| Kernel(GPU)| AbsDiff |\n"); + printf("==========================================================================================================================================\n"); + for (int gpu = 0; gpu < numTotalGpus; gpu++) + { + printf("| GPU %02d DIFF MIN |", gpu); + for (int col = 0; col < NUM_COLUMNS; col++) + printCol[col] ? printf(" %10.3f |", singleMinDiff[gpu][col]) : printf(" |"); + printf("\n"); + } + for (int gpu = 0; gpu < numTotalGpus; gpu++) + { + printf("| GPU %02d DIFF AVG |", gpu); + for (int col = 0; col < NUM_COLUMNS; col++) + printCol[col] ? printf(" %10.3f |", singleSumDiff[gpu][col] / numIterations) : printf(" |"); + printf("\n"); + } + for (int gpu = 0; gpu < numTotalGpus; gpu++) + { + printf("| GPU %02d DIFF MAX |", gpu); + for (int col = 0; col < NUM_COLUMNS; col++) + printCol[col] ? printf(" %10.3f |", singleMaxDiff[gpu][col]) : printf(" |"); + printf("\n"); + } + printf("==========================================================================================================================================\n"); + printf("| ALL GPUs DIFF MIN |"); for (auto x : multiMinDiff) printf(" %10.3f |", x); printf("\n"); + printf("| ALL GPUs DIFF AVG |"); for (auto x : multiSumDiff) printf(" %10.3f |", x / numIterations); printf("\n"); + printf("| ALL GPUs DIFF MAX |"); for (auto x : multiMaxDiff) printf(" %10.3f |", x); printf("\n"); + + sprintf(buff, "timeline_%dx%s_%dx%dblockSize_%dCUTs_Numa%d_Sleep%d.html", numTotalGpus, archName, numBlocks, blockSize, numUpdateThreads, useNuma, numSleepUsec); + printf("Timeline exported to %s\n", buff); + ExportToTimeLine(buff, "Device", "Call", timelineData); + +#ifdef MPI_SUPPORT +end: + MPI_Barrier(MPI_COMM_WORLD); + MPI_Finalize(); +#endif + return 0; +} diff --git a/tools/JitterBench/Makefile b/tools/JitterBench/Makefile new file mode 100644 index 0000000000..4a42f5bef8 --- /dev/null +++ b/tools/JitterBench/Makefile @@ -0,0 +1,25 @@ +ROCM_PATH ?= /opt/rocm +CUDA_PATH ?= /usr/local/cuda +HIPCC = $(ROCM_PATH)/bin/hipcc +NVCC = $(CUDA_PATH)/bin/nvcc + +CCFLAGS = -O3 -lhsa-runtime64 -fopenmp -lnuma +NVFLAGS = -O3 -x cu -lnuma -Xcompiler -fopenmp -gencode=arch=compute_90,code=sm_90 + +ifneq ("$(MPI_DIR)", "") +MPIFLAGS = -DMPI_SUPPORT -I$(MPI_DIR)/include -L$(MPI_DIR)/lib -lmpi +else +MPIFLAGS = +endif + +all: JitterBench + +JitterBench: JitterBench.cpp Common.hpp Timeline.hpp +ifeq ("$(shell test -e $(NVCC) && echo found)", "found") + $(NVCC) $(NVFLAGS) $(MPIFLAGS) $< -o $@ +else + $(HIPCC) $(CCFLAGS) $(MPIFLAGS) $< -o $@ +endif + +clean: + rm -f ./JitterBench diff --git a/tools/JitterBench/Timeline.hpp b/tools/JitterBench/Timeline.hpp new file mode 100644 index 0000000000..503ced6a8a --- /dev/null +++ b/tools/JitterBench/Timeline.hpp @@ -0,0 +1,51 @@ +#pragma once +#include +#include +#include + +struct TimelineData +{ + std::string rowLabel; + std::string barLabel; + std::string toolTip; + uint64_t startTime; + uint64_t stopTime; +}; + +void ExportToTimeLine(std::string outputFilename, + std::string rowLabelName, + std::string barLabelName, + std::vector const& timelineData) +{ + FILE *fp = fopen(outputFilename.c_str(), "w"); + + fprintf(fp, "\n"); + fprintf(fp, "\n"); + fprintf(fp, "
\n"); + fclose(fp); +} diff --git a/tools/JitterBench/runSweep.sh b/tools/JitterBench/runSweep.sh new file mode 100755 index 0000000000..96f968022d --- /dev/null +++ b/tools/JitterBench/runSweep.sh @@ -0,0 +1,12 @@ +#!/bin/bash + +for numBlocks in 1 2 4 8 16 32; do + for blockSize in 64 128 256; do + for numTimers in 0 1; do + for useNuma in 0 1; do + echo "numBlocks=$numBlocks blockSize=$blockSize numTimers=$numTimers useNuma=$useNuma"; + ./LaunchBench $numBlocks $blockSize $numTimers $useNuma &> output.$numBlocks.$blockSize.$numTimers.$useNuma.txt + done; + done; + done; +done; diff --git a/tools/LaunchBench/LaunchBench.cpp b/tools/LaunchBench/LaunchBench.cpp deleted file mode 100644 index 15ed68c3ce..0000000000 --- a/tools/LaunchBench/LaunchBench.cpp +++ /dev/null @@ -1,230 +0,0 @@ -/* -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 -#include -#include -#include "Common.hpp" -#include -#include - -struct SyncData -{ - uint64_t cpuStart; - uint64_t cpuStop; - int32_t xccId; -}; - -__global__ void SyncKernel(volatile uint64_t* cpuTime, - volatile uint32_t* abortFlag, - SyncData* syncData) -{ - // Collect timestamp upon kernel entry - uint64_t cpuStart = *cpuTime; - - // Wait for abort flag to be modified - while (*abortFlag == 0); - - // Collect timestamps after abort flag - uint64_t cpuStop = *cpuTime; - - // Save timestamps - syncData[blockIdx.x].cpuStart = cpuStart; - syncData[blockIdx.x].cpuStop = cpuStop; - GetXccId(syncData[blockIdx.x].xccId); -} - -void UpdateCpuTime(volatile uint64_t* cpuTimestamp, volatile bool& abortThread) -{ - while (!abortThread) - { - *cpuTimestamp = std::chrono::steady_clock::now().time_since_epoch().count(); - } -} - - -int main(int argc, char **argv) -{ - int numBlocks = (argc > 1 ? atoi(argv[1]) : 4); - int numIterations = (argc > 2 ? atoi(argv[2]) : 1); - int numWarmups = (argc > 3 ? atoi(argv[3]) : 100); - int numSleepUsec = (argc > 4 ? atoi(argv[4]) : 20); - - int totalIterations = numWarmups + numIterations; - - int numGpus; - HIP_CALL(hipGetDeviceCount(&numGpus)); - printf("Running %d GPUs with %d block(s) each, %d timed iterations, %d untimed warmup iterations, sleeping for %d usec\n", - numGpus, numBlocks, numIterations, numWarmups, numSleepUsec); - for (int i = 0; i < numGpus; i++) - { - hipDeviceProp_t prop; - HIP_CALL(hipGetDeviceProperties(&prop, i)); - printf("GPU %02d: %s\n", i, prop.gcnArchName); - } - - typedef typename std::ratio_multiply::type MicroSec; - printf("std::chrono::steady_clock precision: %8.3f usec\n", - static_cast(MicroSec::num)/MicroSec::den); - - - // Allocate pinned host memory for CPU timestamp / abort flag - volatile uint64_t* cpuTimestamp; - volatile uint32_t* abortFlag; - HIP_CALL(hipHostMalloc((void**)&cpuTimestamp, sizeof(uint64_t))); - HIP_CALL(hipHostMalloc((void**)&abortFlag, sizeof(uint32_t))); - - // Allocate device memory for collecting timestamps - std::vector syncDataList(numGpus); - std::vectorstreams(numGpus); - for (int i = 0; i < numGpus; i++) - { - HIP_CALL(hipSetDevice(i)); - HIP_CALL(hipMalloc((void**)&syncDataList[i], numIterations * numBlocks * sizeof(SyncData))); - HIP_CALL(hipStreamCreate(&streams[i])); - } - - // Start update thread - // NOTE: NPKit usually runs 1 GPU per process which means 1 update thread per GPU - // However in this case, only a single CPU update thread is used - volatile bool abortThread = false; - std::thread updateThread(UpdateCpuTime, cpuTimestamp, std::ref(abortThread)); - - // Launch one thread per GPU - std::vector cpuAbortTime(totalIterations); - std::vector> cpuStartList(numGpus, std::vector(totalIterations, 0)); - std::vector> cpuReturnList(numGpus, std::vector(totalIterations, 0)); - std::vector> cpuStopList(numGpus, std::vector(totalIterations, 0)); - - uint64_t cpuAbort; - #pragma omp parallel num_threads(numGpus) - { - int deviceId = omp_get_thread_num(); - HIP_CALL(hipSetDevice(deviceId)); - - for (int iteration = 0; iteration < totalIterations; iteration++) - { - // Single thread resets abort flag - #pragma omp single - *abortFlag = 0; - - // Prepare for this iteration - SyncData* syncData = syncDataList[deviceId] + (iteration * numBlocks); - - // Wait for all threads to arrive before launching all kernels - #pragma omp barrier - uint64_t cpuStart = std::chrono::steady_clock::now().time_since_epoch().count(); - SyncKernel<<>>(cpuTimestamp, abortFlag, syncData); - uint64_t cpuReturn = std::chrono::steady_clock::now().time_since_epoch().count(); - - // Busy wait performs more accurately than usleep / sleep_for - if (deviceId == 0) - { - while (std::chrono::steady_clock::now().time_since_epoch().count() - cpuStart < numSleepUsec * 1000); - *abortFlag = 1; - cpuAbort = std::chrono::steady_clock::now().time_since_epoch().count(); - } - - // Wait for kernels to finish - HIP_CALL(hipStreamSynchronize(streams[deviceId])); - uint64_t cpuStop = std::chrono::steady_clock::now().time_since_epoch().count(); - - // Store values (after all timings to avoid false sharing) - cpuStartList[deviceId][iteration] = cpuStart; - cpuReturnList[deviceId][iteration] = cpuReturn; - #pragma omp single - cpuAbortTime[iteration] = cpuAbort; - cpuStopList[deviceId][iteration] = cpuStop; - - #pragma omp barrier - } - } - - abortThread = true; - updateThread.join(); - - for (int iteration = 1; iteration <= numIterations; iteration++) - { - // Ignore warmup iterations - int iter = iteration + numWarmups - 1; - printf("---------------------------------------------------------------------------------------------------\n"); - printf("Iteration %d: (All times in usec)\n", iteration); - - uint64_t origin = cpuStartList[0][iter]; - for (int gpu = 0; gpu < numGpus; gpu++) - { - for (int block = 0; block < numBlocks; block++) - { - origin = std::min(origin, cpuStartList[gpu][iter]); - origin = std::min(origin, syncDataList[gpu][iter * numBlocks + block].cpuStart); - origin = std::min(origin, cpuAbortTime[iter]); - origin = std::min(origin, syncDataList[gpu][iter * numBlocks + block].cpuStop); - origin = std::min(origin, cpuStopList[gpu][iter]); - } - } - - printf("| GPU | BLOCK | XCC | START(CPU) | RETURN(CPU)| START(GPU) | ABORT(CPU) | STOP (GPU) | STOP (CPU) |\n"); - - double minCpuStart, minGpuStart, minCpuReturn, minCpuAbort, minGpuStop, minCpuStop; - double maxCpuStart, maxGpuStart, maxCpuReturn, maxCpuAbort, maxGpuStop, maxCpuStop; - - for (int gpu = 0; gpu < numGpus; gpu++) - { - for (int block = 0; block < numBlocks; block++) - { - int xccId = syncDataList[gpu][iter * numBlocks + block].xccId; - double cpuStart = (cpuStartList[gpu][iter] - origin) / 1000.0; - double gpuStart = (syncDataList[gpu][iter * numBlocks + block].cpuStart - origin) / 1000.0; - double cpuReturn = (cpuReturnList[gpu][iter] - origin) / 1000.0; - double cpuAbort = (cpuAbortTime[iter] - origin) / 1000.0; - double gpuStop = (syncDataList[gpu][iter * numBlocks + block].cpuStop - origin) / 1000.0; - double cpuStop = (cpuStopList[gpu][iter] - origin) / 1000.0; - - minCpuStart = ((gpu == 0 && block == 0) || (minCpuStart > cpuStart)) ? cpuStart : minCpuStart; - maxCpuStart = ((gpu == 0 && block == 0) || (maxCpuStart < cpuStart)) ? cpuStart : maxCpuStart; - minGpuStart = ((gpu == 0 && block == 0) || (minGpuStart > gpuStart)) ? gpuStart : minGpuStart; - maxGpuStart = ((gpu == 0 && block == 0) || (maxGpuStart < gpuStart)) ? gpuStart : maxGpuStart; - minCpuReturn = ((gpu == 0 && block == 0) || (minCpuReturn > cpuReturn)) ? cpuReturn : minCpuReturn; - maxCpuReturn = ((gpu == 0 && block == 0) || (maxCpuReturn < cpuReturn)) ? cpuReturn : maxCpuReturn; - minCpuAbort = ((gpu == 0 && block == 0) || (minCpuAbort > cpuAbort)) ? cpuAbort : minCpuAbort; - maxCpuAbort = ((gpu == 0 && block == 0) || (maxCpuAbort < cpuAbort)) ? cpuAbort : maxCpuAbort; - minGpuStop = ((gpu == 0 && block == 0) || (minGpuStop > gpuStop)) ? gpuStop : minGpuStop; - maxGpuStop = ((gpu == 0 && block == 0) || (maxGpuStop < gpuStop)) ? gpuStop : maxGpuStop; - minCpuStop = ((gpu == 0 && block == 0) || (minCpuStop > gpuStop)) ? gpuStop : minCpuStop; - maxCpuStop = ((gpu == 0 && block == 0) || (maxCpuStop < gpuStop)) ? gpuStop : maxCpuStop; - - printf("| %3d | %3d | %3d | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n", - gpu, block, xccId, cpuStart, cpuReturn, gpuStart, cpuAbort, gpuStop, cpuStop); - } - } - printf("---------------------------------------------------------------------------------------------------\n"); - printf("| MIN | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n", - minCpuStart, minCpuReturn, minGpuStart, minCpuAbort, minGpuStop, minCpuStop); - printf("| MAX | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n", - maxCpuStart, maxCpuReturn, maxGpuStart, maxCpuAbort, maxGpuStop, maxCpuStop); - printf("| DIFF | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n", - maxCpuStart - minCpuStart, maxCpuReturn - minCpuReturn, maxGpuStart - minGpuStart, - maxCpuAbort - minCpuAbort, maxGpuStop - minGpuStop, maxCpuStop - minCpuStop); - } - - return 0; -} diff --git a/tools/LaunchBench/Makefile b/tools/LaunchBench/Makefile deleted file mode 100644 index 36bf5fe0ed..0000000000 --- a/tools/LaunchBench/Makefile +++ /dev/null @@ -1,11 +0,0 @@ -ROCM_DIR = /opt/rocm -HIPCC = $(ROCM_DIR)/bin/hipcc -CCFLAGS = -lhsa-runtime64 -fopenmp - -all: LaunchBench - -LaunchBench: LaunchBench.cpp Common.hpp - $(HIPCC) $(CCFLAGS) $< -o $@ - -clean: - rm -f ./LaunchBench