diff --git a/projects/hip-tests/perftests/memory/hipPerfSampleRate.cpp b/projects/hip-tests/perftests/memory/hipPerfSampleRate.cpp index 040903e67a..15141998d8 100644 --- a/projects/hip-tests/perftests/memory/hipPerfSampleRate.cpp +++ b/projects/hip-tests/perftests/memory/hipPerfSampleRate.cpp @@ -1,308 +1,308 @@ -/* - 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 -#include -#include "test_common.h" -#include -#include - -using namespace std; - -#define NUM_TYPES 3 -vector types= {"float", "float2", "float4"}; -vector typeSizes = {4, 8, 16}; - -#define NUM_SIZES 12 -vector sizes = {1, 2, 4, 8, 16, 32, - 64, 128, 256, 512, 1024, 2048}; - -#define NUM_BUFS 6 -#define MAX_BUFS (1 << (NUM_BUFS - 1)) - -template -__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 = (T)0.0f; - for(int i = 0; i < numBufs; i++) { - tmp += *(*(inBuffer+i)+inputIdx); - } - - if (writeIt*(unsigned int)tmp.x) { - outBuffer[gid] = tmp; - } -}; - -template -__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) { - std::cout << "info: didn't find any GPU! skipping the test!\n"; - passed(); - return; - } - - 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, 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, 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, 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 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"< +#include +#include "test_common.h" +#include +#include + +using namespace std; + +#define NUM_TYPES 3 +vector types= {"float", "float2", "float4"}; +vector typeSizes = {4, 8, 16}; + +#define NUM_SIZES 12 +vector sizes = {1, 2, 4, 8, 16, 32, + 64, 128, 256, 512, 1024, 2048}; + +#define NUM_BUFS 6 +#define MAX_BUFS (1 << (NUM_BUFS - 1)) + +template +__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 = (T)0.0f; + for(int i = 0; i < numBufs; i++) { + tmp += *(*(inBuffer+i)+inputIdx); + } + + if (writeIt*(unsigned int)tmp.x) { + outBuffer[gid] = tmp; + } +}; + +template +__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) { + std::cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return; + } + + 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, 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, 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, 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 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"< -#include -#include "test_common.h" - -typedef struct { - double x; - double y; - double width; -} coordRec; - -static coordRec coords[] = { - {0.0, 0.0, 0.00001}, // All black -}; - -static unsigned int numCoords = sizeof(coords) / sizeof(coordRec); - -__global__ void mandelbrot(uint *out, uint width, float xPos, float yPos, float xStep, - float yStep, uint maxIter) { - - int tid = (blockIdx.x * blockDim.x + threadIdx.x); - int i = tid % width; - int j = tid / width; - float x0 = (float)(xPos + xStep*i); - float y0 = (float)(yPos + yStep*j); - - float x = x0; - float y = y0; - - uint iter = 0; - float tmp; - for (iter = 0; (x*x + y*y <= 4.0f) && (iter < maxIter); iter++) { - tmp = x; - x = fma(-y,y,fma(x,x,x0)); - y = fma(2.0f*tmp,y,y0); - } - - out[tid] = iter; -}; - -class hipPerfDeviceConcurrency { - public: - hipPerfDeviceConcurrency(); - ~hipPerfDeviceConcurrency(); - - void setNumGpus(unsigned int num) { - numDevices = num; - } - unsigned int getNumGpus() { - return numDevices; - } - - void open(void); - void close(void); - void run(unsigned int testCase, int numGpus); - - private: - void setData(void *ptr, unsigned int value); - void checkData(uint *ptr); - - unsigned int numDevices; - unsigned int width_; - unsigned int bufSize; - unsigned int coordIdx; - unsigned long long totalIters = 0; -}; - - -hipPerfDeviceConcurrency::hipPerfDeviceConcurrency() {} - -hipPerfDeviceConcurrency::~hipPerfDeviceConcurrency() {} - -void hipPerfDeviceConcurrency::open(void) { - - - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - setNumGpus(nGpu); - if (nGpu < 1) { - std::cout << "info: didn't find any GPU! skipping the test!\n"; - passed(); - } - - -} - - -void hipPerfDeviceConcurrency::close() { -} - -void hipPerfDeviceConcurrency::run(unsigned int testCase, int numGpus) { - - - static int deviceId; - uint * hPtr[numGpus]; - uint * dPtr[numGpus]; - hipStream_t streams[numGpus]; - int numCUs[numGpus]; - unsigned int maxIter[numGpus]; - unsigned long long expectedIters[numGpus]; - - int threads, threads_per_block, blocks; - float xStep, yStep, xPos, yPos; - - for(int i = 0; i < numGpus; i++) { - - if(testCase != 0) { - deviceId = i; - } - - HIPCHECK(hipSetDevice(deviceId)); - - hipDeviceProp_t props = {0}; - HIPCHECK(hipGetDeviceProperties(&props, i)); - - if (testCase != 0) { - std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs" << " and device ID: " - << i << std::endl; - } - - numCUs[i] = props.multiProcessorCount; - int clkFrequency = 0; - HIPCHECK(hipDeviceGetAttribute(&clkFrequency, hipDeviceAttributeClockRate, i)); - - clkFrequency =(unsigned int)clkFrequency/1000; - - // Maximum iteration count - // maxIter = 8388608 * (engine_clock / 1000).serial execution - maxIter[i] = (unsigned int)(((8388608 * ((float)clkFrequency / 1000)) * numCUs[i]) / 128); - maxIter[i] = (maxIter[i] + 15) & ~15; - - // Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once. - width_ = 256; - - bufSize = width_ * width_ * sizeof(uint); - - // Create streams for concurrency - HIPCHECK(hipStreamCreate(&streams[i])); - - // Allocate memory on the host and device - HIPCHECK(hipHostMalloc((void **)&hPtr[i], bufSize, hipHostMallocDefault)); - setData(hPtr[i], 0xdeadbeef); - HIPCHECK(hipMalloc((uint **)&dPtr[i], bufSize)) - - // Prepare kernel launch parameters - threads = (bufSize/sizeof(uint)); - threads_per_block = 64; - blocks = (threads/threads_per_block) + (threads % threads_per_block); - - coordIdx = testCase % numCoords; - xStep = (float)(coords[coordIdx].width / (double)width_); - yStep = (float)(-coords[coordIdx].width / (double)width_); - xPos = (float)(coords[coordIdx].x - 0.5 * coords[coordIdx].width); - yPos = (float)(coords[coordIdx].y + 0.5 * coords[coordIdx].width); - - // Copy memory from host to device - HIPCHECK(hipMemcpy(dPtr[i], hPtr[i], bufSize, hipMemcpyHostToDevice)); - - } - - // Time the kernel execution - auto all_start = std::chrono::steady_clock::now(); - - for(int i = 0; i < numGpus; i++) { - - if(testCase != 0) { - deviceId = i; - } - - HIPCHECK(hipSetDevice(deviceId)); - - hipLaunchKernelGGL(mandelbrot, dim3(blocks), dim3(threads_per_block), 0, streams[i], - dPtr[i], width_, xPos, yPos, xStep, yStep, maxIter[i]); - - } - - for(int i = 0; i < numGpus; i++) { - HIPCHECK(hipStreamSynchronize(0)); - } - - - auto all_end = std::chrono::steady_clock::now(); - std::chrono::duration all_kernel_time = all_end - all_start; - - for(int i = 0; i < numGpus; i++) { - - if(testCase != 0) { - deviceId = i; - } - HIPCHECK(hipSetDevice(deviceId)); - - // Copy data back from device to the host - HIPCHECK(hipMemcpy(hPtr[i], dPtr[i], bufSize, hipMemcpyDeviceToHost)); - - checkData(hPtr[i]); - expectedIters[i] = width_ * width_ * (unsigned long long) maxIter[i]; - - if (testCase != 0) { - checkData(hPtr[i]); - if(totalIters != expectedIters[i]) { - std::cout << "Incorrect iteration count detected" << std::endl; - } - } - - - HIPCHECK(hipStreamDestroy(streams[i])); - - // Free host and device memory - HIPCHECK(hipFree(hPtr[i])); - HIPCHECK(hipFree(dPtr[i])); - } - - if (testCase != 0) { - std::cout << '\n' << "Measured time for kernel computation on " << numGpus << " device (s): " - << all_kernel_time.count() << " (s) " << '\n' << std::endl; - } - - if(testCase == 0) { - deviceId++; - } - - -} - - -void hipPerfDeviceConcurrency::setData(void *ptr, unsigned int value) { - unsigned int *ptr2 = (unsigned int *)ptr; - for (unsigned int i = 0; i < width_ * width_ ; i++) { - ptr2[i] = value; - } -} - - -void hipPerfDeviceConcurrency::checkData(uint *ptr) { - totalIters = 0; - for (unsigned int i = 0; i < width_ * width_; i++) { - totalIters += ptr[i]; - } -} - - -int main(int argc, char* argv[]) { - hipPerfDeviceConcurrency deviceConcurrency; - - deviceConcurrency.open(); - - int nGpu = deviceConcurrency.getNumGpus(); - - // testCase = 0 refers to warmup kernel run - int testCase = 0; - - for (int i = 0; i < nGpu; i++) { - // Warm-up kernel on all devices - deviceConcurrency.run(testCase, 1); - } - - // Time for kernel on 1 device - deviceConcurrency.run(++testCase, 1); - - // Time for kernel on all available devices - deviceConcurrency.run(++testCase, nGpu); - - passed(); -} +/* + 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 +#include +#include "test_common.h" + +typedef struct { + double x; + double y; + double width; +} coordRec; + +static coordRec coords[] = { + {0.0, 0.0, 0.00001}, // All black +}; + +static unsigned int numCoords = sizeof(coords) / sizeof(coordRec); + +__global__ void mandelbrot(uint *out, uint width, float xPos, float yPos, float xStep, + float yStep, uint maxIter) { + + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % width; + int j = tid / width; + float x0 = (float)(xPos + xStep*i); + float y0 = (float)(yPos + yStep*j); + + float x = x0; + float y = y0; + + uint iter = 0; + float tmp; + for (iter = 0; (x*x + y*y <= 4.0f) && (iter < maxIter); iter++) { + tmp = x; + x = fma(-y,y,fma(x,x,x0)); + y = fma(2.0f*tmp,y,y0); + } + + out[tid] = iter; +}; + +class hipPerfDeviceConcurrency { + public: + hipPerfDeviceConcurrency(); + ~hipPerfDeviceConcurrency(); + + void setNumGpus(unsigned int num) { + numDevices = num; + } + unsigned int getNumGpus() { + return numDevices; + } + + void open(void); + void close(void); + void run(unsigned int testCase, int numGpus); + + private: + void setData(void *ptr, unsigned int value); + void checkData(uint *ptr); + + unsigned int numDevices; + unsigned int width_; + unsigned int bufSize; + unsigned int coordIdx; + unsigned long long totalIters = 0; +}; + + +hipPerfDeviceConcurrency::hipPerfDeviceConcurrency() {} + +hipPerfDeviceConcurrency::~hipPerfDeviceConcurrency() {} + +void hipPerfDeviceConcurrency::open(void) { + + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + setNumGpus(nGpu); + if (nGpu < 1) { + std::cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + } + + +} + + +void hipPerfDeviceConcurrency::close() { +} + +void hipPerfDeviceConcurrency::run(unsigned int testCase, int numGpus) { + + + static int deviceId; + uint * hPtr[numGpus]; + uint * dPtr[numGpus]; + hipStream_t streams[numGpus]; + int numCUs[numGpus]; + unsigned int maxIter[numGpus]; + unsigned long long expectedIters[numGpus]; + + int threads, threads_per_block, blocks; + float xStep, yStep, xPos, yPos; + + for(int i = 0; i < numGpus; i++) { + + if(testCase != 0) { + deviceId = i; + } + + HIPCHECK(hipSetDevice(deviceId)); + + hipDeviceProp_t props = {0}; + HIPCHECK(hipGetDeviceProperties(&props, i)); + + if (testCase != 0) { + std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs" << " and device ID: " + << i << std::endl; + } + + numCUs[i] = props.multiProcessorCount; + int clkFrequency = 0; + HIPCHECK(hipDeviceGetAttribute(&clkFrequency, hipDeviceAttributeClockRate, i)); + + clkFrequency =(unsigned int)clkFrequency/1000; + + // Maximum iteration count + // maxIter = 8388608 * (engine_clock / 1000).serial execution + maxIter[i] = (unsigned int)(((8388608 * ((float)clkFrequency / 1000)) * numCUs[i]) / 128); + maxIter[i] = (maxIter[i] + 15) & ~15; + + // Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once. + width_ = 256; + + bufSize = width_ * width_ * sizeof(uint); + + // Create streams for concurrency + HIPCHECK(hipStreamCreate(&streams[i])); + + // Allocate memory on the host and device + HIPCHECK(hipHostMalloc((void **)&hPtr[i], bufSize, hipHostMallocDefault)); + setData(hPtr[i], 0xdeadbeef); + HIPCHECK(hipMalloc((uint **)&dPtr[i], bufSize)) + + // Prepare kernel launch parameters + threads = (bufSize/sizeof(uint)); + threads_per_block = 64; + blocks = (threads/threads_per_block) + (threads % threads_per_block); + + coordIdx = testCase % numCoords; + xStep = (float)(coords[coordIdx].width / (double)width_); + yStep = (float)(-coords[coordIdx].width / (double)width_); + xPos = (float)(coords[coordIdx].x - 0.5 * coords[coordIdx].width); + yPos = (float)(coords[coordIdx].y + 0.5 * coords[coordIdx].width); + + // Copy memory from host to device + HIPCHECK(hipMemcpy(dPtr[i], hPtr[i], bufSize, hipMemcpyHostToDevice)); + + } + + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + + for(int i = 0; i < numGpus; i++) { + + if(testCase != 0) { + deviceId = i; + } + + HIPCHECK(hipSetDevice(deviceId)); + + hipLaunchKernelGGL(mandelbrot, dim3(blocks), dim3(threads_per_block), 0, streams[i], + dPtr[i], width_, xPos, yPos, xStep, yStep, maxIter[i]); + + } + + for(int i = 0; i < numGpus; i++) { + HIPCHECK(hipStreamSynchronize(0)); + } + + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + for(int i = 0; i < numGpus; i++) { + + if(testCase != 0) { + deviceId = i; + } + HIPCHECK(hipSetDevice(deviceId)); + + // Copy data back from device to the host + HIPCHECK(hipMemcpy(hPtr[i], dPtr[i], bufSize, hipMemcpyDeviceToHost)); + + checkData(hPtr[i]); + expectedIters[i] = width_ * width_ * (unsigned long long) maxIter[i]; + + if (testCase != 0) { + checkData(hPtr[i]); + if(totalIters != expectedIters[i]) { + std::cout << "Incorrect iteration count detected" << std::endl; + } + } + + + HIPCHECK(hipStreamDestroy(streams[i])); + + // Free host and device memory + HIPCHECK(hipFree(hPtr[i])); + HIPCHECK(hipFree(dPtr[i])); + } + + if (testCase != 0) { + std::cout << '\n' << "Measured time for kernel computation on " << numGpus << " device (s): " + << all_kernel_time.count() << " (s) " << '\n' << std::endl; + } + + if(testCase == 0) { + deviceId++; + } + + +} + + +void hipPerfDeviceConcurrency::setData(void *ptr, unsigned int value) { + unsigned int *ptr2 = (unsigned int *)ptr; + for (unsigned int i = 0; i < width_ * width_ ; i++) { + ptr2[i] = value; + } +} + + +void hipPerfDeviceConcurrency::checkData(uint *ptr) { + totalIters = 0; + for (unsigned int i = 0; i < width_ * width_; i++) { + totalIters += ptr[i]; + } +} + + +int main(int argc, char* argv[]) { + hipPerfDeviceConcurrency deviceConcurrency; + + deviceConcurrency.open(); + + int nGpu = deviceConcurrency.getNumGpus(); + + // testCase = 0 refers to warmup kernel run + int testCase = 0; + + for (int i = 0; i < nGpu; i++) { + // Warm-up kernel on all devices + deviceConcurrency.run(testCase, 1); + } + + // Time for kernel on 1 device + deviceConcurrency.run(++testCase, 1); + + // Time for kernel on all available devices + deviceConcurrency.run(++testCase, nGpu); + + passed(); +}