SWDEV-546485 Port and clean up for all tests in catch/perftests/memory folder. (#558)
* SWDEV-546485 Port and clean up for hipPerfBufferCopyRectSpeed * SWDEV-546485 Port and clean up for hipPerfDevMemReadSpeed * SWDEV-546485 Port and clean up for hipPerfDevMemWriteSpeed * SWDEV-546485 Port and clean up for hipPerfHostNumaAlloc * SWDEV-546485 Port and clean up for hipPerfMemcpy * SWDEV-546485 Port and clean up for hipPerfMemMallocCpyFree * SWDEV-546485 Port and clean up for hipPerfMemset * SWDEV-546485 Port and clean up for hipPerfSampleRate * SWDEV-546485 Port and clean up for hipPerfSharedMemReadSpeed * SWDEV-546485 Ported and fixed up segfault for hipPerfMemFill * SWDEV-545485 Returning to unedited stage
Tá an tiomantas seo le fáil i:
Naeisseh, Hadi
tiomanta ag
GitHub
GitHub
tuismitheoir
9fdc9a98b7
tiomantas
04469c0cde
@@ -1,300 +0,0 @@
|
||||
/*
|
||||
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 ../../src/timer.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <stdio.h>
|
||||
#include <assert.h>
|
||||
#include <string.h>
|
||||
#include <complex>
|
||||
|
||||
#include "timer.h"
|
||||
#include "test_common.h"
|
||||
|
||||
// Quiet pesky warnings
|
||||
#ifdef WIN_OS
|
||||
#define SNPRINTF sprintf_s
|
||||
#else
|
||||
#define SNPRINTF snprintf
|
||||
#endif
|
||||
|
||||
#define NUM_SIZES 8
|
||||
//4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10
|
||||
static const unsigned int Sizes[NUM_SIZES] = {4096, 8192, 65536, 262144, 1048576, 4194304, 16777216, 16777216+10};
|
||||
|
||||
static const unsigned int Iterations[2] = {1, 1000};
|
||||
|
||||
#define BUF_TYPES 4
|
||||
// 16 ways to combine 4 different buffer types
|
||||
#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES)
|
||||
|
||||
#define CHECK_RESULT(test, msg) \
|
||||
if ((test)) \
|
||||
{ \
|
||||
printf("\n%s\n", msg); \
|
||||
abort(); \
|
||||
}
|
||||
|
||||
void setData(void *ptr, unsigned int size, char value)
|
||||
{
|
||||
char *ptr2 = (char *)ptr;
|
||||
for (unsigned int i = 0; i < size ; i++)
|
||||
{
|
||||
ptr2[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
void checkData(void *ptr, unsigned int size, char value)
|
||||
{
|
||||
char *ptr2 = (char *)ptr;
|
||||
for (unsigned int i = 0; i < size; i++)
|
||||
{
|
||||
if (ptr2[i] != value)
|
||||
{
|
||||
printf("Data validation failed at %d! Got 0x%08x\n", i, ptr2[i]);
|
||||
printf("Expected 0x%08x\n", value);
|
||||
CHECK_RESULT(true, "Data validation failed!");
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
hipError_t err = hipSuccess;
|
||||
hipDeviceProp_t props = {0};
|
||||
hipGetDeviceProperties(&props, p_gpuDevice);
|
||||
CHECK_RESULT(err != hipSuccess, "hipGetDeviceProperties failed" );
|
||||
printf("Set device to %d : %s\n", p_gpuDevice, props.name);
|
||||
printf("Legend: unp - unpinned(malloc), hM - hipMalloc(device)\n");
|
||||
printf(" hHR - hipHostRegister(pinned), hHM - hipHostMalloc(prePinned)\n");
|
||||
err = hipSetDevice(p_gpuDevice);
|
||||
CHECK_RESULT(err != hipSuccess, "hipSetDevice failed" );
|
||||
|
||||
unsigned int bufSize_;
|
||||
bool hostMalloc[2] = {false};
|
||||
bool hostRegister[2] = {false};
|
||||
bool unpinnedMalloc[2] = {false};
|
||||
unsigned int numIter;
|
||||
void *memptr[2] = {NULL};
|
||||
void *alignedmemptr[2] = {NULL};
|
||||
void* srcBuffer = NULL;
|
||||
void* dstBuffer = NULL;
|
||||
|
||||
int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests;
|
||||
int test = (p_tests == -1) ? 0 : p_tests;
|
||||
|
||||
for(;test <= numTests; test++)
|
||||
{
|
||||
unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES;
|
||||
unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES;
|
||||
bufSize_ = Sizes[test % NUM_SIZES];
|
||||
hostMalloc[0] = hostMalloc[1] = false;
|
||||
hostRegister[0] = hostRegister[1] = false;
|
||||
unpinnedMalloc[0] = unpinnedMalloc[1] = false;
|
||||
srcBuffer = dstBuffer = 0;
|
||||
memptr[0] = memptr[1] = NULL;
|
||||
alignedmemptr[0] = alignedmemptr[1] = NULL;
|
||||
|
||||
size_t width = static_cast<size_t>(sqrt(static_cast<float>(bufSize_)));
|
||||
|
||||
if (srcTest == 3)
|
||||
{
|
||||
hostRegister[0] = true;
|
||||
}
|
||||
else if (srcTest == 2)
|
||||
{
|
||||
hostMalloc[0] = true;
|
||||
}
|
||||
else if (srcTest == 1)
|
||||
{
|
||||
unpinnedMalloc[0] = true;
|
||||
}
|
||||
|
||||
if (dstTest == 1)
|
||||
{
|
||||
unpinnedMalloc[1] = true;
|
||||
}
|
||||
else if (dstTest == 2)
|
||||
{
|
||||
hostMalloc[1] = true;
|
||||
}
|
||||
else if (dstTest == 3)
|
||||
{
|
||||
hostRegister[1] = true;
|
||||
}
|
||||
|
||||
numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)];
|
||||
|
||||
if (hostMalloc[0])
|
||||
{
|
||||
err = hipHostMalloc((void**)&srcBuffer, bufSize_, 0);
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
CHECK_RESULT(err != hipSuccess, "hipHostMalloc failed");
|
||||
}
|
||||
else if (hostRegister[0])
|
||||
{
|
||||
memptr[0] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[0] = (void*)(((size_t)memptr[0] + 4095) & ~4095);
|
||||
srcBuffer = alignedmemptr[0];
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
err = hipHostRegister(srcBuffer, bufSize_, 0);
|
||||
CHECK_RESULT(err != hipSuccess, "hipHostRegister failed");
|
||||
}
|
||||
else if (unpinnedMalloc[0])
|
||||
{
|
||||
memptr[0] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[0] = (void*)(((size_t)memptr[0] + 4095) & ~4095);
|
||||
srcBuffer = alignedmemptr[0];
|
||||
setData(srcBuffer, bufSize_, 0xd0);
|
||||
}
|
||||
else
|
||||
{
|
||||
err = hipMalloc(&srcBuffer, bufSize_);
|
||||
CHECK_RESULT(err != hipSuccess, "hipMalloc failed");
|
||||
err = hipMemset(srcBuffer, 0xd0, bufSize_);
|
||||
CHECK_RESULT(err != hipSuccess, "hipMemset failed");
|
||||
}
|
||||
|
||||
if (hostMalloc[1])
|
||||
{
|
||||
err = hipHostMalloc((void**)&dstBuffer, bufSize_, 0);
|
||||
CHECK_RESULT(err != hipSuccess, "hipHostMalloc failed");
|
||||
}
|
||||
else if (hostRegister[1])
|
||||
{
|
||||
memptr[1] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[1] = (void*)(((size_t)memptr[1] + 4095) & ~4095);
|
||||
dstBuffer = alignedmemptr[1];
|
||||
err = hipHostRegister(dstBuffer, bufSize_, 0);
|
||||
CHECK_RESULT(err != hipSuccess, "hipHostRegister failed");
|
||||
}
|
||||
else if (unpinnedMalloc[1])
|
||||
{
|
||||
memptr[1] = malloc(bufSize_ + 4096);
|
||||
alignedmemptr[1] = (void*)(((size_t)memptr[1] + 4095) & ~4095);
|
||||
dstBuffer = alignedmemptr[1];
|
||||
}
|
||||
else
|
||||
{
|
||||
err = hipMalloc(&dstBuffer, bufSize_);
|
||||
CHECK_RESULT(err != hipSuccess, "hipMalloc failed");
|
||||
}
|
||||
|
||||
CPerfCounter timer;
|
||||
|
||||
//warm up
|
||||
err = hipMemcpy2D(dstBuffer, width, srcBuffer, width, width, width, hipMemcpyDefault);
|
||||
CHECK_RESULT(err, "hipMemcpy2D failed");
|
||||
|
||||
timer.Reset();
|
||||
timer.Start();
|
||||
for (unsigned int i = 0; i < numIter; i++)
|
||||
{
|
||||
err = hipMemcpy2DAsync(dstBuffer, width, srcBuffer, width, width, width, hipMemcpyDefault, NULL);
|
||||
CHECK_RESULT(err, "hipMemcpyAsync2D failed");
|
||||
}
|
||||
err = hipDeviceSynchronize();
|
||||
CHECK_RESULT(err, "hipDeviceSynchronize failed");
|
||||
timer.Stop();
|
||||
double sec = timer.GetElapsedTime();
|
||||
|
||||
// Buffer copy bandwidth in GB/s
|
||||
double perf = ((double)bufSize_*numIter*(double)(1e-09)) / sec;
|
||||
|
||||
const char *strSrc = NULL;
|
||||
const char *strDst = NULL;
|
||||
if (hostMalloc[0])
|
||||
strSrc = "hHM";
|
||||
else if (hostRegister[0])
|
||||
strSrc = "hHR";
|
||||
else if (unpinnedMalloc[0])
|
||||
strSrc = "unp";
|
||||
else
|
||||
strSrc = "hM";
|
||||
|
||||
if (hostMalloc[1])
|
||||
strDst = "hHM";
|
||||
else if (hostRegister[1])
|
||||
strDst = "hHR";
|
||||
else if (unpinnedMalloc[1])
|
||||
strDst = "unp";
|
||||
else
|
||||
strDst = "hM";
|
||||
// Double results when src and dst are both on device
|
||||
if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) &&
|
||||
(!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1]))
|
||||
perf *= 2.0;
|
||||
// Double results when src and dst are both in sysmem
|
||||
if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) &&
|
||||
(hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1]))
|
||||
perf *= 2.0;
|
||||
|
||||
char buf[256];
|
||||
SNPRINTF(buf, sizeof(buf), "HIPPerfBufferCopyRectSpeed[%d]\t(%8d bytes)\ts:%s d:%s\ti:%4d\t(GB/s) perf\t%f",
|
||||
test, bufSize_, strSrc, strDst, numIter, (float)perf);
|
||||
printf("%s\n", buf);
|
||||
|
||||
//Free src
|
||||
if (hostMalloc[0])
|
||||
{
|
||||
hipHostFree(srcBuffer);
|
||||
}
|
||||
else if (hostRegister[0])
|
||||
{
|
||||
hipHostUnregister(srcBuffer);
|
||||
free(memptr[0]);
|
||||
}
|
||||
else if (unpinnedMalloc[0])
|
||||
{
|
||||
free(memptr[0]);
|
||||
}
|
||||
else
|
||||
{
|
||||
hipFree(srcBuffer);
|
||||
}
|
||||
|
||||
//Free dst
|
||||
if (hostMalloc[1])
|
||||
{
|
||||
hipHostFree(dstBuffer);
|
||||
}
|
||||
else if (hostRegister[1])
|
||||
{
|
||||
hipHostUnregister(dstBuffer);
|
||||
free(memptr[1]);
|
||||
}
|
||||
else if (unpinnedMalloc[1])
|
||||
{
|
||||
free(memptr[1]);
|
||||
}
|
||||
else
|
||||
{
|
||||
hipFree(dstBuffer);
|
||||
}
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -1,136 +0,0 @@
|
||||
/*
|
||||
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 <iostream>
|
||||
#include <chrono>
|
||||
#include "test_common.h"
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define arraySize 16
|
||||
|
||||
typedef struct d_uint16 {
|
||||
uint data[arraySize];
|
||||
} d_uint16;
|
||||
|
||||
__global__ void read_kernel(d_uint16 *src, ulong N, uint *dst) {
|
||||
|
||||
size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x ;
|
||||
|
||||
uint tmp = 0;
|
||||
for (size_t i = idx; i < N; i += stride) {
|
||||
for (size_t j = 0; j < arraySize; j++) {
|
||||
tmp += src[i].data[j];
|
||||
}
|
||||
}
|
||||
|
||||
atomicAdd(dst, tmp);
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
d_uint16 *dSrc;
|
||||
d_uint16 *hSrc;
|
||||
uint *dDst;
|
||||
uint *hDst;
|
||||
hipStream_t stream;
|
||||
ulong N = 4 * 1024 * 1024;
|
||||
uint nBytes = N * sizeof(d_uint16);
|
||||
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
cout << "info: didn't find any GPU! skipping the test!\n";
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int device = 0;
|
||||
HIPCHECK(hipSetDevice(device));
|
||||
hipDeviceProp_t props;
|
||||
HIPCHECK(hipGetDeviceProperties(&props, device));
|
||||
cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name <<
|
||||
" with " << props.multiProcessorCount << " CUs" << endl;
|
||||
|
||||
const unsigned threadsPerBlock = 64;
|
||||
const unsigned blocks = props.multiProcessorCount * 4;
|
||||
|
||||
uint inputData = 0x1;
|
||||
int nIter = 1000;
|
||||
|
||||
hSrc = new d_uint16[nBytes];
|
||||
HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
hDst = new uint;
|
||||
hDst[0] = 0;
|
||||
HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
for (int j = 0; j < arraySize; j++) {
|
||||
hSrc[i].data[j] = inputData;
|
||||
}
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(&dSrc, nBytes));
|
||||
HIPCHECK(hipMalloc(&dDst, sizeof(uint)));
|
||||
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst);
|
||||
HIPCHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost));
|
||||
hipDeviceSynchronize();
|
||||
|
||||
if (hDst[0] != (nBytes / sizeof(uint))) {
|
||||
cout << "info: Data validation failed for warm up run!" << endl;
|
||||
cout << "info: expected " << nBytes / sizeof(uint) << " got " << hDst[0] << endl;
|
||||
HIPCHECK(hipErrorUnknown);
|
||||
}
|
||||
|
||||
// measure performance based on host time
|
||||
auto all_start = chrono::steady_clock::now();
|
||||
|
||||
for(int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst);
|
||||
}
|
||||
hipDeviceSynchronize();
|
||||
|
||||
auto all_end = chrono::steady_clock::now();
|
||||
chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count();
|
||||
|
||||
cout << "info: average read speed of " << perf << " GB/s " << "achieved for memory size of " <<
|
||||
nBytes / (1024 * 1024) << " MB" << endl;
|
||||
|
||||
delete [] hSrc;
|
||||
delete hDst;
|
||||
hipFree(dSrc);
|
||||
hipFree(dDst);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -1,126 +0,0 @@
|
||||
/*
|
||||
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 <iostream>
|
||||
#include <chrono>
|
||||
#include "test_common.h"
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define arraySize 16
|
||||
|
||||
typedef struct d_uint16 {
|
||||
uint data[arraySize];
|
||||
} d_uint16;
|
||||
|
||||
__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) {
|
||||
size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = idx; i < N; i += stride) {
|
||||
dst[i] = pval;
|
||||
}
|
||||
};
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
d_uint16 *dDst;
|
||||
d_uint16 *hDst;
|
||||
hipStream_t stream;
|
||||
ulong N = 4 * 1024 * 1024;
|
||||
uint nBytes = N * sizeof(d_uint16);
|
||||
d_uint16 pval;
|
||||
|
||||
for (int i = 0; i < arraySize; i++) {
|
||||
pval.data[i] = 0xabababab;
|
||||
}
|
||||
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
cout << "info: didn't find any GPU! skipping the test!\n";
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int device = 0;
|
||||
HIPCHECK(hipSetDevice(device));
|
||||
hipDeviceProp_t props;
|
||||
HIPCHECK(hipGetDeviceProperties(&props, device));
|
||||
cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name <<
|
||||
" with " << props.multiProcessorCount << " CUs" << endl;
|
||||
|
||||
size_t threadsPerBlock = 64;
|
||||
size_t blocks = props.multiProcessorCount * 4;
|
||||
|
||||
uint inputData = 0xabababab;
|
||||
int nIter = 1000;
|
||||
|
||||
hDst = new d_uint16[nBytes];
|
||||
HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
for (size_t j = 0; j < arraySize; j++) {
|
||||
hDst[i].data[j] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(&dDst, nBytes));
|
||||
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval);
|
||||
HIPCHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost));
|
||||
hipDeviceSynchronize();
|
||||
|
||||
for (uint i = 0; i < N; i++) {
|
||||
for (uint j = 0; j < arraySize; j++) {
|
||||
if (hDst[i].data[j] != inputData) {
|
||||
cout << "info: Data validation failed for warm up run! " << endl;
|
||||
cout << "at index i: " << i << " element j: " << j << endl;
|
||||
cout << hex << "expected 0x" << inputData << " but got 0x" << hDst[i].data[j] << endl;
|
||||
HIPCHECK(hipErrorUnknown);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
auto all_start = chrono::steady_clock::now();
|
||||
for(int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval);
|
||||
}
|
||||
hipDeviceSynchronize();
|
||||
auto all_end = chrono::steady_clock::now();
|
||||
chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count();
|
||||
|
||||
cout << "info: average write speed of " << perf << " GB/s " << "achieved for memory size of " <<
|
||||
nBytes / (1024 * 1024) << " MB" << endl;
|
||||
|
||||
|
||||
delete [] hDst;
|
||||
hipFree(dDst);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -1,190 +0,0 @@
|
||||
/*
|
||||
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_CMD: hipPerfHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvidia
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#include <iostream>
|
||||
#include <time.h>
|
||||
#include <cstdio>
|
||||
#include <unistd.h>
|
||||
#include <sys/types.h>
|
||||
#include <sys/mman.h>
|
||||
#include <sys/stat.h>
|
||||
#include <fcntl.h>
|
||||
#include <numaif.h>
|
||||
#include <iostream>
|
||||
#include <memory>
|
||||
#include <stdexcept>
|
||||
#include <string>
|
||||
#include <array>
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
// To run it correctly, we must not export HIP_VISIBLE_DEVICES.
|
||||
// And we must explicitly link libnuma because of numa api move_pages().
|
||||
#define NUM_PAGES 4
|
||||
char *h = nullptr;
|
||||
char *d_h = nullptr;
|
||||
char *m = nullptr;
|
||||
char *d_m = nullptr;
|
||||
int page_size = 0;
|
||||
const int mode[] = { MPOL_DEFAULT, MPOL_BIND, MPOL_PREFERRED, MPOL_INTERLEAVE };
|
||||
const char* modeStr[] = { "MPOL_DEFAULT", "MPOL_BIND", "MPOL_PREFERRED", "MPOL_INTERLEAVE" };
|
||||
|
||||
std::string exeCommand(const char* cmd) {
|
||||
std::array<char, 128> buff;
|
||||
std::string result;
|
||||
std::unique_ptr<FILE, decltype(&pclose)> pipe(popen(cmd, "r"), pclose);
|
||||
if (!pipe) {
|
||||
return result;
|
||||
}
|
||||
while (fgets(buff.data(), buff.size(), pipe.get()) != nullptr) {
|
||||
result += buff.data();
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
int getCpuAgentCount() {
|
||||
const char* cmd = "cat /proc/cpuinfo | grep \"physical id\" | sort | uniq | wc -l";
|
||||
int cpuAgentCount = std::atoi(exeCommand(cmd).c_str());
|
||||
return cpuAgentCount;
|
||||
}
|
||||
|
||||
bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) {
|
||||
void *pages[NUM_PAGES];
|
||||
int status[NUM_PAGES];
|
||||
int nodes[NUM_PAGES];
|
||||
int ret_code;
|
||||
|
||||
printf("set cpu %d, gpu %d, numaMode %d, hostMallocflags 0x%x\n", cpuId,
|
||||
gpuId, numaMode, hostMallocflags);
|
||||
|
||||
if (cpuId >= 0) {
|
||||
unsigned long nodeMask = 1 << cpuId;
|
||||
unsigned long maxNode = sizeof(nodeMask) * 8;
|
||||
if (set_mempolicy(numaMode, numaMode == MPOL_DEFAULT ? NULL : &nodeMask,
|
||||
numaMode == MPOL_DEFAULT ? 0 : maxNode) == -1) {
|
||||
printf("set_mempolicy() failed with err %d\n", errno);
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (gpuId >= 0) {
|
||||
HIPCHECK(hipSetDevice(gpuId));
|
||||
}
|
||||
|
||||
posix_memalign((void**) &m, page_size, page_size * NUM_PAGES);
|
||||
hipHostRegister(m, page_size * NUM_PAGES, hipHostRegisterMapped);
|
||||
hipHostGetDevicePointer((void**) &d_m, m, 0);
|
||||
|
||||
status[0] = -1;
|
||||
pages[0] = m;
|
||||
for (int i = 1; i < NUM_PAGES; i++) {
|
||||
pages[i] = (char*) pages[0] + page_size;
|
||||
}
|
||||
ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0);
|
||||
printf("Memory (malloc) ret %d at %p (dev %p) is at node: ", ret_code, m, d_m);
|
||||
for (int i = 0; i < NUM_PAGES; i++) {
|
||||
printf("%d ", status[i]); // Don't verify as it's out of our control
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
HIPCHECK(hipHostMalloc((void**) &h, page_size*NUM_PAGES, hostMallocflags));
|
||||
pages[0] = h;
|
||||
for (int i = 1; i < NUM_PAGES; i++) {
|
||||
pages[i] = (char*) pages[0] + page_size;
|
||||
}
|
||||
ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0);
|
||||
d_h = nullptr;
|
||||
if (hostMallocflags & hipHostMallocMapped) {
|
||||
hipHostGetDevicePointer((void**) &d_h, h, 0);
|
||||
printf("Memory (hipHostMalloc) ret %d at %p (dev %p) is at node: ",
|
||||
ret_code, h, d_h);
|
||||
} else {
|
||||
printf("Memory (hipHostMalloc) ret %d at %p is at node: ", ret_code, h);
|
||||
}
|
||||
for (int i = 0; i < NUM_PAGES; i++) {
|
||||
printf("%d ", status[i]); // Always print it even if it's wrong. Verify later
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
HIPCHECK(hipHostFree((void* )h));
|
||||
hipHostUnregister(m);
|
||||
free(m);
|
||||
|
||||
if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) {
|
||||
for (int i = 0; i < NUM_PAGES; i++) {
|
||||
if (status[i] != cpuId) { // Now verify
|
||||
printf("Failed at %d", i);
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
bool runTest(const int &cpuCount, const int &gpuCount,
|
||||
const unsigned int &hostMallocflags, const std::string &str) {
|
||||
printf("%s\n", str.c_str());
|
||||
|
||||
for (int m = 0; m < sizeof(mode) / sizeof(mode[0]); m++) {
|
||||
printf("Testing %s\n", modeStr[m]);
|
||||
|
||||
for (int i = 0; i < cpuCount; i++) {
|
||||
for (int j = 0; j < gpuCount; j++) {
|
||||
if (!test(i, j, mode[m], hostMallocflags)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
int gpuCount = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&gpuCount));
|
||||
int cpuCount = getCpuAgentCount();
|
||||
page_size = getpagesize();
|
||||
printf("Cpu count %d, Gpu count %d, Page size %d\n", cpuCount, gpuCount,
|
||||
page_size);
|
||||
|
||||
if (cpuCount < 0 || gpuCount < 0) {
|
||||
failed("Bad device count\n");
|
||||
return -1;
|
||||
}
|
||||
|
||||
if (!runTest(cpuCount, gpuCount, hipHostMallocDefault | hipHostMallocNumaUser,
|
||||
"Testing hipHostMallocDefault | hipHostMallocNumaUser........................")) {
|
||||
failed("Failed testing hipHostMallocDefault | hipHostMallocNumaUser\n");
|
||||
return -1;
|
||||
}
|
||||
|
||||
if (!runTest(cpuCount, gpuCount, hipHostMallocMapped | hipHostMallocNumaUser,
|
||||
"Testing hipHostMallocMapped | hipHostMallocNumaUser.........................")) {
|
||||
failed("Failed testing hipHostMallocMapped | hipHostMallocNumaUser\n");
|
||||
return -1;
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -1,534 +0,0 @@
|
||||
/*
|
||||
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 "test_common.h"
|
||||
#include <printf/printf_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
|
||||
}
|
||||
}
|
||||
|
||||
__device__ void print_log(int i, double value, double expected) {
|
||||
printf("failed at %d: val=%g, expected=%g\n", i, value, expected);
|
||||
}
|
||||
|
||||
__device__ void print_log(int i, int value, int expected) {
|
||||
printf("failed at %d: val=%d, expected=%d\n", i, value, expected);
|
||||
}
|
||||
|
||||
template<class T>
|
||||
__global__ void vec_verify(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
|
||||
if(x[i] != coef) {
|
||||
print_log(i, x[i], coef);
|
||||
}
|
||||
#else
|
||||
if(x[i] != coef * i) {
|
||||
print_log(i, 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) {
|
||||
failed("No GPU!");
|
||||
} else if (deviceId >= nGpu) {
|
||||
failed("Info: wrong GPU Id %d\n", deviceId);
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(deviceId));
|
||||
memset(&props_, 0, sizeof(props_));
|
||||
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 log_host(const char* title, double GBytes, double sec) {
|
||||
cout << title << " [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl;
|
||||
}
|
||||
|
||||
void log_kernel(const char* title, double GBytes, double sec, double sec_hv, double sec_kv) {
|
||||
cout << title << " [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec
|
||||
<< " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", hostVerify cost "
|
||||
<< setw(10) << sec_hv << " s in bandwidth " << setw(10) << GBytes / sec_hv << " [GB/s]"
|
||||
<< ", kernelVerify cost "<< setw(10) << sec_kv << " s in bandwidth " << setw(10)
|
||||
<< GBytes / sec_kv << " [GB/s]" << 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 hostVerify(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 << "hostVerify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef << endl;
|
||||
failed("failed\n");
|
||||
}
|
||||
#else
|
||||
if(data[i] != coef * i) {
|
||||
cout << "hostVerify 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();
|
||||
}
|
||||
|
||||
void kernelVerify(size_t size, T *data, T coef, double &sec) {
|
||||
size_t num = size / sizeof(T); // Size of elements
|
||||
unsigned blocks = setNumBlocks(num);
|
||||
|
||||
CaptureStream *capture = new CaptureStream(stdout);
|
||||
capture->Begin();
|
||||
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify<T>), dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, 0, data, coef, num); // kernel will be loaded first time
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
capture->End();
|
||||
capture->Truncate(1000); // Don't want too long log if existing
|
||||
std::string device_output = capture->getData();
|
||||
delete capture;
|
||||
if (device_output.length() > 0) {
|
||||
failed("kernelVerify failed:\n%s\n", device_output.c_str());
|
||||
}
|
||||
|
||||
// Now all data verified. The following is to test bandwidth.
|
||||
auto start = chrono::steady_clock::now();
|
||||
|
||||
for (int iter = 0; iter < NUM_ITER; ++iter) {
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify<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
|
||||
}
|
||||
|
||||
bool testLargeBarDeviceMemoryHostFill(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));
|
||||
|
||||
log_host("Largebar: host fill", GBytes, sec);
|
||||
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 (!testLargeBarDeviceMemoryHostFill(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));
|
||||
|
||||
log_host("Managed: host fill", GBytes, sec);
|
||||
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_hv = 0, sec_kv = 0;
|
||||
kernelFill(size, A, coef_, sec);
|
||||
hostVerify(size, A, coef_, sec_hv); // Managed memory can be verified by host
|
||||
kernelVerify(size, A, coef_, sec_kv);
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
log_kernel("Managed: kernel fill", GBytes, sec, sec_hv, sec_kv);
|
||||
|
||||
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));
|
||||
|
||||
log_host("Host: host fill", GBytes, sec);
|
||||
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_hv = 0, sec_kv = 0;
|
||||
kernelFill(size, A, coef_, sec);
|
||||
hostVerify(size, A, coef_, sec_hv);
|
||||
kernelVerify(size, A, coef_, sec_kv);
|
||||
HIPCHECK(hipHostFree(A));
|
||||
|
||||
log_kernel("Host: kernel fill", GBytes, sec, sec_hv, sec_kv);
|
||||
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 function should be via device attribute query*/
|
||||
bool supportDeviceMallocFinegrained() {
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
T *A = nullptr;
|
||||
hipExtMallocWithFlags((void **)&A, sizeof(T), hipDeviceMallocFinegrained);
|
||||
if (!A) {
|
||||
return false;
|
||||
}
|
||||
HIPCHECK(hipFree(A));
|
||||
return true;
|
||||
#else
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
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
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
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) {
|
||||
cout << "failed hipExtMallocWithFlags() with size =" << size << " flags="
|
||||
<< std::hex << flags << endl;
|
||||
return false;
|
||||
}
|
||||
|
||||
double sec = 0;
|
||||
hostFill(size, A, coef_, sec); // Cpu can access this mem
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
log_host("ExtDevice: host fill", GBytes, sec);
|
||||
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) {
|
||||
cout << "failed hipExtMallocWithFlags() with size =" << size << " flags="
|
||||
<< std::hex << flags << endl;
|
||||
return false;
|
||||
}
|
||||
|
||||
double sec = 0, sec_hv = 0, sec_kv = 0;
|
||||
kernelFill(size, A, coef_, sec);
|
||||
hostVerify(size, A, coef_, sec_hv); // Fine grained device memory can be verified by host
|
||||
kernelVerify(size, A, coef_, sec_kv);
|
||||
HIPCHECK(hipFree(A));
|
||||
|
||||
log_kernel("ExtDevice: kernel fill", GBytes, sec, sec_hv, sec_kv);
|
||||
|
||||
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;
|
||||
}
|
||||
#endif
|
||||
|
||||
bool run() {
|
||||
if (supportLargeBar()) {
|
||||
if (!testLargeBar()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (supportManagedMemory()) {
|
||||
if (!testManagedMemory()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (!testHostMemory()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
if (supportDeviceMallocFinegrained()) {
|
||||
if (!testExtDeviceMemory()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
#endif
|
||||
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();
|
||||
}
|
||||
@@ -1,124 +0,0 @@
|
||||
/*
|
||||
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 "test_common.h"
|
||||
#include <iostream>
|
||||
#include <time.h>
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../src/test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#define NUM_SIZE 19 //size up to 16M
|
||||
#define NUM_ITER 500 //Total GPU memory up to 16M*500=8G
|
||||
|
||||
void valSet(int* A, int val, size_t size) {
|
||||
size_t len = size / sizeof(int);
|
||||
for (int i = 0; i < len; i++) {
|
||||
A[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
void setup(size_t *size, int &num, int **pA, const size_t totalGlobalMem) {
|
||||
|
||||
std::cout << "size: ";
|
||||
for (int i = 0; i < num; i++) {
|
||||
size[i] = 1 << (i + 6);
|
||||
if((NUM_ITER + 1) * size[i] > totalGlobalMem) {
|
||||
num = i;
|
||||
break;
|
||||
}
|
||||
std::cout << size[i] << " ";
|
||||
}
|
||||
std::cout << std::endl;
|
||||
*pA = (int*)malloc(size[num - 1]);
|
||||
valSet(*pA, 1, size[num - 1]);
|
||||
}
|
||||
|
||||
void testInit(size_t size, int *A) {
|
||||
int *Ad;
|
||||
clock_t start = clock();
|
||||
hipMalloc(&Ad, size); //hip::init() will be called
|
||||
clock_t end = clock();
|
||||
double uS = (end - start) * 1000000. / CLOCKS_PER_SEC;
|
||||
std::cout << "Initial" << std::endl;
|
||||
std::cout << "hipMalloc(" << size << ") cost " << uS << "us" << std::endl;
|
||||
|
||||
start = clock();
|
||||
hipMemcpy(Ad, A, size, hipMemcpyHostToDevice);
|
||||
hipDeviceSynchronize();
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / CLOCKS_PER_SEC;
|
||||
std::cout << "hipMemcpy(" << size << ") cost " << uS << "us" << std::endl;
|
||||
|
||||
start = clock();
|
||||
hipFree(Ad);
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / CLOCKS_PER_SEC;
|
||||
std::cout << "hipFree(" << size << ") cost " << uS << "us" << std::endl;
|
||||
}
|
||||
|
||||
int main() {
|
||||
double uS;
|
||||
clock_t start, end;
|
||||
size_t size[NUM_SIZE] = { 0 };
|
||||
int *Ad[NUM_ITER] = { nullptr };
|
||||
int *A;
|
||||
hipDeviceProp_t props;
|
||||
memset(&props, 0, sizeof(props));
|
||||
HIPCHECK(hipGetDeviceProperties(&props, 0));
|
||||
std::cout << "totalGlobalMem: " << props.totalGlobalMem << std::endl;
|
||||
|
||||
int num = NUM_SIZE;
|
||||
setup(size, num, &A, props.totalGlobalMem);
|
||||
testInit(size[0], A);
|
||||
|
||||
for (int i = 0; i < num; i++) {
|
||||
std::cout << size[i] << std::endl;
|
||||
start = clock();
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIPCHECK(hipMalloc(&Ad[j], size[i]));
|
||||
}
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC);
|
||||
std::cout << "hipMalloc(" << size[i] << ") cost " << uS << "us" << std::endl;
|
||||
|
||||
start = clock();
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIPCHECK(hipMemcpy(Ad[j], A, size[i], hipMemcpyHostToDevice));
|
||||
}
|
||||
hipDeviceSynchronize();
|
||||
end = clock();
|
||||
uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC);
|
||||
std::cout << "hipMemcpy(" << size[i] << ") cost " << uS << "us" << std::endl;
|
||||
|
||||
start = clock();
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
HIPCHECK(hipFree(Ad[j]));
|
||||
Ad[j] = nullptr;
|
||||
}
|
||||
end = clock();
|
||||
double uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC);
|
||||
std::cout << "hipFree(" << size[i] << ") cost " << uS << "us" << std::endl;
|
||||
}
|
||||
free(A);
|
||||
passed();
|
||||
}
|
||||
@@ -1,114 +0,0 @@
|
||||
/*
|
||||
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 "test_common.h"
|
||||
#include <iostream>
|
||||
#include <chrono>
|
||||
|
||||
#define NUM_SIZE 8
|
||||
#define NUM_ITER 0x40000
|
||||
|
||||
|
||||
using namespace std;
|
||||
|
||||
class hipPerfMemcpy {
|
||||
private:
|
||||
unsigned int numBuffers_;
|
||||
size_t totalSizes_[NUM_SIZE];
|
||||
void setHostBuffer(int *A, int val, size_t size);
|
||||
public:
|
||||
hipPerfMemcpy();
|
||||
~hipPerfMemcpy() {};
|
||||
void open(int deviceID);
|
||||
void run(unsigned int testNumber);
|
||||
};
|
||||
|
||||
hipPerfMemcpy::hipPerfMemcpy() : numBuffers_(0) {
|
||||
for (int i = 0; i < NUM_SIZE; i++) {
|
||||
totalSizes_[i] = 1 << (i + 6);
|
||||
}
|
||||
};
|
||||
|
||||
void hipPerfMemcpy::setHostBuffer(int *A, int val, size_t size) {
|
||||
size_t len = size / sizeof(int);
|
||||
for (int i = 0; i < len; i++) {
|
||||
A[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
void hipPerfMemcpy::open(int deviceId) {
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
failed("No GPU!");
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
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;
|
||||
}
|
||||
|
||||
void hipPerfMemcpy::run(unsigned int testNumber) {
|
||||
int *A, *Ad;
|
||||
A = new int[totalSizes_[testNumber]];
|
||||
setHostBuffer(A, 1, totalSizes_[testNumber]);
|
||||
hipMalloc(&Ad, totalSizes_[testNumber]);
|
||||
|
||||
auto start = chrono::steady_clock::now();
|
||||
|
||||
for (int j = 0; j < NUM_ITER; j++) {
|
||||
hipMemcpy(Ad, A, totalSizes_[testNumber], hipMemcpyHostToDevice);
|
||||
}
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
auto end = chrono::steady_clock::now();
|
||||
chrono::duration<double, micro> diff = end - start;
|
||||
|
||||
cout << "hipPerfMemcpy[" << testNumber << "] " << "Host to Device copy took "
|
||||
<< diff.count() / NUM_ITER << " us for memory size of " << totalSizes_[testNumber]
|
||||
<< " Bytes" << endl;
|
||||
|
||||
delete [] A;
|
||||
HIPCHECK(hipFree(Ad));
|
||||
|
||||
}
|
||||
|
||||
|
||||
int main() {
|
||||
hipPerfMemcpy hipPerfMemcpy;
|
||||
|
||||
int deviceId = 0;
|
||||
hipPerfMemcpy.open(deviceId);
|
||||
|
||||
for (auto testCase = 0; testCase < NUM_SIZE; testCase++) {
|
||||
hipPerfMemcpy.run(testCase);
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
@@ -1,437 +0,0 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../src/test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#include <iostream>
|
||||
#include <chrono>
|
||||
|
||||
static unsigned int sizeList[] = {
|
||||
256, 512, 1024, 2048, 4096, 8192,
|
||||
};
|
||||
|
||||
static unsigned int eleNumList[] = {
|
||||
0x100, 0x400, 0x1000, 0x4000, 0x10000, 0x20000, 0x40000, 0x80000, 0x100000,
|
||||
0x200000, 0x400000, 0x800000, 0x1000000
|
||||
};
|
||||
|
||||
typedef struct _dataType {
|
||||
char memsetval = 0x42;
|
||||
char memsetD8val = 0xDE;
|
||||
int16_t memsetD16val = 0xDEAD;
|
||||
int memsetD32val = 0xDEADBEEF;
|
||||
}dataType;
|
||||
|
||||
#define NUM_ITER 1000
|
||||
|
||||
enum MemsetType {
|
||||
hipMemsetTypeDefault,
|
||||
hipMemsetTypeD8,
|
||||
hipMemsetTypeD16,
|
||||
hipMemsetTypeD32,
|
||||
hipMemsetTypeMax
|
||||
|
||||
};
|
||||
|
||||
using namespace std;
|
||||
|
||||
class hipPerfMemset {
|
||||
private:
|
||||
uint64_t bufSize_;
|
||||
unsigned int num_elements_;
|
||||
unsigned int testNumEle_;
|
||||
unsigned int _numSubTests = 0;
|
||||
unsigned int _numSubTests2D = 0;
|
||||
unsigned int _numSubTests3D = 0;
|
||||
unsigned int num_sizes_ =0;
|
||||
|
||||
public:
|
||||
hipPerfMemset() {
|
||||
num_elements_ = sizeof(eleNumList) / sizeof(unsigned int);
|
||||
_numSubTests = num_elements_ * hipMemsetTypeMax;
|
||||
|
||||
num_sizes_ = sizeof(sizeList) / sizeof(unsigned int);
|
||||
_numSubTests2D = num_sizes_;
|
||||
_numSubTests3D = _numSubTests2D;
|
||||
};
|
||||
|
||||
~hipPerfMemset() {};
|
||||
|
||||
void open(int deviceID);
|
||||
|
||||
template<typename T>
|
||||
void run1D(unsigned int test, T memsetval, enum MemsetType type, bool async);
|
||||
|
||||
template<typename T>
|
||||
void run2D(unsigned int test, T memsetval, enum MemsetType type, bool async);
|
||||
|
||||
template<typename T>
|
||||
void run3D(unsigned int test, T memsetval, enum MemsetType type, bool async);
|
||||
|
||||
uint getNumTests() {
|
||||
return _numSubTests;
|
||||
}
|
||||
|
||||
uint getNumTests2D() {
|
||||
return _numSubTests2D;
|
||||
}
|
||||
uint getNumTests3D() {
|
||||
return _numSubTests3D;
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
void hipPerfMemset::open(int deviceId) {
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
failed("No GPU!");
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(deviceId));
|
||||
hipDeviceProp_t props = {0};
|
||||
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;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void hipPerfMemset::run1D(unsigned int test, T memsetval, enum MemsetType type, bool async) {
|
||||
|
||||
T * A_h;
|
||||
T * A_d;
|
||||
|
||||
testNumEle_ = eleNumList[test % num_elements_];
|
||||
|
||||
bufSize_ = testNumEle_ * sizeof(uint32_t);
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, bufSize_));
|
||||
|
||||
A_h = reinterpret_cast<T*> (malloc(bufSize_));
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
|
||||
|
||||
// Warm-up
|
||||
if (async) {
|
||||
HIPCHECK(hipMemsetAsync((void *)A_d, memsetval, bufSize_, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
} else {
|
||||
HIPCHECK(hipMemset((void *)A_d, memsetval, bufSize_));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
auto start = chrono::high_resolution_clock::now();
|
||||
for (uint i = 0; i < NUM_ITER; i++) {
|
||||
if (type == hipMemsetTypeDefault && !async) {
|
||||
HIPCHECK(hipMemset((void *)A_d, memsetval, bufSize_));
|
||||
}
|
||||
else if (type == hipMemsetTypeDefault && async) {
|
||||
HIPCHECK(hipMemsetAsync(A_d, memsetval, bufSize_, stream));
|
||||
}
|
||||
else if (type == hipMemsetTypeD8 && !async){
|
||||
HIPCHECK(hipMemsetD8((hipDeviceptr_t)A_d, memsetval, bufSize_));
|
||||
}
|
||||
else if (type == hipMemsetTypeD8 && async) {
|
||||
HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memsetval, bufSize_, stream));
|
||||
}
|
||||
else if (type == hipMemsetTypeD16 && !async) {
|
||||
HIPCHECK(hipMemsetD16((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T)));
|
||||
}
|
||||
else if (type == hipMemsetTypeD16 && async) {
|
||||
HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T), stream));
|
||||
}
|
||||
else if (type == hipMemsetTypeD32 && !async) {
|
||||
HIPCHECK(hipMemsetD32((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T)));
|
||||
}
|
||||
else if (type == hipMemsetTypeD32 && async) {
|
||||
HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T), stream));
|
||||
}
|
||||
}
|
||||
if (async) {
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
} else {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
auto end = chrono::high_resolution_clock::now();
|
||||
|
||||
HIPCHECK(hipMemcpy(A_h, A_d, bufSize_, hipMemcpyDeviceToHost) );
|
||||
|
||||
for (int i = 0; i < bufSize_ / sizeof(T); i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
cout << "mismatch at index " << i << " computed: " << static_cast<int> (A_h[i])
|
||||
<< ", memsetval: " << static_cast<int> (memsetval) << endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
HIPCHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
|
||||
auto diff = std::chrono::duration<double>(end - start);
|
||||
auto sec = diff.count();
|
||||
|
||||
auto perf = static_cast<double>((bufSize_ * NUM_ITER * (double)(1e-09)) / sec);
|
||||
|
||||
cout << "[" << setw(2) << test << "] " << setw(5) << bufSize_/1024 << " Kb " << setw(4)
|
||||
<< " typeSize " << (int)sizeof(T) << " : " << setw(7) << perf << " GB/s " << endl;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void hipPerfMemset::run2D(unsigned int test, T memsetval, enum MemsetType type, bool async) {
|
||||
|
||||
bufSize_ = sizeList[test % num_sizes_];
|
||||
|
||||
size_t numH = bufSize_;
|
||||
size_t numW = bufSize_;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW* numH;
|
||||
|
||||
T * A_h;
|
||||
T * A_d;
|
||||
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width ,
|
||||
numH));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
|
||||
for (size_t i=0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
|
||||
|
||||
// Warm-up
|
||||
if (async) {
|
||||
HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
} else {
|
||||
HIPCHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
auto start = chrono::steady_clock::now();
|
||||
|
||||
for (uint i = 0; i < NUM_ITER; i++) {
|
||||
if (type == hipMemsetTypeDefault && !async) {
|
||||
HIPCHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
|
||||
}
|
||||
else if (type == hipMemsetTypeDefault && async) {
|
||||
HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream));
|
||||
}
|
||||
}
|
||||
|
||||
if (async) {
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
} else {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
auto end = chrono::steady_clock::now();
|
||||
|
||||
HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i=0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
cout << "mismatch at index " << i << " computed: " << static_cast<int> (A_h[i])
|
||||
<< ", memsetval: " << static_cast<int> (memsetval) << endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
chrono::duration<double> diff = end - start;
|
||||
|
||||
auto sec = diff.count();
|
||||
|
||||
auto perf = static_cast<double>((sizeElements* NUM_ITER * (double)(1e-09)) / sec);
|
||||
|
||||
cout << " hipPerf2DMemset" << (async ? "Async" : " ") << "[" << test << "] "
|
||||
<< " " << "(GB/s) for " << setw(5) << bufSize_
|
||||
<< " x " << setw(5) << bufSize_ << " bytes : " << setw(7) << perf << endl;
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void hipPerfMemset::run3D(unsigned int test, T memsetval, enum MemsetType type, bool async) {
|
||||
|
||||
bufSize_ = sizeList[test % num_sizes_];
|
||||
|
||||
size_t numH = bufSize_;
|
||||
size_t numW = bufSize_;
|
||||
size_t depth = 10;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH * depth;
|
||||
size_t elements = numW* numH* depth;
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
|
||||
|
||||
T *A_h;
|
||||
|
||||
hipExtent extent = make_hipExtent(width, numH, depth);
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
|
||||
HIPCHECK(hipMalloc3D(&devPitchedPtr, extent));
|
||||
A_h = (char*)malloc(sizeElements);
|
||||
HIPASSERT(A_h != NULL);
|
||||
|
||||
for (size_t i=0; i<elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
// Warm-up
|
||||
if (async) {
|
||||
HIPCHECK(hipMemset3DAsync( devPitchedPtr, memsetval, extent, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
} else {
|
||||
HIPCHECK(hipMemset3D( devPitchedPtr, memsetval, extent));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
auto start = chrono::steady_clock::now();
|
||||
|
||||
for (uint i = 0; i < NUM_ITER; i++) {
|
||||
if (type == hipMemsetTypeDefault && !async) {
|
||||
HIPCHECK(hipMemset3D( devPitchedPtr, memsetval, extent));
|
||||
}
|
||||
else if (type == hipMemsetTypeDefault && async) {
|
||||
HIPCHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
|
||||
}
|
||||
}
|
||||
|
||||
if (async) {
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
} else {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
auto end = chrono::steady_clock::now();
|
||||
|
||||
hipMemcpy3DParms myparms = {0};
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
|
||||
myparms.srcPtr = devPitchedPtr;
|
||||
myparms.extent = extent;
|
||||
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
|
||||
HIPCHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int i=0; i<elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
cout << "mismatch at index " << i << " computed: " << static_cast<int> (A_h[i])
|
||||
<< ", memsetval: " << static_cast<int> (memsetval) << endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
chrono::duration<double> diff = end - start;
|
||||
|
||||
auto sec = diff.count();
|
||||
|
||||
auto perf = static_cast<double>((sizeElements * NUM_ITER * (double)(1e-09)) / sec);
|
||||
|
||||
cout << " hipPerf3DMemset" << (async ? "Async" : " ") << "[" << test << "] " << " "
|
||||
<< "(GB/s) for " << setw(5) << bufSize_ << " x " << setw(5)
|
||||
<< bufSize_ << " x " << depth << " bytes : " << setw(7) << perf << endl;
|
||||
HIPCHECK(hipFree(devPitchedPtr.ptr));
|
||||
free(A_h);
|
||||
}
|
||||
|
||||
int main() {
|
||||
hipPerfMemset hipPerfMemset;
|
||||
|
||||
dataType pattern;
|
||||
int deviceId = 0;
|
||||
hipPerfMemset.open(deviceId);
|
||||
MemsetType type;
|
||||
|
||||
int numTests = hipPerfMemset.getNumTests();
|
||||
int numTests2D = hipPerfMemset.getNumTests2D();
|
||||
int numTests3D = hipPerfMemset.getNumTests3D();
|
||||
|
||||
|
||||
cout << "--------------------- 1D buffer -------------------" << endl;
|
||||
bool async= false;
|
||||
for (uint i = 0; i < 2 ; i++) {
|
||||
cout << endl;
|
||||
|
||||
for (auto testCase = 0; testCase < numTests; testCase++) {
|
||||
if (testCase < sizeof(eleNumList) / sizeof(uint32_t)) {
|
||||
cout << "API: hipMemsetD8" << (async ? "Async " : " ");
|
||||
hipPerfMemset.run1D(testCase, pattern.memsetval, hipMemsetTypeD8, async);
|
||||
}
|
||||
|
||||
else if (testCase < 2 * sizeof(eleNumList) / sizeof(uint32_t)) {
|
||||
cout << "API: hipMemsetD16" << (async ? "Async" : " ");
|
||||
hipPerfMemset.run1D(testCase,pattern.memsetD16val, hipMemsetTypeD16, async);
|
||||
}
|
||||
|
||||
else if (testCase < 3 * sizeof(eleNumList) / sizeof(uint32_t)) {
|
||||
cout << "API: hipMemsetD32" << (async ? "Async" : " ");
|
||||
hipPerfMemset.run1D(testCase,pattern.memsetD32val, hipMemsetTypeD32, async);
|
||||
}
|
||||
|
||||
else {
|
||||
cout << "API: hipMemset" << (async ? "Async " : " ");
|
||||
hipPerfMemset.run1D(testCase,pattern.memsetval, hipMemsetTypeDefault, async);
|
||||
}
|
||||
}
|
||||
async = true;
|
||||
}
|
||||
|
||||
cout << endl;
|
||||
cout << "------------------ 2D buffer arrays ---------------" << endl;
|
||||
|
||||
async = false;
|
||||
for (uint i = 0; i < 2; i++) {
|
||||
cout << endl;
|
||||
for (uint test = 0; test < numTests2D; test++) {
|
||||
hipPerfMemset.run2D(test, pattern.memsetval, hipMemsetTypeDefault, async);
|
||||
}
|
||||
async = true;
|
||||
}
|
||||
|
||||
cout << endl;
|
||||
cout << "------------------ 3D buffer arrays ---------------" << endl;
|
||||
|
||||
async = false;
|
||||
for (uint i = 0; i < 2; i++) {
|
||||
cout << endl;
|
||||
for (uint test =0; test < numTests3D; test++) {
|
||||
hipPerfMemset.run3D(test, pattern.memsetval, hipMemsetTypeDefault, async);
|
||||
}
|
||||
async = true;
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -1,319 +0,0 @@
|
||||
/*
|
||||
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 <iostream>
|
||||
#include <chrono>
|
||||
#include "test_common.h"
|
||||
#include <hip/hip_vector_types.h>
|
||||
#include <vector>
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define NUM_TYPES 3
|
||||
vector<string> types= {"float", "float2", "float4"};
|
||||
vector<unsigned int> typeSizes = {4, 8, 16};
|
||||
|
||||
#define NUM_SIZES 12
|
||||
vector<unsigned int> sizes = {1, 2, 4, 8, 16, 32,
|
||||
64, 128, 256, 512, 1024, 2048};
|
||||
|
||||
#define NUM_BUFS 6
|
||||
#define MAX_BUFS (1 << (NUM_BUFS - 1))
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVIDIA__
|
||||
inline __host__ __device__ void operator+=(float2 &a, float2 b)
|
||||
{
|
||||
a.x += b.x; a.y += b.y;
|
||||
}
|
||||
|
||||
inline __host__ __device__ void operator+=(float4 &a, float4 b)
|
||||
{
|
||||
a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w;
|
||||
}
|
||||
#endif
|
||||
|
||||
template <typename T>
|
||||
__global__ void sampleRate(T * outBuffer, unsigned int inBufSize, unsigned int writeIt,
|
||||
T **inBuffer, int numBufs) {
|
||||
|
||||
uint gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
uint inputIdx = gid % inBufSize;
|
||||
|
||||
T tmp;
|
||||
memset(&tmp, 0, sizeof(T));
|
||||
for(int i = 0; i < numBufs; i++) {
|
||||
tmp += *(*(inBuffer+i)+inputIdx);
|
||||
}
|
||||
|
||||
if (writeIt*(unsigned int)tmp.x) {
|
||||
outBuffer[gid] = tmp;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
__global__ void sampleRateFloat(T * outBuffer, unsigned int inBufSize, unsigned int writeIt,
|
||||
T ** inBuffer, int numBufs) {
|
||||
|
||||
uint gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
uint inputIdx = gid % inBufSize;
|
||||
|
||||
T tmp = (T)0.0f;
|
||||
|
||||
for(int i = 0; i < numBufs; i++) {
|
||||
tmp += *((*inBuffer+i)+inputIdx);
|
||||
}
|
||||
|
||||
if (writeIt*(unsigned int)tmp) {
|
||||
outBuffer[gid] = tmp;
|
||||
}
|
||||
};
|
||||
|
||||
class hipPerfSampleRate {
|
||||
public:
|
||||
hipPerfSampleRate();
|
||||
~hipPerfSampleRate();
|
||||
|
||||
void open(void);
|
||||
void run(unsigned int testCase);
|
||||
void close(void);
|
||||
|
||||
// array of funtion pointers
|
||||
typedef void (hipPerfSampleRate::*funPtr)(void * outBuffer, unsigned int
|
||||
inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks,
|
||||
int threads_per_block);
|
||||
|
||||
// Wrappers
|
||||
void float_kernel(void * outBuffer, unsigned int
|
||||
inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks,
|
||||
int threads_per_block);
|
||||
|
||||
void float2_kernel(void * outBuffer, unsigned int
|
||||
inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks,
|
||||
int threads_per_block);
|
||||
|
||||
void float4_kernel(void * outBuffer, unsigned int
|
||||
inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks,
|
||||
int threads_per_block);
|
||||
|
||||
private:
|
||||
void setData(void *ptr, unsigned int value);
|
||||
void checkData(uint *ptr);
|
||||
|
||||
unsigned int width_;
|
||||
unsigned int bufSize_;
|
||||
unsigned long long totalIters = 0;
|
||||
int numCUs;
|
||||
|
||||
unsigned int outBufSize_;
|
||||
static const unsigned int MAX_ITERATIONS = 25;
|
||||
unsigned int numBufs_;
|
||||
unsigned int typeIdx_;
|
||||
};
|
||||
|
||||
|
||||
hipPerfSampleRate::hipPerfSampleRate() {}
|
||||
|
||||
hipPerfSampleRate::~hipPerfSampleRate() {}
|
||||
|
||||
void hipPerfSampleRate::open(void) {
|
||||
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
failed("No GPU!");
|
||||
}
|
||||
|
||||
int deviceId = 0;
|
||||
hipDeviceProp_t props = {0};
|
||||
props = {0};
|
||||
HIPCHECK(hipSetDevice(deviceId));
|
||||
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 hipPerfSampleRate::close() {
|
||||
|
||||
}
|
||||
|
||||
|
||||
// Wrappers for the kernel launches
|
||||
void hipPerfSampleRate::float_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer,
|
||||
int numBufs, int grids, int blocks, int threads_per_block) {
|
||||
|
||||
hipLaunchKernelGGL(sampleRateFloat<float>, dim3(grids, grids, grids), dim3 (blocks), 0, 0,
|
||||
(float*)outBuffer, inBufSize, writeIt, (float**)inBuffer, numBufs);
|
||||
|
||||
}
|
||||
|
||||
void hipPerfSampleRate::float2_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer,
|
||||
int grids, int blocks, int threads_per_block, int numBufs) {
|
||||
|
||||
hipLaunchKernelGGL(sampleRate<float2>, dim3(grids, grids, grids), dim3(blocks), 0, 0,
|
||||
(float2 *)outBuffer, inBufSize, writeIt, (float2**)inBuffer, numBufs);
|
||||
}
|
||||
|
||||
void hipPerfSampleRate::float4_kernel(void * outBuffer, unsigned int inBufSize,
|
||||
unsigned int writeIt, void **inBuffer,
|
||||
int grids, int blocks, int threads_per_block, int numBufs) {
|
||||
|
||||
hipLaunchKernelGGL(sampleRate<float4>, dim3(grids, grids, grids), dim3(blocks), 0, 0,
|
||||
(float4 *) outBuffer, inBufSize, writeIt, (float4**)inBuffer, numBufs);
|
||||
}
|
||||
|
||||
void hipPerfSampleRate::run(unsigned int test) {
|
||||
|
||||
funPtr p[] = {&hipPerfSampleRate::float_kernel, &hipPerfSampleRate::float2_kernel,
|
||||
&hipPerfSampleRate::float4_kernel};
|
||||
|
||||
// We compute a square domain
|
||||
width_ = sizes[test % NUM_SIZES];
|
||||
typeIdx_ = (test / NUM_SIZES) % NUM_TYPES;
|
||||
bufSize_ = width_ * width_ * typeSizes[typeIdx_];
|
||||
numBufs_ = (1 << (test / (NUM_SIZES * NUM_TYPES)));
|
||||
|
||||
void * hOutPtr;
|
||||
void * dOutPtr;
|
||||
void * hInPtr[numBufs_];
|
||||
void ** dPtr;
|
||||
void * dInPtr[numBufs_];
|
||||
|
||||
outBufSize_ =
|
||||
sizes[NUM_SIZES - 1] * sizes[NUM_SIZES - 1] * typeSizes[NUM_TYPES - 1];
|
||||
|
||||
// Allocate memory on the host and device
|
||||
HIPCHECK(hipHostMalloc((void **)&hOutPtr, outBufSize_, hipHostMallocDefault));
|
||||
setData((void *)hOutPtr, 0xdeadbeef);
|
||||
HIPCHECK(hipMalloc((uint **)&dOutPtr, outBufSize_));
|
||||
|
||||
// Allocate 2D array in Device
|
||||
hipMalloc((void **)&dPtr, numBufs_* sizeof(void *));
|
||||
|
||||
for (uint i = 0; i < numBufs_; i++) {
|
||||
HIPCHECK(hipHostMalloc((void **)&hInPtr[i], bufSize_, hipHostMallocDefault));
|
||||
HIPCHECK(hipMalloc((uint **)&dInPtr[i], bufSize_));
|
||||
setData(hInPtr[i], 0x3f800000);
|
||||
}
|
||||
|
||||
// Populate array of pointers with array addresses
|
||||
hipMemcpy(dPtr, dInPtr, numBufs_* sizeof(void *), hipMemcpyHostToDevice);
|
||||
|
||||
// Copy memory from host to device
|
||||
for (uint i = 0; i < numBufs_; i++) {
|
||||
HIPCHECK(hipMemcpy(dInPtr[i], hInPtr[i], bufSize_, hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(dOutPtr, hOutPtr, outBufSize_, hipMemcpyHostToDevice));
|
||||
|
||||
// Prepare kernel launch parameters
|
||||
// outBufSize_/sizeof(uint) - Grid size in 3D
|
||||
int grids = 64;
|
||||
int blocks = 64;
|
||||
int threads_per_block = 1;
|
||||
|
||||
unsigned int maxIter = MAX_ITERATIONS * (MAX_BUFS / numBufs_);
|
||||
unsigned int sizeDW = width_ * width_;
|
||||
unsigned int writeIt = 0;
|
||||
|
||||
int idx = 0;
|
||||
|
||||
if (!types[typeIdx_].compare("float")) {
|
||||
idx = 0;
|
||||
}
|
||||
else if(!types[typeIdx_].compare("float2")) {
|
||||
idx = 1;
|
||||
}
|
||||
else if(!types[typeIdx_].compare("float4")) {
|
||||
idx = 2;
|
||||
}
|
||||
|
||||
|
||||
// Time the kernel execution
|
||||
auto all_start = std::chrono::steady_clock::now();
|
||||
for (uint i = 0; i < maxIter; i++) {
|
||||
(this->*p[idx]) ((void *)dOutPtr, sizeDW, writeIt, dPtr, numBufs_, grids, blocks,
|
||||
threads_per_block);
|
||||
}
|
||||
|
||||
hipDeviceSynchronize();
|
||||
auto all_end = std::chrono::steady_clock::now();
|
||||
std::chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
double perf = ((double)outBufSize_ * numBufs_ * (double)maxIter * (double)(1e-09)) /
|
||||
all_kernel_time.count();
|
||||
|
||||
cout << "Domain " << sizes[NUM_SIZES - 1] << "x"<< sizes[NUM_SIZES - 1] << " bufs "
|
||||
<< numBufs_ << " " << types[typeIdx_] << " " << width_<<"x"<<width_<< " (GB/s) "
|
||||
<< perf << endl;
|
||||
|
||||
HIPCHECK(hipFree(dOutPtr));
|
||||
|
||||
// Free host and device memory
|
||||
for (uint i = 0; i < numBufs_; i++) {
|
||||
HIPCHECK(hipHostFree(hInPtr[i]));
|
||||
HIPCHECK(hipFree(dInPtr[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostFree(hOutPtr));
|
||||
HIPCHECK(hipFree(dPtr));
|
||||
}
|
||||
|
||||
|
||||
void hipPerfSampleRate::setData(void *ptr, unsigned int value) {
|
||||
unsigned int *ptr2 = (unsigned int *)ptr;
|
||||
for (unsigned int i = 0; i < bufSize_ / sizeof(unsigned int); i++) {
|
||||
ptr2[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void hipPerfSampleRate::checkData(uint *ptr) {
|
||||
for (unsigned int i = 0; i < outBufSize_ / sizeof(float); i++) {
|
||||
if (ptr[i] != (float)numBufs_) {
|
||||
cout << "Data validation failed at "<< i << " Got "<< ptr[i] << ", expected "
|
||||
<< (float)numBufs_;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
hipPerfSampleRate sampleTypes;
|
||||
|
||||
sampleTypes.open();
|
||||
|
||||
for (unsigned int testCase = 0; testCase < 216 ; testCase+=36) {
|
||||
sampleTypes.run(testCase);
|
||||
}
|
||||
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -1,250 +0,0 @@
|
||||
/*
|
||||
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 <iostream>
|
||||
#include <chrono>
|
||||
#include "test_common.h"
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define sharedMemSize1 2048
|
||||
#define sharedMemSize2 256
|
||||
|
||||
__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) {
|
||||
|
||||
size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t lid = threadIdx.x;
|
||||
__shared__ float local[sharedMemSize1];
|
||||
|
||||
float val1 = 0;
|
||||
float val2 = 0;
|
||||
float val3 = 0;
|
||||
float val4 = 0;
|
||||
|
||||
for (int i = 0; i < (sharedMemSize1 / 64); i++) {
|
||||
local[lid + i * 64] = lid;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
val1 += local[lid];
|
||||
val2 += local[lid + 64];
|
||||
val3 += local[lid + 128];
|
||||
val4 += local[lid + 192];
|
||||
val1 += local[lid + 256];
|
||||
val2 += local[lid + 320];
|
||||
val3 += local[lid + 384];
|
||||
val4 += local[lid + 448];
|
||||
val1 += local[lid + 512];
|
||||
val2 += local[lid + 576];
|
||||
val3 += local[lid + 640];
|
||||
val4 += local[lid + 704];
|
||||
val1 += local[lid + 768];
|
||||
val2 += local[lid + 832];
|
||||
val3 += local[lid + 896];
|
||||
val4 += local[lid + 960];
|
||||
val1 += local[lid + 1024];
|
||||
val2 += local[lid + 1088];
|
||||
val3 += local[lid + 1152];
|
||||
val4 += local[lid + 1216];
|
||||
val1 += local[lid + 1280];
|
||||
val2 += local[lid + 1344];
|
||||
val3 += local[lid + 1408];
|
||||
val4 += local[lid + 1472];
|
||||
val1 += local[lid + 1536];
|
||||
val2 += local[lid + 1600];
|
||||
val3 += local[lid + 1664];
|
||||
val4 += local[lid + 1728];
|
||||
val1 += local[lid + 1792];
|
||||
val2 += local[lid + 1856];
|
||||
val3 += local[lid + 1920];
|
||||
val4 += local[lid + 1984];
|
||||
|
||||
if (gid < N) {
|
||||
outBuf[gid] = val1 + val2 + val3 + val4;
|
||||
}
|
||||
};
|
||||
|
||||
__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) {
|
||||
size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t lid = threadIdx.x;
|
||||
__shared__ float local[sharedMemSize2];
|
||||
|
||||
float val0 = 0.0f;
|
||||
float val1 = 0.0f;
|
||||
|
||||
for (int i = 0; i < (sharedMemSize2 / 64); i++) {
|
||||
local[lid + i * 64] = lid;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
#pragma nounroll
|
||||
for (uint i = 0; i < 32; i++) {
|
||||
val0 += local[8 * i + 0];
|
||||
val1 += local[8 * i + 1];
|
||||
val0 += local[8 * i + 2];
|
||||
val1 += local[8 * i + 3];
|
||||
val0 += local[8 * i + 4];
|
||||
val1 += local[8 * i + 5];
|
||||
val0 += local[8 * i + 6];
|
||||
val1 += local[8 * i + 7];
|
||||
}
|
||||
|
||||
if (gid < N) {
|
||||
outBuf[gid] = val0 + val1;
|
||||
}
|
||||
};
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
float *dDst;
|
||||
float *hDst;
|
||||
hipStream_t stream;
|
||||
constexpr uint numSizes = 4;
|
||||
constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216};
|
||||
uint numReads1 = 32;
|
||||
uint numReads2 = 256;
|
||||
uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float);
|
||||
uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float);
|
||||
int nIter = 1000;
|
||||
const unsigned threadsPerBlock = 64;
|
||||
|
||||
int nGpu = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&nGpu));
|
||||
if (nGpu < 1) {
|
||||
cout << "info: didn't find any GPU! skipping the test!\n";
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int device = 0;
|
||||
HIPCHECK(hipSetDevice(device));
|
||||
hipDeviceProp_t props;
|
||||
HIPCHECK(hipGetDeviceProperties(&props, device));
|
||||
cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name
|
||||
<< " with " << props.multiProcessorCount << " CUs" << endl;
|
||||
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
for (int nTest = 0; nTest < numSizes; nTest++) {
|
||||
uint nBytes = Sizes[nTest % numSizes];
|
||||
ulong N = nBytes / sizeof(float);
|
||||
const unsigned blocks = N / threadsPerBlock;
|
||||
|
||||
hDst = new float[nBytes];
|
||||
HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
memset(hDst, 0, nBytes);
|
||||
|
||||
HIPCHECK(hipMalloc(&dDst, nBytes));
|
||||
HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, dDst, N);
|
||||
HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
|
||||
hipDeviceSynchronize();
|
||||
|
||||
int tmp = 0;
|
||||
for (int i = 0; i < N; i++) {
|
||||
if (i % threadsPerBlock == 0) {
|
||||
tmp = 0;
|
||||
}
|
||||
if (hDst[i] != tmp) {
|
||||
cout << "info: Data validation failed for warm up run!" << endl;
|
||||
cout << "info: expected " << tmp << " got " << hDst[i] << endl;
|
||||
HIPCHECK (hipErrorUnknown);
|
||||
}
|
||||
tmp += threadsPerBlock / 2;
|
||||
}
|
||||
|
||||
auto all_start = chrono::steady_clock::now();
|
||||
for (int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, dDst, N);
|
||||
}
|
||||
hipDeviceSynchronize();
|
||||
|
||||
auto all_end = chrono::steady_clock::now();
|
||||
chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = ((double) blocks * threadsPerBlock
|
||||
* (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) * nIter
|
||||
* (double) (1e-09)) / all_kernel_time.count();
|
||||
|
||||
cout << "info: read speed = " << setw(8) << perf << " GB/s for "
|
||||
<< sharedMemSizeBytes1 / 1024 << " KB shared memory"
|
||||
" with " << setw(8) << blocks * threadsPerBlock << " threads, "
|
||||
<< setw(4) << numReads1 << " reads in sharedMemReadSpeed1 kernel" << endl;
|
||||
|
||||
delete[] hDst;
|
||||
hipFree(dDst);
|
||||
}
|
||||
|
||||
|
||||
for (int nTest = 0; nTest < numSizes; nTest++) {
|
||||
uint nBytes = Sizes[nTest % numSizes];
|
||||
ulong N = nBytes / sizeof(float);
|
||||
const unsigned blocks = N / threadsPerBlock;
|
||||
|
||||
hDst = new float[nBytes];
|
||||
HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
memset(hDst, 0, nBytes);
|
||||
|
||||
HIPCHECK(hipMalloc(&dDst, nBytes));
|
||||
HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, dDst, N);
|
||||
HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
|
||||
hipDeviceSynchronize();
|
||||
|
||||
auto all_start = chrono::steady_clock::now();
|
||||
for (int i = 0; i < nIter; i++) {
|
||||
hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks),
|
||||
dim3(threadsPerBlock), 0, stream, dDst, N);
|
||||
}
|
||||
hipDeviceSynchronize();
|
||||
|
||||
auto all_end = chrono::steady_clock::now();
|
||||
chrono::duration<double> all_kernel_time = all_end - all_start;
|
||||
|
||||
// read speed in GB/s
|
||||
double perf = ((double) blocks * threadsPerBlock
|
||||
* (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) * nIter
|
||||
* (double) (1e-09)) / all_kernel_time.count();
|
||||
|
||||
cout << "info: read speed = " << setw(8) << perf << " GB/s for "
|
||||
<< sharedMemSizeBytes2 / 1024 << " KB shared memory"
|
||||
" with " << setw(8) << blocks * threadsPerBlock << " threads, "
|
||||
<< setw(4) << numReads2 << " reads in sharedMemReadSpeed2 kernel" << endl;
|
||||
|
||||
delete[] hDst;
|
||||
hipFree(dDst);
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
passed();
|
||||
}
|
||||
Tagairt in Eagrán Nua
Cuir bac ar úsáideoir