AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Merge 'amd-master-next' into 'amd-npi-next'

Sfoglia il codice sorgente 
Change-Id: I4d7ded0a1316a469f0880490e563c3e3f00cc970


[ROCm/hip commit: d01cbc7220]
This commit is contained in:
Jenkins
2020-08-20 21:09:41 +00:00
parent da9a722013 fc61d76630
commit cd599ef1e0
6 ha cambiato i file con 2123 aggiunte e 79 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
projects/hip/tests/performance/stream/hipPerfStreamConcurrency.cpp
+446
Vedi File
@@ -0,0 +1,446 @@
/*
Copyright (c) 2015-2020 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 nvcc
* TEST: %t
* HIT_END
*/
#include <iostream>
#include <chrono>
#include "test_common.h"
#include <hip/hip_vector_types.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/4);
int j = tid / (width/4);
int4 veci = make_int4(4*i, 4*i+1, 4*i+2, 4*i+3);
int4 vecj = make_int4(j, j, j, j);
float4 x0;
x0.data[0] = (float)(xPos + xStep*veci.data[0]);
x0.data[1] = (float)(xPos + xStep*veci.data[1]);
x0.data[2] = (float)(xPos + xStep*veci.data[2]);
x0.data[3] = (float)(xPos + xStep*veci.data[3]);
float4 y0;
y0.data[0] = (float)(yPos + yStep*vecj.data[0]);
y0.data[1] = (float)(yPos + yStep*vecj.data[1]);
y0.data[2] = (float)(yPos + yStep*vecj.data[2]);
y0.data[3] = (float)(yPos + yStep*vecj.data[3]);
float4 x = x0;
float4 y = y0;
uint iter = 0;
float4 tmp;
int4 stay;
int4 ccount = make_int4(0, 0, 0, 0);
float4 savx = x;
float4 savy = y;
stay.data[0] = (x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= (float)(4.0f);
stay.data[1] = (x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= (float)(4.0f);
stay.data[2] = (x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= (float)(4.0f);
stay.data[3] = (x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= (float)(4.0f);
for (iter = 0; (stay.data[0] | stay.data[1] | stay.data[2] | stay.data[3]) && (iter < maxIter);
iter+=16) {
x = savx;
y = savy;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
// Two iterations
tmp = x*x + x0 - y*y;
y = 2.0f * x * y + y0;
x = tmp*tmp + x0 - y*y;
y = 2.0f * tmp * y + y0;
stay.data[0] = (x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= (float)(4.0f);
stay.data[1] = (x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= (float)(4.0f);
stay.data[2] = (x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= (float)(4.0f);
stay.data[3] = (x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= (float)(4.0f);
savx.data[0] = (bool)(stay.data[0] ? x.data[0] : savx.data[0]);
savx.data[1] = (bool)(stay.data[1] ? x.data[1] : savx.data[1]);
savx.data[2] = (bool)(stay.data[2] ? x.data[2] : savx.data[2]);
savx.data[3] = (bool)(stay.data[3] ? x.data[3] : savx.data[3]);
savy.data[0] = (bool)(stay.data[0] ? y.data[0] : savy.data[0]);
savy.data[1] = (bool)(stay.data[1] ? y.data[1] : savy.data[1]);
savy.data[2] = (bool)(stay.data[2] ? y.data[2] : savy.data[2]);
savy.data[3] = (bool)(stay.data[3] ? y.data[3] : savy.data[3]);
ccount.data[0] -= stay.data[0]*16;
ccount.data[1] -= stay.data[1]*16;
ccount.data[2] -= stay.data[2]*16;
ccount.data[3] -= stay.data[3]*16;
}
// Handle remainder
if (!(stay.data[0] & stay.data[1] & stay.data[2] & stay.data[3]))
{
iter = 16;
do
{
x = savx;
y = savy;
stay.x = ((x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= 4.0f) && (ccount.data[0] < maxIter);
stay.y = ((x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= 4.0f) && (ccount.data[1] < maxIter);
stay.z = ((x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= 4.0f) && (ccount.data[2] < maxIter);
stay.w = ((x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= 4.0f) && (ccount.data[3] < maxIter);
tmp = x;
x = x*x + x0 - y*y;
y = 2.0f*tmp*y + y0;
ccount.data[0] += stay.data[0];
ccount.data[1] += stay.data[1];
ccount.data[2] += stay.data[2];
ccount.data[3] += stay.data[3];
iter--;
savx.data[0] = (stay.data[0] ? x.data[0] : savx.data[0]);
savx.data[1] = (stay.data[1] ? x.data[1] : savx.data[1]);
savx.data[2] = (stay.data[2] ? x.data[2] : savx.data[2]);
savx.data[3] = (stay.data[3] ? x.data[3] : savx.data[3]);
savy.data[0] = (stay.data[0] ? y.data[0] : savy.data[0]);
savy.data[1] = (stay.data[1] ? y.data[1] : savy.data[1]);
savy.data[2] = (stay.data[2] ? y.data[2] : savy.data[2]);
savy.data[3] = (stay.data[3] ? y.data[3] : savy.data[3]);
} while ((stay.data[0] | stay.data[1] | stay.data[2] | stay.data[3]) && iter);
}
uint4 *vecOut = (uint4 *)out;
vecOut[tid].data[0] = (uint)(ccount.data[0]);
vecOut[tid].data[1] = (uint)(ccount.data[1]);
vecOut[tid].data[2] = (uint)(ccount.data[2]);
vecOut[tid].data[3] = (uint)(ccount.data[3]);
}
class hipPerfStreamConcurrency {
public:
hipPerfStreamConcurrency();
~hipPerfStreamConcurrency();
void setNumKernels(unsigned int num) {
numKernels = num;
}
void setNumStreams(unsigned int num) {
numStreams = num;
}
unsigned int getNumStreams() {
return numStreams;
}
unsigned int getNumKernels() {
return numKernels;
}
void open(int deviceID);
void run(unsigned int testCase, unsigned int deviceId);
void close(void);
private:
void setData(void *ptr, unsigned int value);
void checkData(uint *ptr);
unsigned int numKernels;
unsigned int numStreams;
unsigned int width_;
unsigned int bufSize;
unsigned int maxIter;
unsigned int coordIdx;
unsigned long long totalIters;
int numCUs;
};
hipPerfStreamConcurrency::hipPerfStreamConcurrency() {}
hipPerfStreamConcurrency::~hipPerfStreamConcurrency() {}
void hipPerfStreamConcurrency::open(int deviceId) {
int nGpu = 0;
HIPCHECK(hipGetDeviceCount(&nGpu));
if (nGpu < 1) {
std::cout << "info: didn't find any GPU! skipping the test!\n";
passed();
return;
}
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;
numCUs = props.multiProcessorCount;
}
void hipPerfStreamConcurrency::close() {
}
void hipPerfStreamConcurrency::run(unsigned int testCase,unsigned int deviceId) {
int clkFrequency = 0;
unsigned int numStreams = getNumStreams();
unsigned int numKernels = getNumKernels();
HIPCHECK(hipDeviceGetAttribute(&clkFrequency, hipDeviceAttributeClockRate, deviceId));
clkFrequency =(unsigned int)clkFrequency/1000;
// Maximum iteration count
// maxIter = 8388608 * (engine_clock / 1000).serial execution
maxIter = (unsigned int)(((8388608 * ((float)clkFrequency / 1000)) * numCUs) / 128);
maxIter = (maxIter + 15) & ~15;
hipStream_t streams[numStreams];
uint * hPtr[numKernels];
uint * dPtr[numKernels];
// Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once.
width_ = 256;
bufSize = width_ * sizeof(uint);
// Create streams for concurrency
for (uint i = 0; i < numStreams; i++) {
HIPCHECK(hipStreamCreate(&streams[i]));
}
// Allocate memory on the host and device
for (uint i = 0; i < numKernels; i++) {
HIPCHECK(hipHostMalloc((void **)&hPtr[i], bufSize, hipHostMallocDefault));
setData(hPtr[i], 0xdeadbeef);
HIPCHECK(hipMalloc((uint **)&dPtr[i], bufSize))
}
// Prepare kernel launch parameters
int threads = (bufSize/sizeof(uint));
int threads_per_block = 64;
int blocks = (threads/threads_per_block) + (threads % threads_per_block);
coordIdx = testCase % numCoords;
float xStep = (float)(coords[coordIdx].width / (double)width_);
float yStep = (float)(-coords[coordIdx].width / (double)width_);
float xPos = (float)(coords[coordIdx].x - 0.5 * coords[coordIdx].width);
float yPos = (float)(coords[coordIdx].y + 0.5 * coords[coordIdx].width);
// Copy memory asynchronously and concurrently from host to device
for (uint i = 0; i < numKernels; i++) {
HIPCHECK(hipMemcpyHtoDAsync(dPtr[i], hPtr[i], bufSize, streams[i % numStreams]));
}
// Synchronize to make sure all the copies are completed
for(uint i = 0; i < numStreams; i++) {
HIPCHECK(hipStreamSynchronize(streams[i]));
}
// Warm-up kernel with lower iteration
if (testCase == 0) {
maxIter = 256;
}
// Time the kernel execution
auto all_start = std::chrono::steady_clock::now();
for (uint i = 0; i < numKernels; i++) {
hipLaunchKernelGGL(mandelbrot, dim3(blocks), dim3(threads_per_block), 0, streams[i%numStreams],
dPtr[i], width_, xPos, yPos, xStep, yStep, maxIter);
}
// Synchronize all the concurrent streans to have completed execution
for(uint i = 0; i < numStreams; i++) {
HIPCHECK(hipStreamSynchronize(streams[i]));
}
auto all_end = std::chrono::steady_clock::now();
std::chrono::duration<double> all_kernel_time = all_end - all_start;
// Copy data back from device to the host
for(uint i = 0; i < numKernels; i++) {
HIPCHECK(hipMemcpyDtoHAsync(hPtr[i] ,dPtr[i], bufSize, streams[i % numStreams]));
}
if (testCase != 0) {
std::cout <<"Measured time for " << numKernels <<" kernels (s) on " << numStreams <<" stream (s): "
<< all_kernel_time.count() << std::endl;
}
unsigned long long expected =
(unsigned long long)width_ * (unsigned long long)maxIter;
for(uint i = 0 ; i < numStreams; i++) {
HIPCHECK(hipStreamDestroy(streams[i]));
}
// Free host and device memory
for (uint i = 0; i < numKernels; i++) {
HIPCHECK(hipFree(hPtr[i]));
HIPCHECK(hipFree(dPtr[i]));
}
}
void hipPerfStreamConcurrency::setData(void *ptr, unsigned int value) {
unsigned int *ptr2 = (unsigned int *)ptr;
for (unsigned int i = 0; i < width_ ; i++) {
ptr2[i] = value;
}
}
void hipPerfStreamConcurrency::checkData(uint *ptr) {
totalIters = 0;
for (unsigned int i = 0; i < width_; i++) {
totalIters += ptr[i];
}
}
int main(int argc, char* argv[]) {
hipPerfStreamConcurrency streamConcurrency;
int deviceId = 0;
streamConcurrency.open(deviceId);
for (unsigned int testCase = 0; testCase < 5; testCase++) {
switch (testCase) {
case 0:
// Warm-up kernel
streamConcurrency.setNumStreams(1);
streamConcurrency.setNumKernels(1);
break;
case 1:
// default stream executes serially
streamConcurrency.setNumStreams(1);
streamConcurrency.setNumKernels(1);
break;
case 2:
// 2-way concurrency
streamConcurrency.setNumStreams(2);
streamConcurrency.setNumKernels(2);
break;
case 3:
// 4-way concurrency
streamConcurrency.setNumStreams(4);
streamConcurrency.setNumKernels(4);
break;
case 4:
streamConcurrency.setNumStreams(2);
streamConcurrency.setNumKernels(4);
break;
case 5:
break;
default:
break;
}
streamConcurrency.run(testCase, deviceId);
}
passed();
}
projects/hip/tests/performance/stream/hipPerfStreamCreateCopyDestroy.cpp
+133
Vedi File
@@ -0,0 +1,133 @@
/*
Copyright (c) 2015-2020 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 nvcc
* TEST: %t
* HIT_END
*/
#include <iostream>
#include <chrono>
#include "test_common.h"
using namespace std;
#define BufSize 0x1000
#define Iterations 0x100
#define TotalStreams 4
#define TotalBufs 4
class hipPerfStreamCreateCopyDestroy {
private:
unsigned int numBuffers_;
unsigned int numStreams_;
const size_t totalStreams_[TotalStreams];
const size_t totalBuffers_[TotalBufs];
public:
hipPerfStreamCreateCopyDestroy() : numBuffers_(0), numStreams_(0),
totalStreams_{1, 2, 4, 8},
totalBuffers_{1, 100, 1000, 5000} {};
~hipPerfStreamCreateCopyDestroy() {};
void open(int deviceID);
void run(unsigned int testNumber);
};
void hipPerfStreamCreateCopyDestroy::open(int deviceId) {
int nGpu = 0;
HIPCHECK(hipGetDeviceCount(&nGpu));
if (nGpu < 1) {
std::cout << "info: didn't find any GPU! skipping the test!\n";
passed();
return;
}
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 hipPerfStreamCreateCopyDestroy::run(unsigned int testNumber) {
numStreams_ = totalStreams_[testNumber % TotalStreams];
size_t iter = Iterations / (numStreams_ * ((size_t)1 << (testNumber / TotalBufs + 1)));
hipStream_t streams[numStreams_];
numBuffers_ = totalBuffers_[testNumber / TotalBufs];
float* dSrc[numBuffers_];
size_t nBytes = BufSize * sizeof(float);
for (size_t b = 0; b < numBuffers_; ++b) {
HIPCHECK(hipMalloc(&dSrc[b], nBytes));
}
float* hSrc;
hSrc = new float[nBytes];
HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess);
for (size_t i = 0; i < BufSize; i++) {
hSrc[i] = 1.618f + i;
}
auto start = std::chrono::steady_clock::now();
for (size_t i = 0; i < iter; ++i) {
for (size_t s = 0; s < numStreams_; ++s) {
HIPCHECK(hipStreamCreate(&streams[s]));
}
for (size_t s = 0; s < numStreams_; ++s) {
for (size_t b = 0; b < numBuffers_; ++b) {
HIPCHECK(hipMemcpyWithStream(dSrc[b], hSrc, nBytes, hipMemcpyHostToDevice, streams[s]));
}
}
for (size_t s = 0; s < numStreams_; ++s) {
HIPCHECK(hipStreamDestroy(streams[s]));
}
}
auto end = std::chrono::steady_clock::now();
std::chrono::duration<double> diff = end - start;
auto time = static_cast<float>(diff.count() * 1000 / (iter * numStreams_));
cout << "Create+Copy+Destroy time for " << numStreams_ << " streams and "
<< setw(4) << numBuffers_ << " buffers " << " and " << setw(4)
<< iter << " iterations " << time << " (ms) " << endl;
delete [] hSrc;
for (size_t b = 0; b < numBuffers_; ++b) {
HIPCHECK(hipFree(dSrc[b]));
}
}
int main(int argc, char* argv[]) {
hipPerfStreamCreateCopyDestroy streamCCD;
int deviceId = 0;
streamCCD.open(deviceId);
for (auto testCase = 0; testCase < TotalStreams * TotalBufs; testCase++) {
streamCCD.run(testCase);
}
passed();
}
projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
-43
Vedi File
@@ -1,43 +0,0 @@
/*
Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
* Conformance test for checking functionality of
* hipError_t hipGetDevice(int *device);
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
int main() {
int numDevices = 0;
int device;
HIPCHECK(hipGetDeviceCount(&numDevices));
for (int i = 0; i < numDevices; i++) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipGetDevice(&device));
HIPASSERT(device == i);
}
passed();
}
projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
-36
Vedi File
@@ -1,36 +0,0 @@
/*
Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* TEST: %t EXCLUDE_HIP_PLATFORM
* HIT_END
*/
#include "test_common.h"
int main() {
int numDevices = 0;
HIPCHECK(hipGetDeviceCount(&numDevices));
for (int i = 0; i < numDevices; i++) {
HIPCHECK(hipSetDevice(i));
}
HIPASSERT(hipErrorInvalidDevice == hipSetDevice(numDevices));
passed();
}
projects/hip/tests/src/runtimeApi/device/hipSetGetDevice.cpp
+356
Vedi File
@@ -0,0 +1,356 @@
/*
* Copyright (c) 2020-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.
*/
/*
* Test designed to run on Linux based platforms
* Verifies functionality of
* -- hipSetDevice and hipGetDevice with different ROCR_VISIBLE_DEVICES and
* HIP_VISIBLE_DEVICES values set
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST_NAMED: %t hipSetGetDevice-invalidDevice
* TEST_NAMED: %t hipSetGetDevice-allValidDevice
* TEST_NAMED: %t hipSetGetDevice-validDev1 --computeDevCnt 1
* TEST_NAMED: %t hipSetGetDevice-validDev2 --computeDevCnt 2
* TEST_NAMED: %t hipSetGetDevice-validDev3 --computeDevCnt 3
* TEST_NAMED: %t hipSetGetDevice-validDev4 --computeDevCnt 4
* TEST_NAMED: %t hipSetGetDevice-validDev5 --computeDevCnt 5
* TEST_NAMED: %t hipSetGetDevice-validDev6 --computeDevCnt 6
* TEST_NAMED: %t hipSetGetDevice-validDev7 --computeDevCnt 7
* TEST_NAMED: %t hipSetGetDevice-validDev8 --computeDevCnt 8
* HIT_END
*/
#include "test_common.h"
#include <sys/wait.h>
#include <unistd.h>
void getDeviceCount(int *numDevices) {
int fd[2], val = 0;
pipe(fd);
pid_t childPid;
childPid = fork();
if (childPid > 0) { // parent
close(fd[1]);
read(fd[0], &val, sizeof(val));
close(fd[0]);
*numDevices = val;
} else if (childPid == 0) { // child
int devCnt = 0;
close(fd[0]);
unsetenv("ROCR_VISIBLE_DEVICES");
unsetenv("HIP_VISIBLE_DEVICES");
hipGetDeviceCount(&devCnt);
write(fd[1], &devCnt, sizeof(devCnt));
close(fd[1]);
exit(0);
} else {
failed("fork() failed. Exiting the test\n");
}
}
#define MAX_SIZE 30
// Pass either -1 in deviceNumber or invalid device number
bool testInvalidDevice(int numDevices, bool useRocrEnv, int deviceNumber) {
bool testResult = true;
int device;
int tempCount = 0;
int setDeviceErrorCheck = 0;
int getDeviceErrorCheck = 0;
int getDeviceCountErrorCheck = 0;
int fd[2];
pipe(fd);
pid_t cPid;
cPid = fork();
char visibleDeviceString[MAX_SIZE] = {};
snprintf(visibleDeviceString, MAX_SIZE, "%d", deviceNumber);
if (cPid == 0) { // child
hipError_t err;
if (true == useRocrEnv) {
setenv("ROCR_VISIBLE_DEVICES", visibleDeviceString, 1);
} else {
setenv("HIP_VISIBLE_DEVICES", visibleDeviceString, 1);
}
err = hipGetDeviceCount(&tempCount);
if (err == hipErrorNoDevice) {
getDeviceCountErrorCheck = 1;
} else {
printf("hipGetDeviceCount returns wrong value: %u\n", hipError_t(err));
}
for (int i = 0; i < numDevices; i++) {
err = hipSetDevice(i);
if (err == hipErrorInvalidDevice) {
setDeviceErrorCheck+= 1;
} else {
printf("hipSetDevice returns wrong value: %u\n", hipError_t(err));
}
err = hipGetDevice(&device);
if (err == hipErrorNoDevice) {
getDeviceErrorCheck+= 1;
} else {
printf("hipGetDevice returns wrong value: %u\n", hipError_t(err));
}
}
if ((getDeviceCountErrorCheck == 1) && (setDeviceErrorCheck == numDevices)
&& (getDeviceErrorCheck == numDevices)) {
testResult = true;
} else {
printf("Test failed for invalid device\n");
testResult = false;
}
close(fd[0]);
write(fd[1], &testResult, sizeof(testResult));
close(fd[1]);
exit(0);
} else if (cPid > 0) { // parent
close(fd[1]);
read(fd[0], &testResult, sizeof(testResult));
close(fd[0]);
wait(NULL);
} else {
printf("fork() failed\n");
testResult = false;
}
return testResult;
}
int deviceListLength = 1;
int parseExtraArguments(int argc, char* argv[]) {
int i = 0;
for (i = 1; i < argc; i++) {
const char* arg = argv[i];
if (!strcmp(arg, " ")) {
// skip NULL args.
} else if (!strcmp(arg, "--computeDevCnt")) {
if (++i >= argc || !HipTest::parseInt(argv[i], &deviceListLength)) {
failed("Bad deviceListLength argument");
}
} else {
failed("Bad argument");
}
}
return i;
}
bool testValidDevices(int numDevices, bool useRocrEnv, int *deviceList,
int deviceListLength) {
bool testResult = true;
int tempCount = 0;
int device;
int setDeviceErrorCheck = 0;
int getDeviceErrorCheck = 0;
int getDeviceCountErrorCheck = 0;
int *deviceListPtr = deviceList;
char visibleDeviceString[MAX_SIZE] = {};
if ((NULL == deviceList) || ((deviceListLength < 1) ||
deviceListLength > numDevices)) {
printf("Invalid argument for number of devices. Skipping current test\n");
return testResult;
}
for (int i = 0; i < deviceListLength; i++) {
if (NULL == deviceListPtr) {
printf("Invalid gpu index. Skipping current test\n");
return testResult;
}
snprintf(visibleDeviceString + strlen(visibleDeviceString), MAX_SIZE, "%d,",
*deviceListPtr++);
}
visibleDeviceString[strlen(visibleDeviceString)-1] = 0;
int fd[2];
pipe(fd);
pid_t cPid;
cPid = fork();
if (cPid == 0) {
unsetenv("ROCR_VISIBLE_DEVICES");
unsetenv("HIP_VISIBLE_DEVICES");
if (true == useRocrEnv) {
setenv("ROCR_VISIBLE_DEVICES", visibleDeviceString, 1);
} else {
setenv("HIP_VISIBLE_DEVICES", visibleDeviceString, 1);
}
hipError_t err;
err = hipGetDeviceCount(&tempCount);
if (tempCount == deviceListLength) {
getDeviceCountErrorCheck = 1;
} else {
printf("hipGetDeviceCount failed. return value: %u\n", hipError_t(err));
}
for (int i = 0; i < numDevices; i++) {
err = hipSetDevice(i);
if (err == hipErrorInvalidDevice) {
setDeviceErrorCheck+= 1;
}
err = hipGetDevice(&device);
if (err == hipErrorNoDevice) {
getDeviceErrorCheck+= 1;
}
}
if ((getDeviceCountErrorCheck == 1) && (setDeviceErrorCheck ==
(numDevices-deviceListLength)) && (getDeviceErrorCheck == 0)) {
testResult = true;
} else {
printf("Test failed for device count %d\n", deviceListLength);
testResult = false;
}
close(fd[0]);
write(fd[1], &testResult, sizeof(testResult));
close(fd[1]);
exit(0);
} else if (cPid > 0) {
close(fd[1]);
read(fd[0], &testResult, sizeof(testResult));
close(fd[0]);
wait(NULL);
} else {
printf("fork() failed\n");
testResult = false;
}
return testResult;
}
bool testValidDevicesBasic() {
bool testResult = true;
int numDevices = 0;
int device;
int validateCount = 0;
HIPCHECK(hipGetDeviceCount(&numDevices));
printf("Available compute devices in the system: %d\n", numDevices);
for (int i = 0; i < numDevices; i++) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipGetDevice(&device));
if (device == i) {
validateCount+= 1;
}
}
if (numDevices != validateCount) {
testResult = false;
}
return testResult;
}
int main(int argc, char* argv[]) {
bool testResult = true;
int numDevices = 0;
int device;
int deviceList[MAX_SIZE];
int extraArgs = 0;
#ifdef __unix__
getDeviceCount(&numDevices);
if (numDevices == 0) {
failed("No gpus found. exiting\n");
}
printf("Available compute devices in the system: %d\n", numDevices);
extraArgs = HipTest::parseStandardArguments(argc, argv, false);
parseExtraArguments(extraArgs, argv);
if (extraArgs == 1) {
printf("\nRunning test for invalid compute device\n");
// Test setting -1 to ROCR_VISIBLE_DEVICES
testResult &= testInvalidDevice(numDevices, true, -1);
// Test setting -1 to HIP_VISIBLE_DEVICES
testResult &= testInvalidDevice(numDevices, false, -1);
// Test setting invalid device to ROCR_VISIBLE_DEVICES
testResult &= testInvalidDevice(numDevices, true, numDevices);
// Test setting invalide device to HIP_VISIBLE_DEVICES
testResult &= testInvalidDevice(numDevices, false, numDevices);
// Test for all available devices
printf("\nRunning test for all available compute devices\n");
for (int i = 0; i < numDevices; i++) {
deviceList[i] = i;
}
testResult &= testValidDevices(numDevices, true, deviceList, numDevices);
testResult &= testValidDevices(numDevices, false, deviceList, numDevices);
}
// Assigning values to deviceList in reverse order
for (int i=0; i < deviceListLength; i++) {
deviceList[i] = deviceListLength-1-i;
}
// Test for subset of available gpus
if (extraArgs == 3) {
printf("\nRunning test for %d compute devices\n", deviceListLength);
testResult &= testValidDevices(numDevices, true, deviceList,
deviceListLength);
testResult &= testValidDevices(numDevices, false, deviceList,
deviceListLength);
}
#else
printf("Running basic test on Windows\n");
testResult &= testValidDevicesBasic();
#endif
if (testResult == true) {
passed();
} else {
failed("One or more tests failed\n");
}
}
projects/hip/tests/src/runtimeApi/memory/hipMemcpyNegativeMThrdMSize.cpp
+1188
Vedi File
File diff soppresso perché troppo grande Carica Diff