SWDEV-276853 - Add test for mem perf
Add memory filling performance test for large bar device memory, hip managed memory, conherent/non-coherent host memory and fine grained device memory Change-Id: Id3f9bcad45d643e493daf9d5f47b3a012a427178
Bu işleme şunda yer alıyor:
Tao Sang
@@ -0,0 +1,461 @@
|
||||
/*
|
||||
Copyright (c) 2015-present 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 EXCLUDE_HIP_PLATFORM nvidia
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#include <iostream>
|
||||
#include <chrono>
|
||||
#include <sys/time.h>
|
||||
|
||||
#define SIMPLY_ASSIGN 0
|
||||
#define USE_HIPTEST_SETNUMBLOCKS 0
|
||||
|
||||
using namespace std;
|
||||
|
||||
template<class T>
|
||||
__global__ void vec_fill(T *x, T coef, int N) {
|
||||
const int istart = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
const int ishift = blockDim.x * gridDim.x;
|
||||
for (int i = istart; i < N; i += ishift) {
|
||||
#if SIMPLY_ASSIGN
|
||||
x[i] = coef;
|
||||
#else
|
||||
x[i] = coef * i;
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
template<class T>
|
||||
__global__ void daxpy(T *__restrict__ x, T *__restrict__ y,
|
||||
const T coef, int Niter, int N) {
|
||||
const int istart = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
const int ishift = blockDim.x * gridDim.x;
|
||||
for (int iter = 0; iter < Niter; ++iter) {
|
||||
T iv = coef * iter;
|
||||
for (int i = istart; i < N; i += ishift)
|
||||
y[i] = iv * x[i] + y[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<class T>
|
||||
class hipPerfMemFill {
|
||||
private:
|
||||
static constexpr int NUM_START = 27;
|
||||
static constexpr int NUM_SIZE = 5;
|
||||
static constexpr int NUM_ITER = 10;
|
||||
size_t totalSizes_[NUM_SIZE];
|
||||
hipDeviceProp_t props_;
|
||||
const T coef_ = getCoefficient(3.14159);
|
||||
const unsigned int blocksPerCU_;
|
||||
const unsigned int threadsPerBlock_;
|
||||
|
||||
public:
|
||||
hipPerfMemFill(unsigned int blocksPerCU, unsigned int threadsPerBlock) :
|
||||
blocksPerCU_(blocksPerCU), threadsPerBlock_(threadsPerBlock) {
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
totalSizes_[i] = 1ull << (i + NUM_START); // 128M, 256M, 512M, 1024M, 2048M
|
||||
}
|
||||
}
|
||||
|
||||
~hipPerfMemFill() {
|
||||
}
|
||||
|
||||
bool supportLargeBar() {
|
||||
return props_.isLargeBar != 0;
|
||||
}
|
||||
|
||||
bool supportManagedMemory() {
|
||||
return props_.managedMemory != 0;
|
||||
}
|
||||
|
||||
const T getCoefficient(double val) {
|
||||
return static_cast<T>(val);
|
||||
}
|
||||
|
||||
void setHostBuffer(T *A, T val, size_t size) {
|
||||
size_t len = size / sizeof(T);
|
||||
for (int i = 0; i < len; i++) {
|
||||
A[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
void open(int deviceId) {
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
cout << "Info: didn't find any GPU! skipping the test!\n";
|
||||
passed();
|
||||
} else if (deviceId >= nGpu) {
|
||||
failed("Info: wrong GPU Id %d\n", deviceId);
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(deviceId));
|
||||
HIPCHECK(hipGetDeviceProperties(&props_, deviceId));
|
||||
std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x"
|
||||
<< props_.pciBusID << " " << props_.name << " with "
|
||||
<< props_.multiProcessorCount << " CUs, large bar: "
|
||||
<< supportLargeBar() << ", managed memory: " << supportManagedMemory()
|
||||
<< ", DeviceMallocFinegrained: " << supportDeviceMallocFinegrained()
|
||||
<< std::endl;
|
||||
}
|
||||
|
||||
void hostFill(size_t size, T *data, T coef, double &sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
auto start = chrono::steady_clock::now();
|
||||
for (int i = 0; i < num; ++i) {
|
||||
#if SIMPLY_ASSIGN
|
||||
data[i] = coef;
|
||||
#else
|
||||
data[i] = coef * i;
|
||||
#endif
|
||||
}
|
||||
auto end = chrono::steady_clock::now();
|
||||
chrono::duration<double> diff = end - start; // in second
|
||||
sec = diff.count();
|
||||
}
|
||||
|
||||
void kernelFill(size_t size, T *data, T coef, double &sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
unsigned blocks = setNumBlocks(num);
|
||||
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, 0, data, 0, num); // kernel will be loaded first time
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
auto start = chrono::steady_clock::now();
|
||||
|
||||
for (int iter = 0; iter < NUM_ITER; ++iter) {
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, 0, data, coef, num);
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
auto end = chrono::steady_clock::now();
|
||||
chrono::duration<double> diff = end - start; // in second
|
||||
sec = diff.count() / NUM_ITER; // in second
|
||||
}
|
||||
|
||||
void verify(size_t size, T *data, T coef, double &sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
auto start = chrono::steady_clock::now();
|
||||
for (int i = 0; i < num; ++i) {
|
||||
#if SIMPLY_ASSIGN
|
||||
if(data[i] != coef) {
|
||||
cout << "verify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef << endl;
|
||||
failed("failed\n");
|
||||
}
|
||||
#else
|
||||
if(data[i] != coef * i) {
|
||||
cout << "verify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef * i << endl;
|
||||
failed("failed\n");
|
||||
}
|
||||
#endif
|
||||
}
|
||||
auto end = chrono::steady_clock::now();
|
||||
chrono::duration<double> diff = end - start; // in second
|
||||
sec = diff.count();
|
||||
}
|
||||
|
||||
bool testLargeBarHostFill(size_t size) {
|
||||
if (!supportLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
|
||||
T *A;
|
||||
HIPCHECK(hipMalloc(&A, size));
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, sec); // Cpu can access device mem in LB
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
cout << "Largebar: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl;
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testLargeBar() {
|
||||
if (!supportLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
cout << "Test large bar device memory host filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testLargeBarHostFill(totalSizes_[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testManagedMemoryHostFill(size_t size) {
|
||||
if (!supportManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
|
||||
T *A;
|
||||
HIPCHECK(hipMallocManaged(&A, size));
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, sec); // Cpu can access HMM mem
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
cout << "Managed: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testManagedMemoryKernelFill(size_t size) {
|
||||
if (!supportManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
|
||||
T *A;
|
||||
HIPCHECK(hipMallocManaged(&A, size));
|
||||
double sec = 0, sec_v = 0;
|
||||
kernelFill(size, A, coef_, sec);
|
||||
verify(size, A, coef_, sec_v); // Managed memory can be verified by host
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
cout << "Managed: kernel fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", verify: cost "
|
||||
<< setw(10) << sec_v << " s in bandwidth " << setw(10) << GBytes / sec_v << " [GB/s]"
|
||||
<< endl;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testManagedMemory() {
|
||||
if (!supportManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
cout << "Test managed memory host filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testManagedMemoryHostFill(totalSizes_[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
cout << "Test managed memory kernel filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testManagedMemoryKernelFill(totalSizes_[i])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testHostMemoryHostFill(size_t size, unsigned int flags) {
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
T *A;
|
||||
HIPCHECK(hipHostMalloc(&A, size, flags));
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, sec);
|
||||
HIPCHECK(hipHostFree(A));
|
||||
|
||||
cout << "Host: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testHostMemoryKernelFill(size_t size, unsigned int flags) {
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
|
||||
T *A;
|
||||
HIPCHECK(hipHostMalloc((void** ) &A, size, flags));
|
||||
double sec = 0, sec_v = 0;
|
||||
kernelFill(size, A, coef_, sec);
|
||||
verify(size, A, coef_, sec_v);
|
||||
HIPCHECK(hipHostFree(A));
|
||||
|
||||
cout << "Host: kernel fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", verify cost "
|
||||
<< setw(10) << sec_v << " s in bandwidth " << setw(10) << GBytes / sec_v << " [GB/s]"
|
||||
<< endl;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testHostMemory() {
|
||||
cout << "Test coherent host memory host filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
cout << "Test non-coherent host memory host filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocNonCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
cout << "Test coherent host memory kernel filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
cout << "Test non-coherent host memory kernel filling" << endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocNonCoherent)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/* This fuction should be via device attribute query*/
|
||||
bool supportDeviceMallocFinegrained() {
|
||||
T *A = nullptr;
|
||||
hipExtMallocWithFlags((void **)&A, sizeof(T), hipDeviceMallocFinegrained);
|
||||
if (!A) {
|
||||
return false;
|
||||
}
|
||||
HIPCHECK(hipFree(A));
|
||||
return true;
|
||||
}
|
||||
|
||||
unsigned int setNumBlocks(size_t size) {
|
||||
size_t num = size/sizeof(T);
|
||||
|
||||
#if USE_HIPTEST_SETNUMBLOCKS
|
||||
return HipTest::setNumBlocks(blocksPerCU_, threadsPerBlock_,
|
||||
num);
|
||||
#else
|
||||
return (num + threadsPerBlock_ - 1) / threadsPerBlock_;
|
||||
#endif
|
||||
}
|
||||
bool testExtDeviceMemoryHostFill(size_t size, unsigned int flags) {
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
|
||||
T *A = nullptr;
|
||||
HIPCHECK(hipExtMallocWithFlags((void **)&A, size, flags));
|
||||
if (!A) {
|
||||
return false;
|
||||
}
|
||||
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, sec); // Cpu can access this mem
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
cout << "ExtDevice: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl;
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testExtDeviceMemoryKernelFill(size_t size, unsigned int flags) {
|
||||
double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);
|
||||
|
||||
T *A = nullptr;
|
||||
HIPCHECK(hipExtMallocWithFlags((void **)&A, size, flags));
|
||||
if (!A) {
|
||||
return false;
|
||||
}
|
||||
|
||||
double sec = 0, sec_v = 0;
|
||||
kernelFill(size, A, coef_, sec);
|
||||
verify(size, A, coef_, sec_v); // Fine grained device memory can be verified by host
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
cout << "ExtDevice: kernel fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", verify cost "
|
||||
<< setw(10) << sec_v << " s in bandwidth " << setw(10) << GBytes / sec_v << " [GB/s]"
|
||||
<< endl;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool testExtDeviceMemory() {
|
||||
cout << "Test fine grained device memory host filling"
|
||||
<< endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testExtDeviceMemoryHostFill(totalSizes_[i],
|
||||
hipDeviceMallocFinegrained)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
cout << "Test fine grained device memory kernel filling"
|
||||
<< endl;
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
if (!testExtDeviceMemoryKernelFill(totalSizes_[i],
|
||||
hipDeviceMallocFinegrained)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool run() {
|
||||
if (supportLargeBar()) {
|
||||
if (!testLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (supportManagedMemory()) {
|
||||
if (!testManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (!testHostMemory()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (supportDeviceMallocFinegrained()) {
|
||||
if (!testExtDeviceMemory()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
};
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true); // For ::p_gpuDevice, ::blocksPerCU, ::threadsPerBlock
|
||||
cout << "Test int" << endl;
|
||||
hipPerfMemFill<int> hipPerfMemFillInt(::blocksPerCU, ::threadsPerBlock);
|
||||
hipPerfMemFillInt.open(::p_gpuDevice);
|
||||
HIPASSERT(hipPerfMemFillInt.run());
|
||||
|
||||
cout << "Test double" << endl;
|
||||
hipPerfMemFill<double> hipPerfMemFillDouble(::blocksPerCU, ::threadsPerBlock);
|
||||
hipPerfMemFillDouble.open(::p_gpuDevice);
|
||||
HIPASSERT(hipPerfMemFillDouble.run());
|
||||
|
||||
passed();
|
||||
}
|
||||
Yeni konuda referans
Bir kullanıcı engelle