Refactor directed test infrastructue.
- Add hierarchy. Tests now live in directories, each with its own
CMakeFiles.txt. Reduces merge conflicts.
- Change make_hip_executable -> build_hip_executable.
- Refresh docs.
- Enable some tests that were previously built but not run.
Change-Id: I8c5de3c954400bf233904282b8b42861a2b7c536
[ROCm/clr commit: 3feb13c8f6]
This commit is contained in:
Ben Sander
@@ -0,0 +1,7 @@
|
||||
cmake_minimum_required (VERSION 2.6)
|
||||
|
||||
project (runtimeApi)
|
||||
|
||||
add_subdirectory(memory)
|
||||
add_subdirectory(multiThread)
|
||||
add_subdirectory(stream)
|
||||
@@ -0,0 +1,31 @@
|
||||
cmake_minimum_required (VERSION 2.6)
|
||||
|
||||
project (runtime_api)
|
||||
|
||||
include_directories( ${HIPTEST_SOURCE_DIR} )
|
||||
|
||||
build_hip_executable (hipMemset hipMemset.cpp)
|
||||
make_test(hipMemset " " )
|
||||
make_test(hipMemset --N 10 --memsetval 0x42 ) # small copy, just 10 bytes.
|
||||
make_test(hipMemset --N 10013 --memsetval 0x5a ) # oddball size.
|
||||
make_test(hipMemset --N 256M --memsetval 0xa6 ) # big copy
|
||||
|
||||
build_hip_executable (hipMemcpy_simple hipMemcpy_simple.cpp)
|
||||
make_test(hipMemcpy_simple " " )
|
||||
|
||||
build_hip_executable (hipMemcpy hipMemcpy.cpp)
|
||||
make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 )
|
||||
make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 )
|
||||
make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 )
|
||||
|
||||
build_hip_executable (hipMemcpyAsync hipMemcpyAsync.cpp)
|
||||
make_named_test(hipMemcpy_simple "hipMemcpyAsync-simple" --async)
|
||||
#make_test(hipMemcpyAsync " " )
|
||||
|
||||
build_hip_executable (hipMemoryAllocate hipMemoryAllocate.cpp)
|
||||
|
||||
build_hip_executable (hipMemcpyAll hipMemcpyAll.cpp)
|
||||
make_test(hipMemcpyAll " ")
|
||||
|
||||
# Debug synchronization, then enable.
|
||||
make_test(hipMemoryAllocate " ")
|
||||
@@ -0,0 +1,265 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
void printSep()
|
||||
{
|
||||
printf ("======================================================================================\n");
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
//---
|
||||
// Test many different kinds of memory copies.
|
||||
// The subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result.
|
||||
//
|
||||
// IN: numElements controls the number of elements used for allocations.
|
||||
// IN: usePinnedHost : If true, allocate host with hipHostMalloc and is pinned ; else allocate host memory with malloc.
|
||||
// IN: useHostToHost : If true, add an extra host-to-host copy.
|
||||
// IN: useDeviceToDevice : If true, add an extra deviceto-device copy after result is produced.
|
||||
// IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction.
|
||||
//
|
||||
template <typename T>
|
||||
void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, bool useDeviceToDevice, bool useMemkindDefault)
|
||||
{
|
||||
size_t sizeElements = numElements * sizeof(T);
|
||||
printf ("test: %s<%s> size=%lu (%6.2fMB) usePinnedHost:%d, useHostToHost:%d, useDeviceToDevice:%d, useMemkindDefault:%d\n",
|
||||
__func__,
|
||||
TYPENAME(T),
|
||||
sizeElements, sizeElements/1024.0/1024.0,
|
||||
usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault);
|
||||
|
||||
|
||||
T *A_d, *B_d, *C_d;
|
||||
T *A_h, *B_h, *C_h;
|
||||
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, numElements, usePinnedHost);
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
|
||||
T *A_hh = NULL;
|
||||
T *B_hh = NULL;
|
||||
T *C_dd = NULL;
|
||||
|
||||
|
||||
|
||||
if (useHostToHost) {
|
||||
if (usePinnedHost) {
|
||||
HIPCHECK ( hipHostMalloc((void**)&A_hh, sizeElements, hipHostMallocDefault) );
|
||||
HIPCHECK ( hipHostMalloc((void**)&B_hh, sizeElements, hipHostMallocDefault) );
|
||||
} else {
|
||||
A_hh = (T*)malloc(sizeElements);
|
||||
B_hh = (T*)malloc(sizeElements);
|
||||
}
|
||||
|
||||
|
||||
// Do some extra host-to-host copies here to mix things up:
|
||||
HIPCHECK ( hipMemcpy(A_hh, A_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost));
|
||||
HIPCHECK ( hipMemcpy(B_hh, B_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost));
|
||||
|
||||
|
||||
HIPCHECK ( hipMemcpy(A_d, A_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(B_d, B_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpy(A_d, A_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(B_d, B_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, numElements);
|
||||
|
||||
if (useDeviceToDevice) {
|
||||
HIPCHECK ( hipMalloc(&C_dd, sizeElements) );
|
||||
|
||||
// Do an extra device-to-device copies here to mix things up:
|
||||
HIPCHECK ( hipMemcpy(C_dd, C_d, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyDeviceToDevice));
|
||||
|
||||
//Destroy the original C_d:
|
||||
HIPCHECK ( hipMemset(C_d, 0x5A, sizeElements));
|
||||
|
||||
HIPCHECK ( hipMemcpy(C_h, C_dd, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpy(C_h, C_d, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
HIPCHECK ( hipDeviceSynchronize() );
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, numElements);
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, usePinnedHost);
|
||||
|
||||
printf (" %s success\n", __func__);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Try all the 16 possible combinations to memcpytest2 - usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault
|
||||
template<typename T>
|
||||
void memcpytest2_for_type(size_t numElements)
|
||||
{
|
||||
printSep();
|
||||
|
||||
for (int usePinnedHost =0; usePinnedHost<=1; usePinnedHost++) {
|
||||
for (int useHostToHost =0; useHostToHost<=1; useHostToHost++) { // TODO
|
||||
for (int useDeviceToDevice =0; useDeviceToDevice<=1; useDeviceToDevice++) {
|
||||
for (int useMemkindDefault =0; useMemkindDefault<=1; useMemkindDefault++) {
|
||||
memcpytest2<T>(numElements, usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Try many different sizes to memory copy.
|
||||
template<typename T>
|
||||
void memcpytest2_sizes(size_t maxElem=0, size_t offset=0)
|
||||
{
|
||||
printSep();
|
||||
printf ("test: %s<%s>\n", __func__, TYPENAME(T));
|
||||
|
||||
int deviceId;
|
||||
HIPCHECK(hipGetDevice(&deviceId));
|
||||
|
||||
size_t free, total;
|
||||
HIPCHECK(hipMemGetInfo(&free, &total));
|
||||
|
||||
if (maxElem == 0) {
|
||||
maxElem = free/sizeof(T)/5;
|
||||
}
|
||||
|
||||
printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB offset=%lu\n",
|
||||
deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0, offset);
|
||||
|
||||
for (size_t elem=64; elem+offset<=maxElem; elem*=2) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
memcpytest2<T>(elem+offset, 0, 1, 1, 0); // unpinned host
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
memcpytest2<T>(elem+offset, 1, 1, 1, 0); // pinned host
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Create multiple threads to stress multi-thread locking behavior in the allocation/deallocation/tracking logic:
|
||||
template<typename T>
|
||||
void multiThread_1(bool serialize, bool usePinnedHost)
|
||||
{
|
||||
printSep();
|
||||
printf ("test: %s<%s> serialize=%d usePinnedHost=%d\n", __func__, TYPENAME(T), serialize, usePinnedHost);
|
||||
std::thread t1 (memcpytest2<T>,N, usePinnedHost,0,0,0);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
}
|
||||
|
||||
|
||||
std::thread t2 (memcpytest2<T>,N, usePinnedHost,0,0,0);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join();
|
||||
t2.join();
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
|
||||
if (p_tests & 0x1) {
|
||||
printf ("\n\n=== tests&1 (types and different memcpy kinds (H2D, D2H, H2H, D2D)\n");
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
memcpytest2_for_type<float>(N);
|
||||
memcpytest2_for_type<double>(N);
|
||||
memcpytest2_for_type<char>(N);
|
||||
memcpytest2_for_type<int>(N);
|
||||
printf ("===\n\n\n");
|
||||
}
|
||||
|
||||
|
||||
if (p_tests & 0x2) {
|
||||
// Some tests around the 64MB boundary which have historically shown issues:
|
||||
printf ("\n\n=== tests&0x2 (64MB boundary)\n");
|
||||
#if 0
|
||||
// These all pass:
|
||||
memcpytest2<float>(15*1024*1024, 1, 0, 0, 0);
|
||||
memcpytest2<float>(16*1024*1024, 1, 0, 0, 0);
|
||||
memcpytest2<float>(16*1024*1024+16*1024, 1, 0, 0, 0);
|
||||
#endif
|
||||
// Just over 64MB:
|
||||
memcpytest2<float>(16*1024*1024+512*1024, 1, 0, 0, 0);
|
||||
memcpytest2<float>(17*1024*1024+1024, 1, 0, 0, 0);
|
||||
memcpytest2<float>(32*1024*1024, 1, 0, 0, 0);
|
||||
memcpytest2<float>(32*1024*1024, 0, 0, 0, 0);
|
||||
memcpytest2<float>(32*1024*1024, 1, 1, 1, 0);
|
||||
memcpytest2<float>(32*1024*1024, 1, 1, 1, 0);
|
||||
}
|
||||
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
printf ("\n\n=== tests&4 (test sizes and offsets)\n");
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(0,0);
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(0,64);
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(1024*1024, 13);
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(1024*1024, 50);
|
||||
}
|
||||
|
||||
if (p_tests & 0x8) {
|
||||
printf ("\n\n=== tests&8\n");
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
printSep();
|
||||
|
||||
// Simplest cases: serialize the threads, and also used pinned memory:
|
||||
// This verifies that the sub-calls to memcpytest2 are correct.
|
||||
multiThread_1<float>(true, true);
|
||||
|
||||
// Serialize, but use unpinned memory to stress the unpinned memory xfer path.
|
||||
multiThread_1<float>(true, false);
|
||||
|
||||
// Remove serialization, so two threads are performing memory copies in parallel.
|
||||
multiThread_1<float>(false, true);
|
||||
|
||||
// Remove serialization, and use unpinned.
|
||||
multiThread_1<float>(false, false); // TODO
|
||||
printf ("===\n\n\n");
|
||||
}
|
||||
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
@@ -0,0 +1,130 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
|
||||
#include<hip_runtime.h>
|
||||
#include<iostream>
|
||||
#include<assert.h>
|
||||
#include"test_common.h"
|
||||
|
||||
#define len 1024*1024
|
||||
#define size len * sizeof(float)
|
||||
|
||||
template<typename T>
|
||||
void hmemset(T *ptr, T value)
|
||||
{
|
||||
for(int i=0;i<len;i++){
|
||||
ptr[i] = value;
|
||||
}
|
||||
}
|
||||
|
||||
int main(){
|
||||
|
||||
int num;
|
||||
hipGetDeviceCount(&num);
|
||||
if(num < 2)
|
||||
{
|
||||
printf ("warning: Not enough GPUs to run the test, exiting without running.\n");
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
|
||||
float *h0, *h1;
|
||||
float *ph0, *ph1;
|
||||
float *d0, *d1;
|
||||
h0 = new float[len];
|
||||
h1 = new float[len];
|
||||
hmemset(h0, 1.0f);
|
||||
int gpu0 = 0, gpu1 = 1;
|
||||
hipSetDevice(gpu0);
|
||||
hipHostMalloc((void**)&ph0, size);
|
||||
hipMalloc(&d0, size);
|
||||
hipSetDevice(gpu1);
|
||||
hipHostMalloc((void**)&ph1, size);
|
||||
hipMalloc(&d1, size);
|
||||
hipSetDevice(gpu0);
|
||||
|
||||
|
||||
|
||||
hipMemcpy(h1, h0, size, hipMemcpyDefault);
|
||||
hipMemcpy(ph0, h1, size, hipMemcpyDefault);
|
||||
hipMemcpy(ph1, ph0, size, hipMemcpyDefault);
|
||||
assert(h0[0] == ph1[0]);
|
||||
hmemset(ph1, 0.0f);
|
||||
hipMemcpy(h0, ph1, size, hipMemcpyDefault);
|
||||
assert(h0[0] == 0.0f);
|
||||
|
||||
|
||||
|
||||
|
||||
hipSetDevice(gpu0);
|
||||
hmemset(ph0, 2.0f);
|
||||
hipMemcpy(d0, ph0, size, hipMemcpyDefault);
|
||||
hipMemcpy(h0, d0, size, hipMemcpyDefault);
|
||||
|
||||
assert(h0[0] == ph0[0]);
|
||||
hmemset(h0, 3.0f);
|
||||
hipMemcpy(d0, h0, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(ph0, d0, size, hipMemcpyDefault);
|
||||
|
||||
assert(h0[0] == ph0[0]);
|
||||
|
||||
hipSetDevice(gpu1);
|
||||
hmemset(ph1, 2.0f);
|
||||
hipMemcpy(d1, ph1, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(h1, d1, size, hipMemcpyDefault);
|
||||
|
||||
assert(h1[0] == ph1[0]);
|
||||
hmemset(h1, 3.0f);
|
||||
hipMemcpy(d1, h1, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(ph1, d1, size, hipMemcpyDefault);
|
||||
|
||||
assert(h1[0] == ph1[0]);
|
||||
|
||||
hipSetDevice(gpu0);
|
||||
hmemset(ph0, 4.0f);
|
||||
hipMemcpy(d0, ph0, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(ph0, d0, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(h0, d0, size, hipMemcpyDefault);
|
||||
|
||||
assert(ph0[0] == 4.0f);
|
||||
assert(h0[0] == 4.0f);
|
||||
|
||||
hipSetDevice(gpu1);
|
||||
hmemset(ph1, 5.0f);
|
||||
hipMemcpy(d1, ph1, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(ph1, d1, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(h1, d1, size, hipMemcpyDefault);
|
||||
|
||||
assert(ph1[0] == 5.0f);
|
||||
assert(h1[0] == 5.0f);
|
||||
|
||||
hipSetDevice(gpu0);
|
||||
hipMemcpy(d0, ph1, size, hipMemcpyDefault);
|
||||
|
||||
hipMemcpy(d1, d0, size, hipMemcpyDefault);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,372 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
// Test under-development. Calls async mem-copy API, experiment with functionality.
|
||||
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
unsigned p_streams = 2;
|
||||
|
||||
|
||||
void simpleNegTest()
|
||||
{
|
||||
printf ("testing: %s\n",__func__);
|
||||
hipError_t e;
|
||||
float *A_malloc, *A_pinned, *A_d;
|
||||
|
||||
size_t Nbytes = N*sizeof(float);
|
||||
A_malloc = (float*)malloc(Nbytes);
|
||||
HIPCHECK(hipHostMalloc((void**)&A_pinned, Nbytes, hipHostMallocDefault));
|
||||
A_d = NULL;
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPASSERT(A_d != NULL);
|
||||
// Can't use default with async copy
|
||||
e = hipMemcpyAsync(A_pinned, A_d, Nbytes, hipMemcpyDefault, NULL);
|
||||
// HIPASSERT (e == hipSuccess);
|
||||
|
||||
|
||||
// Not sure what happens here, the memory must be pinned.
|
||||
e = hipMemcpyAsync(A_malloc, A_d, Nbytes, hipMemcpyHostToDevice, NULL);
|
||||
|
||||
printf (" async memcpy of A_malloc to A_d. Result=%d\n", e);
|
||||
//HIPASSERT (e==hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
class Pinned;
|
||||
class Unpinned;
|
||||
|
||||
template <typename T> struct HostTraits;
|
||||
|
||||
template<>
|
||||
struct HostTraits<Pinned>
|
||||
{
|
||||
static const char *Name() { return "Pinned"; } ;
|
||||
|
||||
static void *Alloc(size_t sizeBytes) {
|
||||
void *p;
|
||||
HIPCHECK(hipHostMalloc((void**)&p, sizeBytes, hipHostMallocDefault));
|
||||
return p;
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
template<typename T>
|
||||
__global__ void
|
||||
addK (hipLaunchParm lp, T *A, T K, size_t numElements)
|
||||
{
|
||||
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
|
||||
size_t stride = hipBlockDim_x * hipGridDim_x ;
|
||||
|
||||
for (size_t i=offset; i<numElements; i+=stride) {
|
||||
A[i] = A[i] + K;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
//---
|
||||
//Tests propert dependency resolution between H2D and D2H commands in same stream:
|
||||
//IN: numInflight : number of copies inflight at any time:
|
||||
//IN: numPongs = number of iterations to run (iteration)
|
||||
template<typename T, class AllocType>
|
||||
void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs, bool doHostSide)
|
||||
{
|
||||
HIPASSERT(numElements % numInflight == 0); // Must be evenly divisible.
|
||||
size_t Nbytes = numElements*sizeof(T);
|
||||
size_t eachCopyElements = numElements / numInflight;
|
||||
size_t eachCopyBytes = eachCopyElements * sizeof(T);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
|
||||
printf ("-----------------------------------------------------------------------------------------------\n");
|
||||
printf ("testing: %s<%s> Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
|
||||
__func__, HostTraits<AllocType>::Name(), Nbytes, (double)(Nbytes)/1024.0/1024.0, numPongs, numInflight, eachCopyElements, eachCopyBytes);
|
||||
|
||||
T *A_h = NULL;
|
||||
T *A_d = NULL;
|
||||
|
||||
A_h = (T*)(HostTraits<AllocType>::Alloc(Nbytes));
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
|
||||
// Initialize the host array:
|
||||
const T initValue = 13;
|
||||
const T deviceConst = 2;
|
||||
const T hostConst = 10000;
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
A_h[i] = initValue + i;
|
||||
}
|
||||
|
||||
|
||||
for (int k=0; k<numPongs; k++ ) {
|
||||
for (int i=0; i<numInflight; i++) {
|
||||
HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
|
||||
HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
|
||||
}
|
||||
|
||||
hipLaunchKernel(addK<T>, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d, 2, numElements);
|
||||
|
||||
for (int i=0; i<numInflight; i++ ) {
|
||||
HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
|
||||
HIPCHECK(hipMemcpyAsync(&A_h[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
if (doHostSide) {
|
||||
assert(0);
|
||||
#if 0
|
||||
hipEvent_t e;
|
||||
HIPCHECK(hipEventCreate(&e));
|
||||
#endif
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
A_h[i] += hostConst;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
|
||||
// Verify we copied back all the data correctly:
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
T gold = initValue + i;
|
||||
// Perform calcs in same order as test above to replicate FP order-of-operations:
|
||||
for (int k=0; k<numPongs; k++) {
|
||||
gold += deviceConst;
|
||||
if (doHostSide) {
|
||||
gold += hostConst;
|
||||
}
|
||||
}
|
||||
|
||||
if (gold != A_h[i]) {
|
||||
std::cout << i << ": gold=" << gold << " out=" << A_h[i] << std::endl;
|
||||
HIPASSERT(gold == A_h[i]);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipHostFree(A_h));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Send many async copies to the same stream.
|
||||
//This requires runtime to keep track of many outstanding commands, and in the case of HCC requires growing/tracking the signal pool:
|
||||
template<typename T>
|
||||
void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies, bool syncBetweenCopies)
|
||||
{
|
||||
size_t Nbytes = numElements*sizeof(T);
|
||||
size_t eachCopyElements = numElements / numCopies;
|
||||
size_t eachCopyBytes = eachCopyElements * sizeof(T);
|
||||
|
||||
printf ("-----------------------------------------------------------------------------------------------\n");
|
||||
printf ("testing: %s Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
|
||||
__func__, Nbytes, (double)(Nbytes)/1024.0/1024.0, numCopies, eachCopyElements, eachCopyBytes);
|
||||
|
||||
T *A_d;
|
||||
T *A_h1, *A_h2;
|
||||
|
||||
HIPCHECK(hipHostMalloc((void**)&A_h1, Nbytes, hipHostMallocDefault));
|
||||
HIPCHECK(hipHostMalloc((void**)&A_h2, Nbytes, hipHostMallocDefault));
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
|
||||
for (int i=0; i<numElements; i++) {
|
||||
A_h1[i] = 3.14f + static_cast<T> (i);
|
||||
}
|
||||
|
||||
|
||||
//stream=0; // fixme TODO
|
||||
|
||||
|
||||
for (int i=0; i<numCopies; i++)
|
||||
{
|
||||
HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
|
||||
HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h1[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
|
||||
}
|
||||
|
||||
if (syncBetweenCopies) {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
for (int i=0; i<numCopies; i++)
|
||||
{
|
||||
HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
|
||||
HIPCHECK(hipMemcpyAsync(&A_h2[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
|
||||
// Verify we copied back all the data correctly:
|
||||
for (int i=0; i<numElements; i++) {
|
||||
HIPASSERT(A_h1[i] == A_h2[i]);
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipHostFree(A_h1));
|
||||
HIPCHECK(hipHostFree(A_h2));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Classic example showing how to overlap data transfer with compute.
|
||||
//We divide the work into "chunks" and create a stream for each chunk.
|
||||
//Each chunk then runs a H2D copy, followed by kernel execution, followed by D2H copyback.
|
||||
//Work in separate streams is independent which enables concurrency.
|
||||
|
||||
// IN: nStreams : number of streams to use for the test
|
||||
// IN :useNullStream - use NULL stream. Synchronizes everything.
|
||||
// IN: useSyncMemcpyH2D - use sync memcpy (no overlap) for H2D
|
||||
// IN: useSyncMemcpyD2H - use sync memcpy (no overlap) for D2H
|
||||
void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemcpyH2D, bool useSyncMemcpyD2H)
|
||||
{
|
||||
|
||||
size_t Nbytes = N*sizeof(int);
|
||||
printf ("testing: %s(useNullStream=%d, useSyncMemcpyH2D=%d, useSyncMemcpyD2H=%d) ",__func__, useNullStream, useSyncMemcpyH2D, useSyncMemcpyD2H);
|
||||
printf ("Nbytes=%zu (%6.1f MB)\n", Nbytes, (double)(Nbytes)/1024.0/1024.0);
|
||||
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, true);
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
|
||||
hipStream_t *stream = (hipStream_t*)malloc(sizeof(hipStream_t) * nStreams);
|
||||
if (useNullStream) {
|
||||
nStreams = 1;
|
||||
stream[0] = NULL;
|
||||
} else {
|
||||
for (int i = 0; i < nStreams; ++i) {
|
||||
HIPCHECK (hipStreamCreate(&stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
size_t workLeft = N;
|
||||
size_t workPerStream = N / nStreams;
|
||||
for (int i = 0; i < nStreams; ++i) {
|
||||
size_t work = (workLeft < workPerStream) ? workLeft : workPerStream;
|
||||
size_t workBytes = work * sizeof(int);
|
||||
|
||||
size_t offset = i*workPerStream;
|
||||
HIPASSERT(A_d + offset < A_d + Nbytes);
|
||||
HIPASSERT(B_d + offset < B_d + Nbytes);
|
||||
HIPASSERT(C_d + offset < C_d + Nbytes);
|
||||
if (useSyncMemcpyH2D) {
|
||||
HIPCHECK ( hipMemcpy(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpyAsync(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK ( hipMemcpyAsync(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice, stream[i]));
|
||||
};
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i], &A_d[offset], &B_d[offset], &C_d[offset], work);
|
||||
|
||||
if (useSyncMemcpyD2H) {
|
||||
HIPCHECK ( hipMemcpy(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpyAsync(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost, stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, true);
|
||||
|
||||
free(stream);
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
//Parse arguments specific to this test.
|
||||
void parseMyArguments(int argc, char *argv[])
|
||||
{
|
||||
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
|
||||
|
||||
// parse args for this test:
|
||||
for (int i = 1; i < more_argc; i++) {
|
||||
const char *arg = argv[i];
|
||||
|
||||
if (!strcmp(arg, "--streams")) {
|
||||
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) {
|
||||
failed("Bad streams argument");
|
||||
}
|
||||
} else {
|
||||
failed("Bad argument '%s'", arg);
|
||||
}
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
parseMyArguments(argc, argv);
|
||||
|
||||
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
if (p_tests & 0x01) {
|
||||
simpleNegTest();
|
||||
}
|
||||
|
||||
if (p_tests & 0x02) {
|
||||
hipStream_t stream;
|
||||
HIPCHECK (hipStreamCreate(&stream));
|
||||
|
||||
test_manyInflightCopies<float>(stream, 1024, 16, true);
|
||||
test_manyInflightCopies<float>(stream, 1024, 4, true); // verify we re-use the same entries instead of growing pool.
|
||||
test_manyInflightCopies<float>(stream, 1024*8, 64, false);
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
if (p_tests & 0x04) {
|
||||
test_chunkedAsyncExample(p_streams, true, true, true); // Easy sync version
|
||||
test_chunkedAsyncExample(p_streams, false, true, true); // Easy sync version
|
||||
test_chunkedAsyncExample(p_streams, false, false, true); // Some async
|
||||
test_chunkedAsyncExample(p_streams, false, false, false); // All async
|
||||
}
|
||||
|
||||
if (p_tests & 0x08) {
|
||||
hipStream_t stream;
|
||||
HIPCHECK (hipStreamCreate(&stream));
|
||||
|
||||
// test_pingpong<int, Pinned>(stream, 1024*1024*32, 1, 1, false);
|
||||
// test_pingpong<int, Pinned>(stream, 1024*1024*32, 1, 10, false);
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
@@ -0,0 +1,15 @@
|
||||
#include"test_common.h"
|
||||
|
||||
#define SIZE 1024*1024
|
||||
|
||||
int main(){
|
||||
float *A, *Ad;
|
||||
HIPCHECK(hipHostMalloc((void**)&A,SIZE, hipHostMallocDefault));
|
||||
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
for(int i=0;i<SIZE;i++){
|
||||
HIPCHECK(hipMemcpyAsync(Ad, A, SIZE, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,174 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
bool p_async = false;
|
||||
|
||||
// ****************************************************************************
|
||||
hipError_t memcopy(void * dst, const void *src, size_t sizeBytes, enum hipMemcpyKind kind)
|
||||
{
|
||||
if (p_async) {
|
||||
return hipMemcpyAsync(dst, src, sizeBytes, kind, NULL);
|
||||
} else {
|
||||
return hipMemcpy(dst, src, sizeBytes, kind);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
// Test simple H2D copies and back.
|
||||
// Designed to stress a small number of simple smoke tests
|
||||
void simpleTest1()
|
||||
{
|
||||
printf ("test: %s\n", __func__);
|
||||
size_t Nbytes = N*sizeof(int);
|
||||
printf ("N=%zu Nbytes=%6.2fMB\n", N, Nbytes/1024.0/1024.0);
|
||||
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
printf ("A_d=%p B_d=%p C_d=%p A_h=%p B_h=%p C_h=%p\n", A_d, B_d, C_d, A_h, B_d, C_h);
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HIPCHECK ( memcopy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK ( memcopy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, N);
|
||||
|
||||
HIPCHECK ( memcopy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK (hipDeviceReset());
|
||||
|
||||
printf (" %s success\n", __func__);
|
||||
}
|
||||
|
||||
|
||||
template <typename T>
|
||||
void simpleTest2(size_t numElements, bool usePinnedHost)
|
||||
{
|
||||
size_t sizeElements = numElements * sizeof(T);
|
||||
size_t alignment = 4096;
|
||||
printf ("test: %s<%s> numElements=%zu sizeElements=%zu bytes\n", __func__, TYPENAME(T), numElements, sizeElements);
|
||||
|
||||
T *A_d, *A_h1, *A_h2;
|
||||
|
||||
if (usePinnedHost) {
|
||||
HIPCHECK ( hipHostMalloc((void**)&A_h1, sizeElements, hipHostMallocDefault) );
|
||||
HIPCHECK ( hipHostMalloc((void**)&A_h2, sizeElements, hipHostMallocDefault) );
|
||||
} else {
|
||||
A_h1 = (T*)aligned_alloc(alignment, sizeElements);
|
||||
HIPASSERT(A_h1);
|
||||
A_h2 = (T*)aligned_alloc(alignment, sizeElements);
|
||||
HIPASSERT(A_h1);
|
||||
}
|
||||
|
||||
// Alloc device array:
|
||||
HIPCHECK ( hipMalloc(&A_d, sizeElements) );
|
||||
|
||||
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
A_h1[i] = 3.14f+ 1000*i;
|
||||
A_h2[i] = 12345678.0 + i; // init output with something distincctive, to ensure we replace it.
|
||||
}
|
||||
|
||||
HIPCHECK(memcopy(A_d, A_h1, sizeElements, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPCHECK(memcopy(A_h2, A_d, sizeElements, hipMemcpyDeviceToHost));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
HIPASSERT(A_h1[i] == A_h2[i]);
|
||||
}
|
||||
|
||||
HIPCHECK(hipFree(A_d));
|
||||
if (usePinnedHost) {
|
||||
HIPCHECK(hipHostFree(A_h1));
|
||||
HIPCHECK(hipHostFree(A_h2));
|
||||
} else {
|
||||
free(A_h1);
|
||||
free(A_h2);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//Parse arguments specific to this test.
|
||||
void parseMyArguments(int argc, char *argv[])
|
||||
{
|
||||
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
|
||||
|
||||
// parse args for this test:
|
||||
for (int i = 1; i < more_argc; i++) {
|
||||
const char *arg = argv[i];
|
||||
|
||||
if (!strcmp(arg, "--async")) {
|
||||
p_async = true;
|
||||
|
||||
} else {
|
||||
failed("Bad argument '%s'", arg);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
parseMyArguments(argc, argv);
|
||||
|
||||
printf ("info: set device to %d, tests=%x\n", p_gpuDevice, p_tests);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
|
||||
if (p_tests & 0x1) {
|
||||
printf ("\n\n=== tests&1\n");
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
simpleTest1();
|
||||
printf ("===\n\n\n");
|
||||
}
|
||||
|
||||
if (p_tests & 0x2) {
|
||||
printf ("\n\n=== tests&2 (copy ping-pong, pinned host)\n");
|
||||
simpleTest2<float>(N, true/*usePinnedHost*/);
|
||||
simpleTest2<char>(N, true/*usePinnedHost*/);
|
||||
}
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
printf ("\n\n=== tests&4 (copy ping-pong, unpinned host)\n");
|
||||
simpleTest2<char>(N, false/*usePinnedHost*/);
|
||||
simpleTest2<float>(N, false/*usePinnedHost*/);
|
||||
}
|
||||
|
||||
hipDeviceSynchronize();
|
||||
hipDeviceReset();
|
||||
|
||||
int v;
|
||||
hipDriverGetVersion(&v);
|
||||
|
||||
passed();
|
||||
};
|
||||
@@ -0,0 +1,50 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#include"test_common.h"
|
||||
|
||||
#define SIZE 1024*1024*256
|
||||
|
||||
int main(){
|
||||
float *Ad, *B, *Bd, *Bm, *C, *Cd;
|
||||
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(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(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);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,59 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
// Simple test for memset.
|
||||
// Also serves as a template for other tests.
|
||||
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
size_t Nbytes = N*sizeof(char);
|
||||
|
||||
printf ("N=%zu memsetval=%2x device=%d\n", N, memsetval, p_gpuDevice);
|
||||
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
|
||||
HIPCHECK ( hipMalloc(&A_d, Nbytes) );
|
||||
A_h = (char*)malloc(Nbytes);
|
||||
|
||||
HIPCHECK ( hipMemset(A_d, memsetval, Nbytes) );
|
||||
|
||||
HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i=0; i<N; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
failed("mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
@@ -0,0 +1,15 @@
|
||||
cmake_minimum_required (VERSION 2.6)
|
||||
|
||||
project (runtime_api)
|
||||
|
||||
include_directories( ${HIPTEST_SOURCE_DIR} )
|
||||
|
||||
build_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp)
|
||||
build_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp)
|
||||
build_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp)
|
||||
|
||||
#make_test(hipMultiThreadStreams1 " " ) Fails if 0x3 specified, passes otherwise.
|
||||
make_test(hipMultiThreadStreams2 " " )
|
||||
make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-serial" --tests 0x1)
|
||||
make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4)
|
||||
make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10)
|
||||
@@ -0,0 +1,138 @@
|
||||
#include <hip_runtime_api.h>
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
// Create a lot of streams and then destroy 'em.
|
||||
void createThenDestroyStreams(int iterations, int burstSize)
|
||||
{
|
||||
hipStream_t *streams = new hipStream_t[burstSize];
|
||||
|
||||
for (int i=0; i<iterations; i++) {
|
||||
if (p_verbose & 0x1) {
|
||||
printf ("%s iter=%d, create %d then destroy %d\n", __func__, i, burstSize, burstSize);
|
||||
}
|
||||
for (int j=0; j<burstSize; j++) {
|
||||
if (p_verbose & 0x2) {
|
||||
printf (" %d.%d streamCreate\n", i, j);
|
||||
}
|
||||
HIPCHECK( hipStreamCreate(&streams[j]));
|
||||
}
|
||||
for (int j=0; j<burstSize; j++) {
|
||||
if (p_verbose & 0x2) {
|
||||
printf (" %d.%d streamDestroy\n", i, j);
|
||||
}
|
||||
HIPCHECK( hipStreamDestroy(streams[j]));
|
||||
}
|
||||
}
|
||||
|
||||
delete streams;
|
||||
}
|
||||
|
||||
|
||||
void waitStreams(int iterations)
|
||||
{
|
||||
// Repeatedly sync and wait for all streams to complete.
|
||||
// TO make this interesting, the test has other threads repeatedly adding and removing streams to the device.
|
||||
for (int i=0; i<iterations; i++) {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// Create 3 streams, all creating and destroying streams on the same device.
|
||||
// Some create many queue, some not many.
|
||||
//
|
||||
void multiThread_pyramid(bool serialize, int iters)
|
||||
{
|
||||
printf ("%s creating %d streams\n", __func__, iters*100);
|
||||
std::thread t1 (createThenDestroyStreams, iters*1, 100);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
printf("t1 done\n");
|
||||
}
|
||||
|
||||
std::thread t2 (createThenDestroyStreams, iters*10, 10);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
printf("t2 done\n");
|
||||
}
|
||||
|
||||
std::thread t3 (createThenDestroyStreams, iters*100, 1);
|
||||
if (serialize) {
|
||||
t3.join();
|
||||
printf("t3 done\n");
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join();
|
||||
t2.join();
|
||||
t3.join();
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
// Create 3 streams, all creating and destroying streams on the same device.
|
||||
// Try to keep number of streams near zero, to cause problems.
|
||||
void multiThread_nearzero(bool serialize, int iters)
|
||||
{
|
||||
printf ("%s creating %d streams x 3 threads\n", __func__, iters);
|
||||
std::thread t1 (createThenDestroyStreams, iters, 1);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
printf("t1 done\n");
|
||||
}
|
||||
|
||||
std::thread t2 (createThenDestroyStreams, iters, 1);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
printf("t2 done\n");
|
||||
}
|
||||
|
||||
std::thread t3 (waitStreams, iters*50);
|
||||
if (serialize) {
|
||||
t3.join();
|
||||
printf("t3 done\n");
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join(); printf ("t1 done\n");
|
||||
t2.join(); printf ("t2 done\n");
|
||||
t3.join(); printf ("t3 done\n");
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
// Serial version, just call once:
|
||||
if (p_tests & 0x1) {
|
||||
printf ("\ntest 0x1 : serial createThenDestroyStreams(10) \n");
|
||||
createThenDestroyStreams(10, 10);
|
||||
};
|
||||
|
||||
/*disable, this takess a while and if the next one works then no need to run serial*/
|
||||
if (1 && (p_tests & 0x2)) {
|
||||
printf ("\ntest 0x2 : serialized multiThread_pyramid(1) \n");
|
||||
multiThread_pyramid(true, 3);
|
||||
}
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
printf ("\ntest 0x4 : parallel multiThread_pyramid(1) \n");
|
||||
multiThread_pyramid(false, 3);
|
||||
}
|
||||
|
||||
//if (p_tests & 0x8) {
|
||||
// printf ("test 0x8 : multiThread_pyramid(100) \n");
|
||||
// multiThread_pyramid(false, 100);
|
||||
// }
|
||||
|
||||
if (p_tests & 0x10) {
|
||||
printf ("\ntest 0x10 : parallel multiThread_nearzero(1000) \n");
|
||||
multiThread_nearzero(false, 1000);
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,144 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
|
||||
void printSep()
|
||||
{
|
||||
printf ("======================================================================================\n");
|
||||
}
|
||||
|
||||
//---
|
||||
// Test simple H2D copies and back.
|
||||
// Designed to stress a small number of simple smoke tests
|
||||
|
||||
template<
|
||||
typename T=float,
|
||||
class P=HipTest::Unpinned,
|
||||
class C=HipTest::Memcpy
|
||||
>
|
||||
void simpleVectorCopy(size_t numElements, int iters, hipStream_t stream)
|
||||
{
|
||||
using HipTest::MemTraits;
|
||||
|
||||
std::thread::id pid = std::this_thread::get_id();
|
||||
|
||||
printf ("test: %s <%s> %s %s\n", __func__, TYPENAME(T), P::str(), C::str());
|
||||
size_t Nbytes = numElements*sizeof(T);
|
||||
printf ("numElements=%zu Nbytes=%6.2fMB\n", numElements, Nbytes/1024.0/1024.0);
|
||||
|
||||
T *A_d, *B_d, *C_d;
|
||||
T *A_h, *B_h, *C_h;
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, P::isPinned);
|
||||
|
||||
|
||||
for (int i=0; i<iters; i++) {
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
|
||||
MemTraits<C>::Copy(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream);
|
||||
MemTraits<C>::Copy(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream);
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, numElements);
|
||||
|
||||
MemTraits<C>::Copy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream);
|
||||
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, numElements);
|
||||
}
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, P::isPinned);
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
|
||||
std::cout <<" pid" << pid << " success\n";
|
||||
}
|
||||
|
||||
template<typename T, class C>
|
||||
void test_multiThread_1(std::string testName, hipStream_t stream0, hipStream_t stream1, bool serialize)
|
||||
{
|
||||
printSep();
|
||||
printf ("%s\n", __func__);
|
||||
std::cout << testName << std::endl;
|
||||
|
||||
// Test 2 threads operating on same stream:
|
||||
std::thread t1 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 100/*iters*/, stream0);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
}
|
||||
std::thread t2 (simpleVectorCopy<T, HipTest::Pinned, C>, 2000000/*mb*/, 100/*iters*/, stream1);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join();
|
||||
t2.join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
};
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
|
||||
if (p_tests & 0x1) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK (hipStreamCreate(&stream));
|
||||
|
||||
simpleVectorCopy<float, HipTest::Pinned, HipTest::MemcpyAsync> (2000000/*mb*/, 10/*iters*/, stream);
|
||||
simpleVectorCopy<float, HipTest::Pinned, HipTest::Memcpy> (2000000/*mb*/, 10/*iters*/, stream);
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
hipStream_t stream0, stream1;
|
||||
HIPCHECK (hipStreamCreate(&stream0));
|
||||
HIPCHECK (hipStreamCreate(&stream1));
|
||||
|
||||
if (p_tests & 0x2) {
|
||||
// Easy tests to verify the test works - these don't allow overlap between the threads:
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread NULL with serialized", NULL, NULL, true);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread two streams serialized", stream0, stream1, true);
|
||||
}
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with NULL stream", NULL, NULL, false);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with two streams", stream0, stream1, false);
|
||||
test_multiThread_1<float, HipTest::MemcpyAsync> ("Multithread with one stream", stream0, stream0, false);
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
@@ -0,0 +1,108 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include<iostream>
|
||||
#include"test_common.h"
|
||||
#include<thread>
|
||||
#define N 1000
|
||||
|
||||
template<typename T>
|
||||
__global__ void Inc(hipLaunchParm lp, T *Array){
|
||||
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
|
||||
Array[tx] = Array[tx] + T(1);
|
||||
}
|
||||
|
||||
void run1(size_t size, hipStream_t stream){
|
||||
float *Ah, *Bh, *Cd, *Dd, *Eh;
|
||||
|
||||
HIPCHECK(hipHostMalloc((void**)&Ah, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipHostMalloc((void**)&Bh, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipMalloc(&Cd, size));
|
||||
HIPCHECK(hipMalloc(&Dd, size));
|
||||
HIPCHECK(hipHostMalloc((void**)&Eh, size, hipHostMallocDefault));
|
||||
|
||||
for(int i=0;i<N;i++){
|
||||
Ah[i] = 1.0f;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream));
|
||||
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
|
||||
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream));
|
||||
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT(Eh[10] == Ah[10] + 1.0f);
|
||||
}
|
||||
|
||||
|
||||
void run(size_t size, hipStream_t stream1, hipStream_t stream2){
|
||||
float *Ah, *Bh, *Cd, *Dd, *Eh;
|
||||
float *Ahh, *Bhh, *Cdd, *Ddd, *Ehh;
|
||||
|
||||
HIPCHECK(hipHostMalloc((void**)&Ah, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipHostMalloc((void**)&Bh, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipMalloc(&Cd, size));
|
||||
HIPCHECK(hipMalloc(&Dd, size));
|
||||
HIPCHECK(hipHostMalloc((void**)&Eh, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipHostMalloc((void**)&Ahh, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipHostMalloc((void**)&Bhh, size, hipHostMallocDefault));
|
||||
HIPCHECK(hipMalloc(&Cdd, size));
|
||||
HIPCHECK(hipMalloc(&Ddd, size));
|
||||
HIPCHECK(hipHostMalloc((void**)&Ehh, size, hipHostMallocDefault));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1));
|
||||
HIPCHECK(hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2));
|
||||
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1));
|
||||
HIPCHECK(hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2));
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream1, Cd);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream2, Cdd);
|
||||
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1));
|
||||
HIPCHECK(hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2));
|
||||
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1));
|
||||
HIPCHECK(hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT(Eh[10] = Ah[10] + 1.0f);
|
||||
HIPASSERT(Ehh[10] = Ahh[10] + 1.0f);
|
||||
}
|
||||
|
||||
int main(int argc, char **argv){
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
|
||||
hipStream_t stream[3];
|
||||
for(int i=0;i<3;i++){
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
const size_t size = N * sizeof(float);
|
||||
|
||||
std::thread t1(run1, size, stream[0]);
|
||||
std::thread t2(run1, size, stream[0]);
|
||||
std::thread t3(run, size, stream[1], stream[2]);
|
||||
t1.join();
|
||||
// std::cout<<"T1"<<std::endl;
|
||||
t2.join();
|
||||
// std::cout<<"T2"<<std::endl;
|
||||
t3.join();
|
||||
passed();
|
||||
}
|
||||
|
||||
@@ -0,0 +1,13 @@
|
||||
cmake_minimum_required (VERSION 2.6)
|
||||
|
||||
project (runtime_api)
|
||||
|
||||
include_directories( ${HIPTEST_SOURCE_DIR} )
|
||||
build_hip_executable (hipAPIStreamEnable hipAPIStreamEnable.cpp)
|
||||
build_hip_executable (hipAPIStreamDisable hipAPIStreamDisable.cpp)
|
||||
build_hip_executable (hipStreamL5 hipStreamL5.cpp)
|
||||
|
||||
# TODO - seg fault
|
||||
#make_test(hipAPIStreamEnable " ")
|
||||
#make_test(hipAPIStreamDisable " ")
|
||||
make_test(hipStreamL5 " ")
|
||||
@@ -0,0 +1,61 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include<iostream>
|
||||
#include"test_common.h"
|
||||
|
||||
const int NN = 1 << 21;
|
||||
|
||||
__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
|
||||
int tid = hipThreadIdx_x;
|
||||
if(tid < 1){
|
||||
for(int i=0;i<n;i++){
|
||||
x[i] = sqrt(pow(3.14159,i));
|
||||
}
|
||||
y[tid] = y[tid] + 1.0f;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void nKernel(hipLaunchParm lp, float *y){
|
||||
int tid = hipThreadIdx_x;
|
||||
y[tid] = y[tid] + 1.0f;
|
||||
}
|
||||
|
||||
int main(){
|
||||
const int num_streams = 8;
|
||||
hipStream_t streams[num_streams];
|
||||
float *data[num_streams], *yd, *xd;
|
||||
float y = 1.0f, x = 1.0f;
|
||||
HIPCHECK(hipMalloc((void**)&yd, sizeof(float)));
|
||||
HIPCHECK(hipMalloc((void**)&xd, sizeof(float)));
|
||||
HIPCHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice));
|
||||
for(int i=0;i<num_streams;i++){
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
HIPCHECK(hipMalloc(&data[i], NN * sizeof(float)));
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(1), dim3(1), 0, streams[i], data[i], xd, N);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(nKernel), dim3(1), dim3(1), 0, 0, yd);
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
|
||||
HIPCHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
|
||||
std::cout<<x<<" "<<y<<std::endl;
|
||||
HIPASSERT(x == y);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,63 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
//#define HIP_API_PER_THREAD_DEFAULT_STREAM
|
||||
|
||||
#include<iostream>
|
||||
#include"test_common.h"
|
||||
|
||||
const int NN = 1 << 21;
|
||||
|
||||
__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
|
||||
int tid = hipThreadIdx_x;
|
||||
if(tid < 1){
|
||||
for(int i=0;i<n;i++){
|
||||
x[i] = sqrt(pow(3.14159,i));
|
||||
}
|
||||
y[tid] = y[tid] + 1.0f;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void nKernel(hipLaunchParm lp, float *y){
|
||||
int tid = hipThreadIdx_x;
|
||||
y[tid] = y[tid] + 1.0f;
|
||||
}
|
||||
|
||||
int main(){
|
||||
const int num_streams = 8;
|
||||
hipStream_t streams[num_streams];
|
||||
float *data[num_streams], *yd, *xd;
|
||||
float y = 1.0f, x = 1.0f;
|
||||
HIPCHECK(hipMalloc((void**)&yd, sizeof(float)));
|
||||
HIPCHECK(hipMalloc((void**)&xd, sizeof(float)));
|
||||
HIPCHECK(hipMemcpy(yd, &y, sizeof(float), hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(xd, &x, sizeof(float), hipMemcpyHostToDevice));
|
||||
for(int i=0;i<num_streams;i++){
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
HIPCHECK(hipMalloc(&data[i], NN * sizeof(float)));
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(1), dim3(1), 0, streams[i], data[i], xd, N);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(nKernel), dim3(1), dim3(1), 0, 0, yd);
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
|
||||
HIPCHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
|
||||
std::cout<<x<<" "<<y<<std::endl;
|
||||
HIPASSERT(x<y);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,121 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#ifndef HIPSTREAM_H
|
||||
#define HIPSTREAM_H
|
||||
#include<hip_runtime.h>
|
||||
|
||||
#define NUM_STREAMS 4
|
||||
|
||||
/*
|
||||
* H2H - 1
|
||||
* H2D - 2
|
||||
* KER - 3
|
||||
* D2D - 4
|
||||
* D2H - 5
|
||||
*/
|
||||
|
||||
template<typename T>
|
||||
void H2HAsync(T *Dst, T *Src, size_t size, hipStream_t stream){
|
||||
HIPCHECK(hipMemcpyAsync(Dst, Src, size, hipMemcpyHostToHost, stream));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void H2DAsync(T *Dst, T *Src, size_t size, hipStream_t stream){
|
||||
HIPCHECK(hipMemcpyAsync(Dst, Src, size, hipMemcpyHostToDevice, stream));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void D2DAsync(T *Dst, T *Src, size_t size, hipStream_t stream){
|
||||
HIPCHECK(hipMemcpyAsync(Dst, Src, size, hipMemcpyDeviceToDevice, stream));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void D2HAsync(T *Dst, T *Src, size_t size, hipStream_t stream){
|
||||
HIPCHECK(hipMemcpyAsync(Dst, Src, size, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void H2H(T *Dst, T *Src, size_t size){
|
||||
HIPCHECK(hipMemcpy(Dst, Src, size, hipMemcpyHostToHost));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void H2D(T *Dst, T *Src, size_t size){
|
||||
HIPCHECK(hipMemcpy(Dst, Src, size, hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void D2D(T *Dst, T *Src, size_t size){
|
||||
HIPCHECK(hipMemcpy(Dst, Src, size, hipMemcpyDeviceToDevice));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void D2H(T *Dst, T *Src, size_t size){
|
||||
HIPCHECK(hipMemcpy(Dst, Src, size, hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
__global__ void Inc(hipLaunchParm lp, T *In){
|
||||
int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
|
||||
In[tx] = In[tx] + 1;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void initArrays(T **Ad, T **Ah,
|
||||
size_t N, bool usePinnedHost=false){
|
||||
size_t NBytes = N * sizeof(T);
|
||||
if(Ad){
|
||||
HIPCHECK( hipMalloc(Ad, NBytes));
|
||||
}
|
||||
if(usePinnedHost){
|
||||
HIPCHECK( hipHostMalloc((void**)Ah, NBytes, hipHostMallocDefault));
|
||||
}
|
||||
else{
|
||||
*Ah = new T[N];
|
||||
HIPASSERT(*Ah != NULL);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void initArrays(T **Ad, size_t N,
|
||||
bool deviceMemory = false,
|
||||
bool usePinnedHost = false){
|
||||
size_t NBytes = N * sizeof(T);
|
||||
if(deviceMemory){
|
||||
HIPCHECK( hipMalloc(Ad, NBytes));
|
||||
}else{
|
||||
if(usePinnedHost){
|
||||
HIPCHECK(hipHostMalloc((void**)Ad, NBytes, hipHostMallocDefault));
|
||||
}else{
|
||||
*Ad = new T[N];
|
||||
HIPASSERT(*Ad != NULL);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void setArray(T* Array, int N, T val){
|
||||
for(int i=0;i<N;i++){
|
||||
Array[i] = val;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,813 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#include "hipStream.h"
|
||||
|
||||
/*
|
||||
The naming of tests is done by assigning a number to
|
||||
type of disptach possible on stream.
|
||||
The following are possible stream dispatches:
|
||||
1. H2H - hipMemcpyHostToHost : indexed as 1
|
||||
2. H2D - hipMemcpyHostToDevice : indexed as 2
|
||||
3. Ker - Kernel Dispatch : indexed as 3
|
||||
4. D2D - hipMemcpyDeviceToDevice : indexed as 4
|
||||
5. D2H - hipMemcpyDeviceToHost : indexed as 5
|
||||
For example,
|
||||
a test for Ker, D2D, D2H, H2H, H2D is given as test34512();
|
||||
Note that all memory copies are Async.
|
||||
|
||||
invalid{
|
||||
*WARNING: The commented out assertions are failing cases.
|
||||
According to my observation, they are happening with tests
|
||||
which end in HostToHost and take data from previous
|
||||
dispatch in the stream. This also include disjoint data passes.
|
||||
The list of failing tests are:
|
||||
test23451<float>();
|
||||
test32451<float>();
|
||||
test42351<float>();
|
||||
|
||||
For disjoint data passed:
|
||||
test24513
|
||||
test25134
|
||||
test34512
|
||||
}
|
||||
*/
|
||||
|
||||
template<typename T>
|
||||
void test12345(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Ad, *Bd;
|
||||
initArrays(&Ad, &Ah, N, true);
|
||||
initArrays(&Bd, &Bh, N, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
H2DAsync(Ad, Bh, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Ch, Bd, size, stream);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Ah[10] + T(1)== Ch[10]);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test13452(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Dh, N, T(2));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Ch, Bd, size, stream);
|
||||
H2DAsync(Cd, Ch, size, stream);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
D2H(Eh,Cd,size);
|
||||
|
||||
HIPASSERT(Ah[10] == Bh[10]);
|
||||
HIPASSERT(Eh[10] == Dh[10] + T(1));
|
||||
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test14523(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Dh, N, T(2));
|
||||
|
||||
H2D(Ad,Dh,size);
|
||||
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Ch, Bd, size, stream);
|
||||
H2DAsync(Cd, Ch, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
D2H(Eh, Cd, size);
|
||||
|
||||
HIPASSERT(Ah[10] == Bh[10]);
|
||||
HIPASSERT(Ch[10] + T(1) == Eh[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test15234(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Dh, N, T(2));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
D2HAsync(Ch, Ad, size, stream);
|
||||
H2DAsync(Bd, Ch, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
|
||||
D2H(Eh, Cd, size);
|
||||
|
||||
HIPASSERT(Ah[10] == Bh[10]);
|
||||
HIPASSERT(Eh[10] == Dh[10] + T(1));
|
||||
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test23451(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Ad, *Bd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
|
||||
H2DAsync(Ad, Ah, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Bh, Bd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT(Ah[10] + T(1) == Ch[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test24513(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Dh, N, T(2));
|
||||
|
||||
H2D(Cd, Dh, size);
|
||||
|
||||
H2DAsync(Ad, Ah, size, stream);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Bh, Bd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
D2H(Eh, Cd, size);
|
||||
|
||||
HIPASSERT(Eh[0] == Dh[0] + T(1));
|
||||
HIPASSERT(Ah[0] == Ch[0]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test25134(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Dh, N, T(2));
|
||||
|
||||
H2D(Bd, Dh, size);
|
||||
|
||||
H2DAsync(Ad, Ah, size, stream);
|
||||
D2HAsync(Bh, Ad, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
|
||||
D2H(Eh, Cd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Ah[10] == Ch[10]);
|
||||
HIPASSERT(Dh[10] + T(1) == Eh[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test21345(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh, *Ch, *Dh;
|
||||
T *Ad, *Bd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, true);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Bh, N, T(2));
|
||||
|
||||
H2DAsync(Ad, Ah, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Dh, Bd, size, stream);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT( Bh[10] == Ch[10] );
|
||||
HIPASSERT( Ah[10] + T(1) == Dh[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test34512(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Bh, *Ch, *Dh;
|
||||
T *Ah, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, true);
|
||||
initArrays(&Ah, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
|
||||
H2D(Ad, Ah, size);
|
||||
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Bh, Bd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
H2DAsync(Cd, Ch, size, stream);
|
||||
|
||||
D2H(Dh, Cd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT( Ah[10] + T(1) == Dh[10] );
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test35124(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh;
|
||||
T *Ch, *Dh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, false);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
D2HAsync(Ah, Ad, size, stream);
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
H2DAsync(Bd, Bh, size, stream);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
|
||||
D2H(Ch, Cd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Dh[10] + T(1) == Ch[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test31245(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
setArray(Ah, N, T(2));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
H2DAsync(Bd, Bh, size, stream);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
D2HAsync(Ch, Cd, size, stream);
|
||||
|
||||
D2H(Eh, Ad, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Dh[10] + T(1) == Eh[10]);
|
||||
HIPASSERT(Bh[10] == Ch[10]);
|
||||
}
|
||||
|
||||
|
||||
template<typename T>
|
||||
void test32451(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ah, N, T(1));
|
||||
setArray(Eh, N, T(2));
|
||||
|
||||
H2D(Ad, Eh, size);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
|
||||
H2DAsync(Bd, Ah, size, stream);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
D2HAsync(Bh, Cd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
D2H(Dh, Ad, size);
|
||||
|
||||
HIPASSERT(Ah[10] == Ch[10]);
|
||||
HIPASSERT(Eh[10] + T(1) == Dh[10]);
|
||||
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test45123(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh;
|
||||
T *Ch, *Dh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, false);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Ah, Bd, size, stream);
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
H2DAsync(Cd, Bh, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
|
||||
D2H(Ch, Cd, size);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Dh[10] + T(1) == Ch[10]);
|
||||
}
|
||||
|
||||
|
||||
template<typename T>
|
||||
void test41235(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh;
|
||||
T *Ch;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Ch, N, T(1));
|
||||
|
||||
H2D(Ad, Ch, size);
|
||||
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
D2HAsync(Ah, Bd, size, stream);
|
||||
H2DAsync(Cd, Ah, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
|
||||
D2HAsync(Bh, Cd, size, stream);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Ch[10] + T(1) == Bh[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test42351(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(2));
|
||||
setArray(Ah, N, T(1));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
H2DAsync(Cd, Ah, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
|
||||
D2HAsync(Bh, Cd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
|
||||
D2H(Eh, Bd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT(Dh[10] == Eh[10]);
|
||||
HIPASSERT(Ah[10] + T(1) == Ch[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test43512(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh;
|
||||
T *Ch, *Dh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, false);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
D2DAsync(Bd, Ad, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
|
||||
D2HAsync(Ah, Bd, size, stream);
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
H2DAsync(Cd, Bh, size, stream);
|
||||
|
||||
D2H(Ch, Cd, size);
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT( Dh[10] + T(1) == Ch[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test51234(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh;
|
||||
T *Ch, *Dh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, false);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
|
||||
D2HAsync(Ah, Ad, size, stream);
|
||||
H2HAsync(Bh, Ah, size, stream);
|
||||
H2DAsync(Bd, Bh, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
|
||||
D2H(Ch, Cd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Ch[10] == Dh[10] + T(1));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test52341(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh;
|
||||
T *Ad, *Bd, *Cd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
|
||||
setArray(Eh, N, T(1));
|
||||
setArray(Bh, N, T(2));
|
||||
|
||||
H2D(Ad, Eh, size);
|
||||
|
||||
D2HAsync(Ah, Ad, size, stream);
|
||||
H2DAsync(Bd, Ah, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
|
||||
D2H(Dh, Cd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
HIPASSERT(Eh[10] + T(1) == Dh[10]);
|
||||
HIPASSERT(Ch[10] == Bh[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test53412(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
const size_t size = sizeof(T) * N;
|
||||
|
||||
T *Ah, *Bh, *Ch, *Dh;
|
||||
T *Eh, *Fh, *Gh;
|
||||
T *Ad, *Bd, *Cd, *Dd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, true);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Fh, N, false, false);
|
||||
initArrays(&Gh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
initArrays(&Dd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
setArray(Eh, N, T(2));
|
||||
setArray(Bh, N, T(3));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
H2D(Bd, Eh, size);
|
||||
|
||||
D2HAsync(Ah, Ad, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
H2DAsync(Dd, Ch, size, stream);
|
||||
|
||||
D2H(Fh, Cd, size);
|
||||
D2H(Gh, Dd, size);
|
||||
|
||||
HIPASSERT(Ah[10] == Dh[10]);
|
||||
HIPASSERT(Eh[10] + T(1) == Fh[10]);
|
||||
HIPASSERT(Bh[10] == Gh[10]);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void test54123(){
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
const size_t size = N * sizeof(T);
|
||||
|
||||
T *Ah, *Bh, *Ch;
|
||||
T *Dh, *Eh, *Fh, *Gh;
|
||||
T *Ad, *Bd, *Cd, *Dd;
|
||||
|
||||
initArrays(&Ah, N, false, true);
|
||||
initArrays(&Bh, N, false, true);
|
||||
initArrays(&Ch, N, false, true);
|
||||
initArrays(&Dh, N, false, false);
|
||||
initArrays(&Eh, N, false, false);
|
||||
initArrays(&Fh, N, false, false);
|
||||
initArrays(&Gh, N, false, false);
|
||||
initArrays(&Ad, N, true, false);
|
||||
initArrays(&Bd, N, true, false);
|
||||
initArrays(&Cd, N, true, false);
|
||||
initArrays(&Dd, N, true, false);
|
||||
|
||||
setArray(Dh, N, T(1));
|
||||
setArray(Eh, N, T(1));
|
||||
setArray(Bh, N, T(1));
|
||||
|
||||
H2D(Ad, Dh, size);
|
||||
H2D(Bd, Eh, size);
|
||||
|
||||
D2HAsync(Ah, Ad, size, stream);
|
||||
D2DAsync(Cd, Bd, size, stream);
|
||||
H2HAsync(Ch, Bh, size, stream);
|
||||
H2DAsync(Dd, Ch, size, stream);
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Dd);
|
||||
|
||||
D2H(Fh, Cd, size);
|
||||
D2H(Gh, Dd, size);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPASSERT(Dh[10] == Ah[10]);
|
||||
HIPASSERT(Eh[10] == Fh[10]);
|
||||
HIPASSERT(Bh[10] + T(1) == Gh[10]);
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
test12345<float>();
|
||||
test13452<float>();
|
||||
test14523<float>();
|
||||
test15234<float>();
|
||||
|
||||
test23451<float>();
|
||||
test24513<float>();
|
||||
test25134<float>();
|
||||
test21345<float>();
|
||||
|
||||
test34512<float>();
|
||||
test35124<float>();
|
||||
test31245<float>();
|
||||
test32451<float>();
|
||||
|
||||
test45123<float>();
|
||||
test41235<float>();
|
||||
test42351<float>();
|
||||
test43512<float>();
|
||||
|
||||
test51234<float>();
|
||||
test52341<float>();
|
||||
test53412<float>();
|
||||
test54123<float>();
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user