Update tests.
Fix some NVCC issues.
Add hipStreamSync2, record_event tests.
[ROCm/clr commit: 863b7c3f56]
Этот коммит содержится в:
Ben Sander
@@ -16,7 +16,7 @@ IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTI
|
||||
THE USE OR OTHER DEALINGS IN THE SOFTWARE. */
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s
|
||||
* BUILD: %t %s EXCLUDE_HIP_PLATFORM nvcc
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
@@ -0,0 +1,149 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 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 ../../test_common.cpp
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
enum SyncMode {
|
||||
syncNone,
|
||||
syncNullStream,
|
||||
syncOtherStream,
|
||||
};
|
||||
|
||||
|
||||
const char *syncModeString(int syncMode) {
|
||||
switch (syncMode) {
|
||||
case syncNone:
|
||||
return "syncNone";
|
||||
case syncNullStream:
|
||||
return "syncNullStream";
|
||||
case syncOtherStream:
|
||||
return "syncOtherStream";
|
||||
default:
|
||||
return "unknown";
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode)
|
||||
{
|
||||
printf ("\ntest: syncMode=%s\n", syncModeString(syncMode));
|
||||
|
||||
size_t sizeBytes = numElements * sizeof(int);
|
||||
|
||||
int count =100;
|
||||
int init0 = 0;
|
||||
HIPCHECK(hipMemset(C_d, init0, sizeBytes));
|
||||
for (int i=0; i<numElements; i++) {
|
||||
C_h[i] = -1; // initialize
|
||||
}
|
||||
|
||||
hipStream_t stream = 0;
|
||||
|
||||
unsigned flags=0;
|
||||
if (syncMode == syncOtherStream) {
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream, flags));
|
||||
}
|
||||
|
||||
hipEvent_t neverCreated=0;
|
||||
hipEvent_t start, stop, neverRecorded;
|
||||
HIPCHECK(hipEventCreate(&start));
|
||||
HIPCHECK(hipEventCreate(&stop));
|
||||
HIPCHECK(hipEventCreate(&neverRecorded));
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
|
||||
// sandwhich a kernel:
|
||||
HIPCHECK(hipEventRecord(start, stream));
|
||||
hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, stream, C_d, C_h, numElements, count);
|
||||
HIPCHECK(hipEventRecord(stop, stream));
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(stream)); // wait for recording to finish...
|
||||
|
||||
float t;
|
||||
HIPCHECK_API(hipEventElapsedTime(&t, neverCreated, stop), hipErrorInvalidResourceHandle);
|
||||
HIPCHECK_API(hipEventElapsedTime(&t, start, neverCreated), hipErrorInvalidResourceHandle);
|
||||
|
||||
HIPCHECK_API(hipEventElapsedTime(&t, neverRecorded, stop), hipErrorInvalidResourceHandle);
|
||||
HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidResourceHandle);
|
||||
|
||||
HIPCHECK(hipEventElapsedTime(&t, start, stop));
|
||||
assert (t>0.0f);
|
||||
printf ("time=%6.2f\n", t);
|
||||
|
||||
HIPCHECK(hipEventElapsedTime(&t, stop, start));
|
||||
assert (t<0.0f);
|
||||
printf ("negtime=%6.2f\n", t);
|
||||
|
||||
HIPCHECK(hipEventElapsedTime(&t, start, start));
|
||||
assert (t==0.0f);
|
||||
HIPCHECK(hipEventElapsedTime(&t, stop, stop));
|
||||
assert (t==0.0f);
|
||||
|
||||
|
||||
if (stream) {
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
HIPCHECK(hipEventDestroy(start));
|
||||
HIPCHECK(hipEventDestroy(stop));
|
||||
|
||||
printf ("test: OK \n");
|
||||
}
|
||||
|
||||
|
||||
|
||||
void runTests(int64_t numElements)
|
||||
{
|
||||
size_t sizeBytes = numElements * sizeof(int);
|
||||
|
||||
printf ("test: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes, sizeBytes/1024.0/1024.0);
|
||||
|
||||
|
||||
int *C_h, *C_d;
|
||||
HIPCHECK(hipMalloc(&C_d, sizeBytes));
|
||||
HIPCHECK(hipHostMalloc(&C_h, sizeBytes));
|
||||
|
||||
|
||||
{
|
||||
test (C_d, C_h, numElements, syncNone);
|
||||
test (C_d, C_h, numElements, syncNullStream);
|
||||
test (C_d, C_h, numElements, syncOtherStream);
|
||||
//test (C_d, C_h, numElements, syncDevice);
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipFree(C_d));
|
||||
HIPCHECK(hipHostFree(C_h));
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true /*failOnUndefinedArg*/);
|
||||
|
||||
runTests(4000000);
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -21,11 +21,12 @@
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <vector>
|
||||
#include"test_common.h"
|
||||
|
||||
#define LEN 1024*1024
|
||||
|
||||
@@ -1,132 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 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
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include"test_common.h"
|
||||
|
||||
#define NUM_ELEMENTS 1024*1024*64
|
||||
#define SIZE NUM_ELEMENTS*sizeof(int)
|
||||
|
||||
int p_count = 4;
|
||||
|
||||
|
||||
void multiGpuHostAlloc(int allocDevice)
|
||||
{
|
||||
|
||||
int numDevices;
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
printf ("info: trying multiGpuHostAlloc with allocDevice=%d numDevices=%d\n", allocDevice, numDevices);
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(allocDevice));
|
||||
|
||||
int *Ah, *Ch;
|
||||
hipHostMalloc((void**)&Ah, SIZE);
|
||||
hipHostMalloc((void**)&Ch, SIZE);
|
||||
|
||||
const int init = -1;
|
||||
for (size_t i=0; i<NUM_ELEMENTS; i++) {
|
||||
Ah[i] = init;
|
||||
Ch[i] = -2;
|
||||
}
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, NUM_ELEMENTS);
|
||||
|
||||
// The host memory allocations should be visible on all of the devices - verify by launching a kernel here that reads those devices:
|
||||
for (int i=0; i<numDevices; i++) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, 0/*_stream*/, Ah, Ch, NUM_ELEMENTS, p_count);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
};
|
||||
|
||||
|
||||
int expected = init + p_count;
|
||||
for (size_t i=0; i<NUM_ELEMENTS; i++) {
|
||||
if (Ch[i] != expected) {
|
||||
failed("for Ch[%zu] (%d) != expected(%d)\n", i, Ch[i], expected);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
|
||||
//assert(more_argc == 0);
|
||||
|
||||
{
|
||||
float *Ad, *B, *Bd, *Bm, *C, *Cd, *ptr_0;
|
||||
B = (float*)malloc(SIZE);
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipHostMalloc((void**)&B, SIZE);
|
||||
hipHostMalloc((void**)&Bd, SIZE, hipHostMallocDefault);
|
||||
hipHostMalloc((void**)&Bm, SIZE, hipHostMallocMapped);
|
||||
hipHostMalloc((void**)&C, SIZE, hipHostMallocMapped);
|
||||
|
||||
hipHostGetDevicePointer((void**)&Cd, C, 0/*flags*/);
|
||||
|
||||
HIPCHECK_API(hipMalloc((void**)&ptr_0,0), hipSuccess);
|
||||
|
||||
HIPCHECK_API(hipFree(Ad) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(Ad) , hipErrorInvalidValue);
|
||||
|
||||
HIPCHECK_API(hipFree(B) , hipErrorInvalidDevicePointer); // try to hipFree on malloced memory
|
||||
HIPCHECK_API(hipFree(Bd) , hipErrorInvalidDevicePointer);
|
||||
HIPCHECK_API(hipFree(Bm) , hipErrorInvalidDevicePointer);
|
||||
HIPCHECK_API(hipFree(ptr_0) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(Bd) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(Bm) , hipSuccess);
|
||||
|
||||
HIPCHECK_API(hipFree(C) , hipErrorInvalidDevicePointer);
|
||||
HIPCHECK_API(hipHostFree(C) , hipSuccess);
|
||||
|
||||
|
||||
HIPCHECK_API(hipFree(NULL) , hipSuccess);
|
||||
HIPCHECK_API(hipHostFree(NULL) , hipSuccess);
|
||||
|
||||
{
|
||||
// Some negative testing - request a too-big allocation and verify it fails:
|
||||
// Someday when we support virtual memory may need to refactor these:
|
||||
size_t tooBig = 128LL*1024*1024*1024*1024; // 128 TB;
|
||||
void *p;
|
||||
HIPCHECK_API ( hipMalloc(&p, tooBig), hipErrorMemoryAllocation );
|
||||
HIPCHECK_API ( hipHostMalloc(&p, tooBig), hipErrorMemoryAllocation );
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
{
|
||||
int numDevices;
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
multiGpuHostAlloc(0);
|
||||
if (numDevices > 1)
|
||||
{
|
||||
multiGpuHostAlloc(1);
|
||||
}
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,169 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 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 ../../test_common.cpp
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
enum SyncMode {
|
||||
syncNone,
|
||||
syncNullStream,
|
||||
syncOtherStream,
|
||||
syncMarkerThenOtherStream,
|
||||
syncMarkerThenOtherNonBlockingStream,
|
||||
syncDevice
|
||||
};
|
||||
|
||||
|
||||
const char *syncModeString(int syncMode) {
|
||||
switch (syncMode) {
|
||||
case syncNone:
|
||||
return "syncNone";
|
||||
case syncNullStream:
|
||||
return "syncNullStream";
|
||||
case syncOtherStream:
|
||||
return "syncOtherStream";
|
||||
case syncMarkerThenOtherStream:
|
||||
return "syncMarkerThenOtherStream";
|
||||
case syncMarkerThenOtherNonBlockingStream:
|
||||
return "syncMarkerThenOtherNonBlockingStream";
|
||||
case syncDevice:
|
||||
return "syncDevice";
|
||||
default:
|
||||
return "unknown";
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expectMismatch)
|
||||
{
|
||||
printf ("\ntest: syncMode=%s expectMismatch=%d\n", syncModeString(syncMode), expectMismatch);
|
||||
|
||||
size_t sizeBytes = numElements * sizeof(int);
|
||||
|
||||
int count =100;
|
||||
int init0 = 0;
|
||||
HIPCHECK(hipMemset(C_d, init0, sizeBytes));
|
||||
for (int i=0; i<numElements; i++) {
|
||||
C_h[i] = -1; // initialize
|
||||
}
|
||||
|
||||
hipStream_t otherStream = 0;
|
||||
unsigned flags = (syncMode == syncMarkerThenOtherNonBlockingStream) ? hipStreamNonBlocking : hipStreamDefault;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&otherStream, flags));
|
||||
hipEvent_t e;
|
||||
HIPCHECK(hipEventCreate(&e));
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
// Launch kernel into null stream, should result in C_h == count.
|
||||
hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, 0 /*stream*/, C_d, C_h, numElements, count);
|
||||
|
||||
switch (syncMode) {
|
||||
case syncNone:
|
||||
break;
|
||||
case syncNullStream:
|
||||
HIPCHECK(hipStreamSynchronize(0)); // wait on host for null stream:
|
||||
break;
|
||||
case syncOtherStream:
|
||||
// Does this synchronize with the null stream?
|
||||
HIPCHECK(hipStreamSynchronize(otherStream));
|
||||
break;
|
||||
case syncMarkerThenOtherStream:
|
||||
case syncMarkerThenOtherNonBlockingStream:
|
||||
HIPCHECK(hipEventRecord(e, otherStream)); // this may wait for NULL stream depending hipStreamNonBlocking flag above
|
||||
HIPCHECK(hipStreamSynchronize(otherStream));
|
||||
break;
|
||||
case syncDevice:
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
break;
|
||||
default:
|
||||
assert(0);
|
||||
};
|
||||
|
||||
int mismatches = 0;
|
||||
int expected = init0 + count;
|
||||
for (int i=0; i<numElements; i++) {
|
||||
bool compareEqual = (C_h[i] == expected);
|
||||
if (!compareEqual) {
|
||||
mismatches ++;
|
||||
if (!expectMismatch) {
|
||||
printf ("C_h[%d] (%d) != %d\n", i, C_h[i], expected);
|
||||
assert(C_h[i] == expected);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (expectMismatch) {
|
||||
assert (mismatches > 0);
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipStreamDestroy(otherStream));
|
||||
HIPCHECK(hipEventDestroy(e));
|
||||
|
||||
printf ("test: OK - %d mismatches (%6.2f%%)\n", mismatches, ((double)(mismatches)*100.0)/numElements);
|
||||
}
|
||||
|
||||
|
||||
void testEventRecord()
|
||||
{
|
||||
}
|
||||
|
||||
|
||||
void runTests(int64_t numElements)
|
||||
{
|
||||
size_t sizeBytes = numElements * sizeof(int);
|
||||
|
||||
printf ("\n\ntest: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes, sizeBytes/1024.0/1024.0);
|
||||
|
||||
|
||||
int *C_h, *C_d;
|
||||
HIPCHECK(hipMalloc(&C_d, sizeBytes));
|
||||
HIPCHECK(hipHostMalloc(&C_h, sizeBytes));
|
||||
|
||||
|
||||
{
|
||||
test (C_d, C_h, numElements, syncNone, true /*expectMismatch*/);
|
||||
test (C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/);
|
||||
test (C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/);
|
||||
test (C_d, C_h, numElements, syncDevice, false /*expectMismatch*/);
|
||||
test (C_d, C_h, numElements, syncMarkerThenOtherStream, false /*expectMismatch*/);
|
||||
test (C_d, C_h, numElements, syncMarkerThenOtherNonBlockingStream, true /*expectMismatch*/);
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipFree(C_d));
|
||||
HIPCHECK(hipHostFree(C_h));
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true /*failOnUndefinedArg*/);
|
||||
|
||||
runTests(40000000);
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -201,6 +201,20 @@ addCountReverse( const T *A_d,
|
||||
}
|
||||
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
memsetReverse( T *C_d, T val,
|
||||
int64_t NELEM)
|
||||
{
|
||||
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
|
||||
size_t stride = hipBlockDim_x * hipGridDim_x ;
|
||||
|
||||
for (int64_t i=NELEM-stride+offset; i>=0; i-=stride) {
|
||||
C_d[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
template <typename T>
|
||||
void setDefaultData(size_t numElements, T *A_h, T* B_h, T *C_h)
|
||||
{
|
||||
|
||||
Ссылка в новой задаче
Block a user