AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Apply .clangformat to all repo source files

Browse Source 
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344


[ROCm/clr commit: 9e47fccc89]
This commit is contained in:
Maneesh Gupta
2018-03-12 11:29:03 +05:30
parent ecbb701440
commit 46ddefedee
293 changed files with 43980 additions and 45830 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp
+34 -34
View File
@@ -22,46 +22,46 @@ THE SOFTWARE.
* HIT_END
*/
#include<iostream>
#include"test_common.h"
#include"hip/math_functions.h"
#include <iostream>
#include "test_common.h"
#include "hip/math_functions.h"
const int NN = 1 << 21;
__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
int tid = threadIdx.x;
if(tid < 1){
for(int i=0;i<n;i++){
x[i] = sqrt(powf(3.14159,i));
}
y[tid] = y[tid] + 1.0f;
}
__global__ void kernel(hipLaunchParm lp, float* x, float* y, int n) {
int tid = threadIdx.x;
if (tid < 1) {
for (int i = 0; i < n; i++) {
x[i] = sqrt(powf(3.14159, i));
}
y[tid] = y[tid] + 1.0f;
}
}
__global__ void nKernel(hipLaunchParm lp, float *y){
int tid = threadIdx.x;
y[tid] = y[tid] + 1.0f;
__global__ void nKernel(hipLaunchParm lp, float* y) {
int tid = threadIdx.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);
}
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();
HIPCHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
std::cout << x << " " << y << std::endl;
HIPASSERT(x == y);
passed();
}
projects/clr/hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp
+34 -34
View File
@@ -24,46 +24,46 @@ THE SOFTWARE.
//#define HIP_API_PER_THREAD_DEFAULT_STREAM
#include<iostream>
#include"test_common.h"
#include"hip/math_functions.h"
#include <iostream>
#include "test_common.h"
#include "hip/math_functions.h"
const int NN = 1 << 21;
__global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
int tid = threadIdx.x;
if(tid < 1){
for(int i=0;i<n;i++){
x[i] = sqrt(powf(3.14159,i));
}
y[tid] = y[tid] + 1.0f;
}
__global__ void kernel(hipLaunchParm lp, float* x, float* y, int n) {
int tid = threadIdx.x;
if (tid < 1) {
for (int i = 0; i < n; i++) {
x[i] = sqrt(powf(3.14159, i));
}
y[tid] = y[tid] + 1.0f;
}
}
__global__ void nKernel(hipLaunchParm lp, float *y){
int tid = threadIdx.x;
y[tid] = y[tid] + 1.0f;
__global__ void nKernel(hipLaunchParm lp, float* y) {
int tid = threadIdx.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);
}
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();
HIPCHECK(hipMemcpy(&x, xd, sizeof(float), hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(&y, yd, sizeof(float), hipMemcpyDeviceToHost));
std::cout << x << " " << y << std::endl;
HIPASSERT(x < y);
passed();
}
projects/clr/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp
+87 -127
View File
@@ -27,39 +27,31 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "test_common.h"
#include <vector>
unsigned p_streams =16;
int p_repeat = 10;
int p_db = 0;
unsigned p_streams = 16;
int p_repeat = 10;
int p_db = 0;
template <typename T>
__global__ void
vectorADDRepeat(hipLaunchParm lp,
const T *A_d,
const T *B_d,
T *C_d,
size_t NELEM,
int repeat)
{
__global__ void vectorADDRepeat(hipLaunchParm lp, const T* A_d, const T* B_d, T* C_d, size_t NELEM,
int repeat) {
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
size_t stride = blockDim.x * gridDim.x;
for (int j=1; j<=repeat;j++) {
for (size_t i=offset; i<NELEM; i+=stride) {
C_d[i] = A_d[i]*j + B_d[i]*j;
for (int j = 1; j <= repeat; j++) {
for (size_t i = offset; i < NELEM; i += stride) {
C_d[i] = A_d[i] * j + B_d[i] * j;
}
};
}
//------
// Structure for one stream - includes the stream + data buffers that are used by the stream.
template <typename T>
class Streamer {
public:
Streamer(size_t numElements, bool useNullStream=false);
public:
Streamer(size_t numElements, bool useNullStream = false);
~Streamer();
void enqueAsync();
void queryUntilComplete();
@@ -69,26 +61,24 @@ public:
void D2H();
public:
T *_A_h;
T *_B_h;
T *_C_h;
public:
T* _A_h;
T* _B_h;
T* _C_h;
T *_A_d;
T *_B_d;
T *_C_d;
T* _A_d;
T* _B_d;
T* _C_d;
hipStream_t _stream;
hipEvent_t _event;
hipEvent_t _event;
size_t _numElements;
size_t _numElements;
};
template <typename T>
Streamer<T>::Streamer(size_t numElements, bool useNullStream) :
_numElements(numElements)
{
HipTest::initArrays (&_A_d, &_B_d, &_C_d, &_A_h, &_B_h, &_C_h, numElements, true);
Streamer<T>::Streamer(size_t numElements, bool useNullStream) : _numElements(numElements) {
HipTest::initArrays(&_A_d, &_B_d, &_C_d, &_A_h, &_B_h, &_C_h, numElements, true);
if (useNullStream) {
_stream = 0x0;
@@ -98,82 +88,65 @@ Streamer<T>::Streamer(size_t numElements, bool useNullStream) :
HIPCHECK(hipEventCreate(&_event));
H2D();
};
template <typename T>
void Streamer<T>::H2D()
{
HIPCHECK(hipMemcpy(_A_d, _A_h, _numElements*sizeof(T), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(_B_d, _B_h, _numElements*sizeof(T), hipMemcpyHostToDevice));
void Streamer<T>::H2D() {
HIPCHECK(hipMemcpy(_A_d, _A_h, _numElements * sizeof(T), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(_B_d, _B_h, _numElements * sizeof(T), hipMemcpyHostToDevice));
}
template <typename T>
void Streamer<T>::D2H()
{
HIPCHECK(hipMemcpy(_C_h, _C_d, _numElements*sizeof(T), hipMemcpyDeviceToHost));
void Streamer<T>::D2H() {
HIPCHECK(hipMemcpy(_C_h, _C_d, _numElements * sizeof(T), hipMemcpyDeviceToHost));
}
template <typename T>
void Streamer<T>::reset()
{
void Streamer<T>::reset() {
HipTest::setDefaultData(_numElements, _A_h, _B_h, _C_h);
H2D();
}
template <typename T>
void Streamer<T>::enqueAsync()
{
printf ("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0);
void Streamer<T>::enqueAsync() {
printf("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements,
_numElements * sizeof(T) / 1024.0 / 1024.0);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
hipLaunchKernel(
vectorADDRepeat,
dim3(blocks),
dim3(threadsPerBlock),
0,
_stream,
static_cast<const T*>(_A_d),
static_cast<const T*>(_B_d),
_C_d,
_numElements,
p_repeat);
hipLaunchKernel(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0, _stream,
static_cast<const T*>(_A_d), static_cast<const T*>(_B_d), _C_d, _numElements,
p_repeat);
}
template <typename T>
void Streamer<T>::queryUntilComplete()
{
void Streamer<T>::queryUntilComplete() {
int numQueries = 0;
hipError_t e = hipSuccess;
do {
numQueries++;
e = hipStreamQuery(_stream);
} while (e != hipSuccess) ;
} while (e != hipSuccess);
printf ("completed after %d queries\n", numQueries);
printf("completed after %d queries\n", numQueries);
};
//---
//Parse arguments specific to this test.
void parseMyArguments(int argc, char *argv[])
{
// 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];
const char* arg = argv[i];
if (!strcmp(arg, "--streams")) {
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) {
failed("Bad streams argument");
failed("Bad streams argument");
}
} else if (!strcmp(arg, "--repeat") || (!strcmp(arg, "-r"))) {
if (++i >= argc || !HipTest::parseInt(argv[i], &p_repeat)) {
failed("Bad repeat argument");
failed("Bad repeat argument");
}
} else {
failed("Bad argument '%s'", arg);
@@ -182,70 +155,61 @@ void parseMyArguments(int argc, char *argv[])
};
void
printBuffer(std::string name, int *f, size_t numElements)
{
void printBuffer(std::string name, int* f, size_t numElements) {
std::cout << name << "\n";
for (size_t i=0; i<numElements; i++) {
printf ("%5zu: %d\n", i, f[i]);
for (size_t i = 0; i < numElements; i++) {
printf("%5zu: %d\n", i, f[i]);
}
}
//---
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
HipTest::parseStandardArguments(argc, argv, false);
parseMyArguments(argc, argv);
typedef Streamer<int> IntStreamer;
std::vector<IntStreamer *> streamers;
std::vector<IntStreamer*> streamers;
size_t numElements = N;
int *expected_H = (int*)malloc(numElements*sizeof(int));
int* expected_H = (int*)malloc(numElements * sizeof(int));
auto nullStreamer = new IntStreamer(numElements, true);
// Expected resultr - last streamer runs vectorADDRepeat, then nullstreamer adds lastStreamer->_C_d + lastStreamer->_C_d
for (size_t i=0; i<numElements; i++) {
expected_H[i] = ((nullStreamer->_A_h[i])*p_repeat + (nullStreamer->_B_h[i]) * p_repeat) *2;
// Expected resultr - last streamer runs vectorADDRepeat, then nullstreamer adds
// lastStreamer->_C_d + lastStreamer->_C_d
for (size_t i = 0; i < numElements; i++) {
expected_H[i] =
((nullStreamer->_A_h[i]) * p_repeat + (nullStreamer->_B_h[i]) * p_repeat) * 2;
}
for (int i=0; i<p_streams; i++) {
IntStreamer * s = new IntStreamer(numElements);
for (int i = 0; i < p_streams; i++) {
IntStreamer* s = new IntStreamer(numElements);
streamers.push_back(s);
}
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
for (int s=1; s<p_streams; s++) {
if (p_tests & (1<<s)) {
printf ("==> Test %x runAsnc, #streams=%d\n", (1<<s), s);
for (int s = 1; s < p_streams; s++) {
if (p_tests & (1 << s)) {
printf("==> Test %x runAsnc, #streams=%d\n", (1 << s), s);
nullStreamer->reset();
for (int i=0; i<s; i++) {
for (int i = 0; i < s; i++) {
streamers[i]->enqueAsync();
}
auto lastStreamer = streamers[s - 1];
// Dispatch to NULL stream, should wait for prior async activity to complete before beginning:
hipLaunchKernel(
vectorADDRepeat,
dim3(blocks),
dim3(threadsPerBlock),
0,
0/*nullstream*/,
static_cast<const int*>(lastStreamer->_C_d),
static_cast<const int*>(lastStreamer->_C_d),
nullStreamer->_C_d,
numElements,
1/*repeat*/);
// Dispatch to NULL stream, should wait for prior async activity to complete before
// beginning:
hipLaunchKernel(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0,
0 /*nullstream*/, static_cast<const int*>(lastStreamer->_C_d),
static_cast<const int*>(lastStreamer->_C_d), nullStreamer->_C_d,
numElements, 1 /*repeat*/);
if (p_db) {
@@ -263,12 +227,12 @@ int main(int argc, char *argv[])
}
for (int s=1; s<p_streams; s+=2) {
unsigned tmask = (0x10000 | (1<<s));
for (int s = 1; s < p_streams; s += 2) {
unsigned tmask = (0x10000 | (1 << s));
if (p_tests & tmask) {
nullStreamer->reset();
printf ("==> Test %x runAsnc-odd-only, #streams=%d\n", tmask, s);
for (int i=0; i<s; i++) {
printf("==> Test %x runAsnc-odd-only, #streams=%d\n", tmask, s);
for (int i = 0; i < s; i++) {
// RUn just odd streams so we have some empty ones to examine/optimize:
if (i & 0x1) {
streamers[i]->enqueAsync();
@@ -276,18 +240,12 @@ int main(int argc, char *argv[])
}
auto lastStreamer = streamers[s - 1];
// Dispatch to NULL stream, should wait for prior async activity to complete before beginning:
hipLaunchKernel(
vectorADDRepeat,
dim3(blocks),
dim3(threadsPerBlock),
0,
0/*nullstream*/,
static_cast<const int*>(lastStreamer->_C_d),
static_cast<const int*>(lastStreamer->_C_d),
nullStreamer->_C_d,
numElements,
1/*repeat*/);
// Dispatch to NULL stream, should wait for prior async activity to complete before
// beginning:
hipLaunchKernel(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0,
0 /*nullstream*/, static_cast<const int*>(lastStreamer->_C_d),
static_cast<const int*>(lastStreamer->_C_d), nullStreamer->_C_d,
numElements, 1 /*repeat*/);
nullStreamer->D2H();
@@ -298,26 +256,28 @@ int main(int argc, char *argv[])
}
// Expected resultr - last streamer runs vectorADDRepeat
for (size_t i=0; i<numElements; i++) {
expected_H[i] = ((nullStreamer->_A_h[i])*p_repeat + (nullStreamer->_B_h[i]) * p_repeat);
for (size_t i = 0; i < numElements; i++) {
expected_H[i] = ((nullStreamer->_A_h[i]) * p_repeat + (nullStreamer->_B_h[i]) * p_repeat);
}
if (p_tests & 0x20000) {
assert (p_streams >=2); // need a couple streams in order to run this test.
assert(p_streams >= 2); // need a couple streams in order to run this test.
nullStreamer->reset();
printf ("\n==> Test hipStreamSynchronize with defaultStream \n");
printf("\n==> Test hipStreamSynchronize with defaultStream \n");
// Enqueue a long-running job to stream1
streamers[0]->enqueAsync();
// Check to see if synchronizing on a null stream synchronizes all other streams or just the null stream.
// This function follows null stream semantics and will wait for all other blocking streams before returning.
// This will wait on the host
// Check to see if synchronizing on a null stream synchronizes all other streams or just the
// null stream. This function follows null stream semantics and will wait for all other
// blocking streams before returning. This will wait on the host
HIPCHECK(hipStreamSynchronize(0));
// Copy with stream1, this could go async if the streamSync doesn't synchronize ALL the streams.
HIPCHECK(hipMemcpyAsync(streamers[0]->_C_h, streamers[0]->_C_d, streamers[0]->_numElements*sizeof(int), hipMemcpyDeviceToHost, streamers[1]->_stream));
// Copy with stream1, this could go async if the streamSync doesn't synchronize ALL the
// streams.
HIPCHECK(hipMemcpyAsync(streamers[0]->_C_h, streamers[0]->_C_d,
streamers[0]->_numElements * sizeof(int), hipMemcpyDeviceToHost,
streamers[1]->_stream));
HIPCHECK(hipDeviceSynchronize());
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStream.h
+64 -68
View File
@@ -24,97 +24,93 @@ THE SOFTWARE.
#define NUM_STREAMS 4
/*
* H2H - 1
* H2D - 2
* KER - 3
* D2D - 4
* D2H - 5
*/
* 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 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 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 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 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 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 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 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>
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 = threadIdx.x + blockIdx.x * blockDim.x;
In[tx] = In[tx] + 1;
template <typename T>
__global__ void Inc(hipLaunchParm lp, T* In) {
int tx = threadIdx.x + blockIdx.x * blockDim.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, 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 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;
}
template <typename T>
void setArray(T* Array, int N, T val) {
for (int i = 0; i < N; i++) {
Array[i] = val;
}
}
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp
+15 -19
View File
@@ -29,16 +29,15 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "test_common.h"
#ifdef __HIP_PLATFORM_HCC__
#ifdef __HIP_PLATFORM_HCC__
#define HIPRT_CB
#endif
__global__ void vector_square(float *C_d, float *A_d, size_t N)
{
__global__ void vector_square(float* C_d, float* A_d, size_t N) {
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
size_t stride = blockDim.x * gridDim.x;
for (size_t i=offset; i<N; i+=stride) {
for (size_t i = offset; i < N; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
@@ -46,31 +45,28 @@ __global__ void vector_square(float *C_d, float *A_d, size_t N)
float *A_h, *C_h;
bool cbDone = false;
static void HIPRT_CB Callback(hipStream_t stream, hipError_t status, void *userData)
{
for (size_t i=0; i<N; i++) {
static void HIPRT_CB Callback(hipStream_t stream, hipError_t status, void* userData) {
for (size_t i = 0; i < N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
warn("Data mismatch %zu", i);
}
}
printf ("PASSED!\n");
printf("PASSED!\n");
cbDone = true;
}
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
float *A_d, *C_d;
size_t Nbytes = N * sizeof(float);
A_h = (float*)malloc(Nbytes);
HIPCHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
HIPCHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
C_h = (float*)malloc(Nbytes);
HIPCHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
HIPCHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess);
// Fill with Phi + i
for (size_t i=0; i<N; i++)
{
A_h[i] = 1.618f + i;
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
}
HIPCHECK(hipMalloc(&A_d, Nbytes));
@@ -83,11 +79,11 @@ int main(int argc, char *argv[])
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0, mystream, C_d, A_d, N);
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0, mystream, C_d, A_d,
N);
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
HIPCHECK(hipStreamAddCallback(mystream, Callback, NULL, 0));
while(!cbDone)
std::this_thread::sleep_for(std::chrono::milliseconds(10));
while (!cbDone) std::this_thread::sleep_for(std::chrono::milliseconds(10));
}
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp
+3 -4
View File
@@ -26,18 +26,17 @@ THE SOFTWARE.
#include "test_common.h"
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
hipStream_t stream;
unsigned int flags;
HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
HIPCHECK(hipStreamGetFlags(stream, &flags));
HIPASSERT(flags == 0);
HIPASSERT(flags == 0);
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
HIPCHECK(hipStreamGetFlags(stream, &flags));
HIPASSERT(flags == 1);
HIPASSERT(flags == 1);
HIPCHECK(hipStreamDestroy(stream));
passed();
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamL5.cpp
+554 -560
View File
File diff suppressed because it is too large Load Diff
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamSync2.cpp
+33 -40
View File
@@ -36,7 +36,7 @@ enum SyncMode {
};
const char *syncModeString(int syncMode) {
const char* syncModeString(int syncMode) {
switch (syncMode) {
case syncNone:
return "syncNone";
@@ -56,9 +56,8 @@ const char *syncModeString(int syncMode) {
};
void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expectMismatch)
{
void test(unsigned testMask, int* C_d, int* C_h, int64_t numElements, SyncMode syncMode,
bool expectMismatch) {
// This test sends a long-running kernel to the null stream, then tests to see if the
// specified synchronization technique is effective.
//
@@ -75,20 +74,21 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode s
if (!(testMask & p_tests)) {
return;
}
printf ("\ntest 0x%02x: syncMode=%s expectMismatch=%d\n",
testMask, syncModeString(syncMode), expectMismatch);
printf("\ntest 0x%02x: syncMode=%s expectMismatch=%d\n", testMask, syncModeString(syncMode),
expectMismatch);
size_t sizeBytes = numElements * sizeof(int);
int count =100;
int count = 100;
int init0 = 0;
HIPCHECK(hipMemset(C_d, init0, sizeBytes));
for (int i=0; i<numElements; i++) {
C_h[i] = -1; // initialize
for (int i = 0; i < numElements; i++) {
C_h[i] = -1; // initialize
}
hipStream_t otherStream = 0;
unsigned flags = (syncMode == syncMarkerThenOtherNonBlockingStream) ? hipStreamNonBlocking : hipStreamDefault;
unsigned flags = (syncMode == syncMarkerThenOtherNonBlockingStream) ? hipStreamNonBlocking
: hipStreamDefault;
HIPCHECK(hipStreamCreateWithFlags(&otherStream, flags));
hipEvent_t stop, otherStreamEvent;
HIPCHECK(hipEventCreate(&stop));
@@ -97,17 +97,9 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode s
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*/,
static_cast<const int*>(C_d),
C_h,
numElements,
count);
HIPCHECK(hipEventRecord(stop, 0/*default*/));
hipLaunchKernelGGL(HipTest::addCountReverse, dim3(blocks), dim3(threadsPerBlock), 0,
0 /*stream*/, static_cast<const int*>(C_d), C_h, numElements, count);
HIPCHECK(hipEventRecord(stop, 0 /*default*/));
switch (syncMode) {
case syncNone:
@@ -137,26 +129,26 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode s
hipError_t done = hipEventQuery(stop);
if (expectMismatch) {
assert (done == hipErrorNotReady);
assert(done == hipErrorNotReady);
} else {
assert (done == hipSuccess);
assert(done == hipSuccess);
}
int mismatches = 0;
int expected = init0 + count;
for (int i=0; i<numElements; i++) {
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);
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);
assert(mismatches > 0);
}
@@ -166,15 +158,16 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode s
HIPCHECK(hipDeviceSynchronize());
printf ("test: OK - %d mismatches (%6.2f%%)\n", mismatches, ((double)(mismatches)*100.0)/numElements);
printf("test: OK - %d mismatches (%6.2f%%)\n", mismatches,
((double)(mismatches)*100.0) / numElements);
}
void runTests(int64_t numElements)
{
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);
printf("\n\ntest: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes,
sizeBytes / 1024.0 / 1024.0);
int *C_h, *C_d;
@@ -183,18 +176,19 @@ void runTests(int64_t numElements)
{
test (0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/);
test (0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/);
test (0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/);
test (0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/);
test(0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/);
test(0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/);
test(0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/);
test(0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/);
// Sending a marker to to null stream may synchronize the otherStream
// - other created with hipStreamNonBlocking=0 : synchronization, should match
// - other created with hipStreamNonBlocking=1 : no synchronization, may mismatch
test (0x10, C_d, C_h, numElements, syncMarkerThenOtherStream, false /*expectMismatch*/);
test(0x10, C_d, C_h, numElements, syncMarkerThenOtherStream, false /*expectMismatch*/);
// TODO - review why this test seems flaky
//test (0x20, C_d, C_h, numElements, syncMarkerThenOtherNonBlockingStream, true /*expectMismatch*/);
// test (0x20, C_d, C_h, numElements, syncMarkerThenOtherNonBlockingStream, true
// /*expectMismatch*/);
}
@@ -203,8 +197,7 @@ void runTests(int64_t numElements)
}
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
// Can' destroy the default stream:// TODO - move to another test
HIPCHECK_API(hipStreamDestroy(0), hipErrorInvalidResourceHandle);
projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp
+145 -176
View File
@@ -34,69 +34,63 @@ unsigned p_db = 0;
unsigned p_count = 100;
//------
// Structure for one stream;
template <typename T>
class Streamer {
#define COMMAND_ADD_FORWARD 0
#define COMMAND_ADD_REVERSE 1
#define COMMAND_COPY 2
#define COMMAND_COPY 2
public:
Streamer(int deviceId, T *input, size_t numElements, int commandType);
public:
Streamer(int deviceId, T* input, size_t numElements, int commandType);
~Streamer();
void runAsyncAfter(Streamer<T> *depStreamer, bool waitSameStream=false);
void runAsyncAfter(Streamer<T>* depStreamer, bool waitSameStream = false);
void runAsyncWaitSameStream();
void queryUntilComplete();
size_t check(int streamerNum, T initValue, T expectedOffset, bool expectPass=true);
size_t check(int streamerNum, T initValue, T expectedOffset, bool expectPass = true);
void copyToHost(hipStream_t copyStream);
hipEvent_t event() { return _event; };
int deviceId() const { return _deviceId; };
size_t mismatchCount() const { return _mismatchCount; };
T *C_d() { return _C_d; };
T* C_d() { return _C_d; };
// How much does this streamer add to A[i] after running runAsyncAfter
int expectedAdd() const { return (_commandType == COMMAND_COPY) ? 0 : p_count; };
int _commandType; // 0=addReverse, 1=addFwd, 2=move
private:
int _commandType; // 0=addReverse, 1=addFwd, 2=move
private:
T* _C_h;
T *_C_h;
T *_preA_d; // if input is on another device, this is pointer to that memory.
T *_A_d;
T *_C_d;
T* _preA_d; // if input is on another device, this is pointer to that memory.
T* _A_d;
T* _C_d;
hipStream_t _stream;
hipEvent_t _event;
hipEvent_t _event;
int _deviceId;
size_t _numElements;
int _deviceId;
size_t _numElements;
size_t _mismatchCount;
size_t _mismatchCount;
};
template <typename T>
Streamer<T>::Streamer(int deviceId, T * A_d, size_t numElements, int commandType) :
_preA_d(NULL),
_A_d(A_d),
_deviceId(deviceId),
_numElements(numElements),
_commandType(commandType)
{
Streamer<T>::Streamer(int deviceId, T* A_d, size_t numElements, int commandType)
: _preA_d(NULL),
_A_d(A_d),
_deviceId(deviceId),
_numElements(numElements),
_commandType(commandType) {
size_t sizeElements = numElements * sizeof(int);
//if (commandType == 0) _commandType = 1; // TODO - remove me
// if (commandType == 0) _commandType = 1; // TODO - remove me
HIPCHECK(hipSetDevice(_deviceId));
@@ -106,7 +100,7 @@ Streamer<T>::Streamer(int deviceId, T * A_d, size_t numElements, int commandType
if (attr.device != deviceId) {
// source is on another device, we will need to copy later.
// So save original source pointer and allocate local space.
printf ("info: source for streamer on another device, will insert memcpy\n");
printf("info: source for streamer on another device, will insert memcpy\n");
_preA_d = A_d;
HIPCHECK(hipMalloc(&_A_d, sizeElements));
HIPCHECK(hipMemset(_A_d, -3, sizeElements));
@@ -120,18 +114,14 @@ Streamer<T>::Streamer(int deviceId, T * A_d, size_t numElements, int commandType
HIPCHECK(hipStreamCreate(&_stream));
HIPCHECK(hipEventCreate(&_event));
};
template <typename T>
Streamer<T>::~Streamer()
{
Streamer<T>::~Streamer() {
HIPCHECK(hipSetDevice(_deviceId));
printf ("info: ~Streamer\n");
printf("info: ~Streamer\n");
if (_preA_d) {
HIPCHECK(hipFree(_preA_d));
}
@@ -144,11 +134,11 @@ Streamer<T>::~Streamer()
template <typename T>
void Streamer<T>::runAsyncAfter(Streamer<T> *depStreamer, bool waitSameStream)
{
void Streamer<T>::runAsyncAfter(Streamer<T>* depStreamer, bool waitSameStream) {
HIPCHECK(hipSetDevice(_deviceId));
if (p_db) {
printf ("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0);
printf("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements,
_numElements * sizeof(T) / 1024.0 / 1024.0);
}
if (depStreamer) {
@@ -157,131 +147,117 @@ void Streamer<T>::runAsyncAfter(Streamer<T> *depStreamer, bool waitSameStream)
if (_preA_d) {
// _preA_d is on another device, so copy to local device so kernel can access it:
HIPCHECK(hipMemcpyAsync(_A_d, _preA_d, _numElements * sizeof(T), hipMemcpyDeviceToDevice, _stream));
HIPCHECK(hipMemcpyAsync(_A_d, _preA_d, _numElements * sizeof(T), hipMemcpyDeviceToDevice,
_stream));
}
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
if (_commandType == COMMAND_ADD_REVERSE) {
hipLaunchKernelGGL(
HipTest::addCountReverse,
dim3(blocks),
dim3(threadsPerBlock),
0,
_stream,
static_cast<const T*>(_A_d),
_C_d,
static_cast<int64_t>(_numElements),
static_cast<int>(p_count));
hipLaunchKernelGGL(HipTest::addCountReverse, dim3(blocks), dim3(threadsPerBlock), 0,
_stream, static_cast<const T*>(_A_d), _C_d,
static_cast<int64_t>(_numElements), static_cast<int>(p_count));
} else if (_commandType == COMMAND_ADD_FORWARD) {
hipLaunchKernelGGL(
HipTest::addCount,
dim3(blocks),
dim3(threadsPerBlock),
0,
_stream,
static_cast<const T*>(_A_d),
_C_d,
_numElements,
static_cast<int>(p_count));
hipLaunchKernelGGL(HipTest::addCount, dim3(blocks), dim3(threadsPerBlock), 0, _stream,
static_cast<const T*>(_A_d), _C_d, _numElements,
static_cast<int>(p_count));
} else if (_commandType == COMMAND_COPY) {
HIPCHECK(hipMemcpyAsync(_C_d, _A_d, _numElements * sizeof(T), hipMemcpyDeviceToDevice, _stream));
HIPCHECK(
hipMemcpyAsync(_C_d, _A_d, _numElements * sizeof(T), hipMemcpyDeviceToDevice, _stream));
} else {
assert(0); // bad command type
assert(0); // bad command type
}
HIPCHECK(hipEventRecord(_event, _stream));
if (waitSameStream) {
HIPCHECK(hipStreamWaitEvent(_stream, _event, 0)); // this is essentially a no-op, but make sure it doesn't crash
HIPCHECK(hipStreamWaitEvent(
_stream, _event, 0)); // this is essentially a no-op, but make sure it doesn't crash
}
}
template <typename T>
void Streamer<T>::queryUntilComplete()
{
void Streamer<T>::queryUntilComplete() {
HIPCHECK(hipSetDevice(_deviceId));
int numQueries = 0;
hipError_t e = hipSuccess;
do {
numQueries++;
e = hipStreamQuery(_stream);
} while (e != hipSuccess) ;
} while (e != hipSuccess);
printf ("info: hipStreamQuery completed after %d queries\n", numQueries);
printf("info: hipStreamQuery completed after %d queries\n", numQueries);
};
// If copyStream is !nullptr it is used for the copy.
template <typename T>
void Streamer<T>::copyToHost(hipStream_t copyStream)
{
void Streamer<T>::copyToHost(hipStream_t copyStream) {
if (p_db) {
printf ("db: copy back to host\n");
printf("db: copy back to host\n");
}
HIPCHECK(hipSetDevice(_deviceId));
HIPCHECK(hipMemcpyAsync(_C_h, _C_d, _numElements*sizeof(T), hipMemcpyDeviceToHost, copyStream ? copyStream : _stream));
HIPCHECK(hipStreamSynchronize(copyStream ? copyStream:_stream));
HIPCHECK(hipMemcpyAsync(_C_h, _C_d, _numElements * sizeof(T), hipMemcpyDeviceToHost,
copyStream ? copyStream : _stream));
HIPCHECK(hipStreamSynchronize(copyStream ? copyStream : _stream));
}
template <typename T>
size_t Streamer<T>::check(int streamerNum, T initValue, T expectedOffset, bool expectPass)
{
size_t Streamer<T>::check(int streamerNum, T initValue, T expectedOffset, bool expectPass) {
T expected = initValue + expectedOffset;
if (p_db) {
printf ("db: check\n");
printf("db: check\n");
}
_mismatchCount = 0;
for (size_t i=0; i<_numElements; i++) {
for (size_t i = 0; i < _numElements; i++) {
if (_C_h[i] != expected) {
_mismatchCount++;
if (expectPass) {
fprintf(stderr, "for streamer:%d _C_h[%zu] (%d) != expected(%d)\n", streamerNum, i, _C_h[i], expected);
fprintf(stderr, "for streamer:%d _C_h[%zu] (%d) != expected(%d)\n", streamerNum,
i, _C_h[i], expected);
if (_mismatchCount > 10) {
failed("for streamer:%d _C_h[%zu] (%d) != expected(%d)\n", streamerNum, i, _C_h[i], expected);
failed("for streamer:%d _C_h[%zu] (%d) != expected(%d)\n", streamerNum, i,
_C_h[i], expected);
}
}
}
}
if (!expectPass && (_mismatchCount ==0)) {
// the test should run kernels long enough that if we don't correctly wait for them to finish then an error is reported.
//failed("for streamer:%d we expected inavalid synchronization to lead to mismatch but none was detected. Increase --N to sensitize sync.\n", streamerNum);
if (!expectPass && (_mismatchCount == 0)) {
// the test should run kernels long enough that if we don't correctly wait for them to
// finish then an error is reported.
// failed("for streamer:%d we expected inavalid synchronization to lead to mismatch but
// none was detected. Increase --N to sensitize sync.\n", streamerNum);
}
return _mismatchCount;
}
//---
//Parse arguments specific to this test.
void parseMyArguments(int argc, char *argv[])
{
N = 64*1024*1024;
// Parse arguments specific to this test.
void parseMyArguments(int argc, char* argv[]) {
N = 64 * 1024 * 1024;
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];
const char* arg = argv[i];
if (!strcmp(arg, "--streams")) {
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) {
failed("Bad streams argument");
failed("Bad streams argument");
}
} else if (!strcmp(arg, "--count")) {
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_count)) {
failed("Bad count argument");
failed("Bad count argument");
}
} else if (!strcmp(arg, "--db")) {
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_db)) {
failed("Bad db argument");
failed("Bad db argument");
}
} else {
failed("Bad argument '%s'", arg);
@@ -293,80 +269,73 @@ void parseMyArguments(int argc, char *argv[])
typedef Streamer<int> IntStreamer;
void runStreamerLoop(std::vector<IntStreamer *> &streamers)
{
for (int i=0; i<streamers.size(); i++) {
streamers[i]->runAsyncAfter(i ? streamers[i-1] : NULL);
void runStreamerLoop(std::vector<IntStreamer*>& streamers) {
for (int i = 0; i < streamers.size(); i++) {
streamers[i]->runAsyncAfter(i ? streamers[i - 1] : NULL);
}
}
void checkAll(int initValue, std::vector<IntStreamer *> &streamers, std::vector<hipStream_t> &sideStreams, bool expectPass=true)
{
size_t mismatchCount=0;
void checkAll(int initValue, std::vector<IntStreamer*>& streamers,
std::vector<hipStream_t>& sideStreams, bool expectPass = true) {
size_t mismatchCount = 0;
// Copy in reverse order to catch anything not yet finished...
for (int i=streamers.size()-1; i>=0; i--) {
streamers[i]->copyToHost(sideStreams.empty() ? NULL : sideStreams[streamers[i]->deviceId()]);
for (int i = streamers.size() - 1; i >= 0; i--) {
streamers[i]->copyToHost(sideStreams.empty() ? NULL
: sideStreams[streamers[i]->deviceId()]);
}
int expected = 0;
// Check in forward order so we can find first mismatch:
for (int i=0; i<streamers.size(); i++) {
for (int i = 0; i < streamers.size(); i++) {
expected += streamers[i]->expectedAdd();
mismatchCount += streamers[i]->check(i+1, initValue, expected, expectPass);
mismatchCount += streamers[i]->check(i + 1, initValue, expected, expectPass);
}
if (!expectPass && (mismatchCount==0)) {
// the test should run kernels long enough that if we don't correctly wait for them to finish then an error is reported.
failed("we expected inavalid synchronization to lead to mismatch but none was detected. Increase --count to sensitize sync.\n");
if (!expectPass && (mismatchCount == 0)) {
// the test should run kernels long enough that if we don't correctly wait for them to
// finish then an error is reported.
failed(
"we expected inavalid synchronization to lead to mismatch but none was detected. "
"Increase --count to sensitize sync.\n");
}
}
#define RUN_SYNC_TEST(_enableBit, _streamers, _sync, _expectPass)\
if (p_tests & (_enableBit)) {\
printf ("==> Test %02x runAsyncAfter sync=%s\n", (_enableBit), #_sync);\
runStreamerLoop(_streamers);\
(_sync);\
checkAll (initValue, _streamers, sideStreams, _expectPass);\
#define RUN_SYNC_TEST(_enableBit, _streamers, _sync, _expectPass) \
if (p_tests & (_enableBit)) { \
printf("==> Test %02x runAsyncAfter sync=%s\n", (_enableBit), #_sync); \
runStreamerLoop(_streamers); \
(_sync); \
checkAll(initValue, _streamers, sideStreams, _expectPass); \
}
//---
// A family of sync functions which somehow wait for inflight activity to finish:
void sync_none(void) {};
void sync_none(void){};
void sync_allDevices(int numDevices)
{
for (int d=0; d<numDevices; d++) {
void sync_allDevices(int numDevices) {
for (int d = 0; d < numDevices; d++) {
HIPCHECK(hipSetDevice(d));
HIPCHECK(hipDeviceSynchronize());
}
}
void sync_queryAllUntilComplete(std::vector<IntStreamer *> streamers)
{
for (int i=streamers.size()-1; i>=0; i--) {
void sync_queryAllUntilComplete(std::vector<IntStreamer*> streamers) {
for (int i = streamers.size() - 1; i >= 0; i--) {
streamers[i]->queryUntilComplete();
};
}
void sync_streamWaitEvent(hipEvent_t lastEvent, int sideDeviceId, hipStream_t sideStream, bool waitHere)
{
void sync_streamWaitEvent(hipEvent_t lastEvent, int sideDeviceId, hipStream_t sideStream,
bool waitHere) {
HIPCHECK(hipSetDevice(sideDeviceId));
// wait on the last event in the stream of chained streamers:
@@ -379,31 +348,28 @@ void sync_streamWaitEvent(hipEvent_t lastEvent, int sideDeviceId, hipStream_t si
}
//---
int main(int argc, char *argv[])
{
int main(int argc, char* argv[]) {
HipTest::parseStandardArguments(argc, argv, false);
parseMyArguments(argc, argv);
size_t numElements = N;
size_t sizeElements = numElements * sizeof(int);
printf("info: sizeof arrays = %zu elements (%6.3f MB)\n", numElements, sizeElements / 1024.0/1024.0);
printf("info: sizeof arrays = %zu elements (%6.3f MB)\n", numElements,
sizeElements / 1024.0 / 1024.0);
printf("info: streams=%d count=%d\n", p_streams, p_count);
assert (sizeElements <= std::numeric_limits<int64_t>::max());
assert(sizeElements <= std::numeric_limits<int64_t>::max());
int initValue = 1000;
int * initArray_d, *initArray_h;
int *initArray_d, *initArray_h;
HIPCHECK(hipMalloc(&initArray_d, sizeElements));
HIPCHECK(hipHostMalloc(&initArray_h, sizeElements));
for (size_t i=0; i<numElements; i++) {
for (size_t i = 0; i < numElements; i++) {
initArray_h[i] = initValue;
}
HIPCHECK(hipMemcpy(initArray_d, initArray_h, sizeElements, hipMemcpyHostToDevice));
@@ -411,31 +377,29 @@ int main(int argc, char *argv[])
int numDevices;
HIPCHECK(hipGetDeviceCount(&numDevices));
numDevices = min(2, numDevices); // multi-GPU to 2 device.
numDevices = min(2, numDevices); // multi-GPU to 2 device.
std::vector<IntStreamer *> streamers;
std::vector<IntStreamer *> streamersDev0; // streamers for first device.
std::vector<IntStreamer*> streamers;
std::vector<IntStreamer*> streamersDev0; // streamers for first device.
for (int d=0; d<numDevices/*TODO*/; d++) {
for (int i=0; i<p_streams; i++) {
int command = (i%2) ? COMMAND_ADD_FORWARD : COMMAND_ADD_REVERSE;
IntStreamer * s = new IntStreamer(d, i ? streamers.back()->C_d() : initArray_d, numElements, command);
for (int d = 0; d < numDevices /*TODO*/; d++) {
for (int i = 0; i < p_streams; i++) {
int command = (i % 2) ? COMMAND_ADD_FORWARD : COMMAND_ADD_REVERSE;
IntStreamer* s =
new IntStreamer(d, i ? streamers.back()->C_d() : initArray_d, numElements, command);
streamers.push_back(s);
if (d==0) {
if (d == 0) {
streamersDev0.push_back(s);
}
}
}
// A sideband stream channel that is independent from above.
// Used to check to ensure the WaitEvent or other synchronization is working correctly since by default sideStream is
// asynchronous wrt the other streams.
// Used to check to ensure the WaitEvent or other synchronization is working correctly since by
// default sideStream is asynchronous wrt the other streams.
std::vector<hipStream_t> sideStreams;
for (int d=0; d<numDevices; d++) {
for (int d = 0; d < numDevices; d++) {
hipStream_t s;
HIPCHECK(hipStreamCreate(&s));
sideStreams.push_back(s);
@@ -444,27 +408,31 @@ int main(int argc, char *argv[])
// Tests on first GPU:
//
// This test has no synchronization - make sure it mismatches so we can ensure the other tests properyl prevent the mismatch:
// This test has no synchronization - make sure it mismatches so we can ensure the other tests
// properyl prevent the mismatch:
RUN_SYNC_TEST(0x01, streamersDev0, sync_none(), false);
RUN_SYNC_TEST(0x02, streamersDev0, sync_allDevices(numDevices), true);
RUN_SYNC_TEST(0x04, streamersDev0, sync_queryAllUntilComplete(streamersDev0), true);
RUN_SYNC_TEST(0x08, streamersDev0, sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false), true);
RUN_SYNC_TEST(0x02, streamersDev0, sync_allDevices(numDevices), true);
RUN_SYNC_TEST(0x04, streamersDev0, sync_queryAllUntilComplete(streamersDev0), true);
RUN_SYNC_TEST(0x08, streamersDev0,
sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false),
true);
if (numDevices > 1) {
// Sync on second device for activity running on device 0:
RUN_SYNC_TEST(0x10, streamersDev0, sync_streamWaitEvent(streamersDev0.back()->event(), 1, sideStreams[1], true), true);
RUN_SYNC_TEST(0x10, streamersDev0,
sync_streamWaitEvent(streamersDev0.back()->event(), 1, sideStreams[1], true),
true);
}
// Tests on all GPUs:
// RUN_SYNC_TEST(0x100, streamers, sync_streamWaitEvent(streamers.back()->event(), 0, sideStreams[0], false), true);
// RUN_SYNC_TEST(0x100, streamers, sync_streamWaitEvent(streamers.back()->event(), 0,
// sideStreams[0], false), true);
if (p_tests & 0x1000) {
printf ("==> Test 0x1000 simple null stream tests\n");
printf("==> Test 0x1000 simple null stream tests\n");
// try some null stream:
hipStreamQuery(0);
@@ -480,7 +448,7 @@ int main(int argc, char *argv[])
HIPCHECK(hipEventRecord(e1, s1))
HIPCHECK(hipStreamWaitEvent(hipStream_t(0), e1, 0/*flags*/));
HIPCHECK(hipStreamWaitEvent(hipStream_t(0), e1, 0 /*flags*/));
HIPCHECK(hipStreamDestroy(s1));
HIPCHECK(hipEventDestroy(e1));
@@ -493,38 +461,39 @@ int main(int argc, char *argv[])
HIPCHECK(hipEventRecord(e1, hipStream_t(0)))
HIPCHECK(hipStreamWaitEvent(s1, e1, 0/*flags*/));
HIPCHECK(hipStreamWaitEvent(s1, e1, 0 /*flags*/));
HIPCHECK(hipStreamDestroy(s1));
HIPCHECK(hipEventDestroy(e1));
}
}
// Insert small wrinkle here, insert a wait on event just recorded, all in the same stream.
if (p_tests & 0x2000) {
printf ("==> Test 0x2000 runAsyncWaitSameStream\n");
for (int i=0; i<streamersDev0.size(); i++) {
streamersDev0[i]->runAsyncAfter(i ? streamersDev0[i-1] : NULL, true/*waitSameStream*/);
printf("==> Test 0x2000 runAsyncWaitSameStream\n");
for (int i = 0; i < streamersDev0.size(); i++) {
streamersDev0[i]->runAsyncAfter(i ? streamersDev0[i - 1] : NULL,
true /*waitSameStream*/);
}
sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false);
checkAll (initValue, streamersDev0, sideStreams);
checkAll(initValue, streamersDev0, sideStreams);
}
// Change Adds to copies to stimulate different case with event followign copy:
for (auto &s : streamers) {
if (s->_commandType == COMMAND_ADD_FORWARD)
s->_commandType = COMMAND_COPY;
for (auto& s : streamers) {
if (s->_commandType == COMMAND_ADD_FORWARD) s->_commandType = COMMAND_COPY;
}
if (p_tests & 0x4000 ) {
printf ("test: %x alternating memcpy/count-reverse followed by event\n", p_tests);
RUN_SYNC_TEST(0x4000, streamersDev0, sync_queryAllUntilComplete(streamersDev0), true);
RUN_SYNC_TEST(0x8000, streamersDev0, sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false), true);
if (p_tests & 0x4000) {
printf("test: %x alternating memcpy/count-reverse followed by event\n", p_tests);
RUN_SYNC_TEST(0x4000, streamersDev0, sync_queryAllUntilComplete(streamersDev0), true);
RUN_SYNC_TEST(0x8000, streamersDev0,
sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false),
true);
}