changed __global__ attribute
1. Moved around tests and added them to HIT
Change-Id: I5d75280c42a5af852670ebabc7305ee56721ec7b
[ROCm/clr commit: f009f3533d]
This commit is contained in:
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
@@ -50,7 +50,7 @@ void help(char *argv[])
|
||||
};
|
||||
|
||||
|
||||
static hipError_t myHipMemcpy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream, bool async)
|
||||
static hipError_t myHipMemcpy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream, bool async)
|
||||
{
|
||||
if (async) {
|
||||
hipError_t e = hipMemcpyAsync(dest, src, sizeBytes, kind, stream);
|
||||
@@ -78,7 +78,7 @@ void parseMyArguments(int argc, char *argv[])
|
||||
p_mirrorPeers = true;
|
||||
} else if (!strcmp(arg, "--peerDevice")) {
|
||||
if (++i >= argc || !HipTest::parseInt(argv[i], &p_peerDevice)) {
|
||||
failed("Bad peerDevice argument");
|
||||
failed("Bad peerDevice argument");
|
||||
}
|
||||
} else {
|
||||
failed("Bad argument '%s'", arg);
|
||||
@@ -101,7 +101,7 @@ void syncBothDevices()
|
||||
|
||||
|
||||
// Sets globals g_currentDevice, g_peerDevice
|
||||
void setupPeerTests()
|
||||
void setupPeerTests()
|
||||
{
|
||||
int deviceCnt;
|
||||
|
||||
@@ -159,17 +159,17 @@ void enablePeerFirst(bool useAsyncCopy)
|
||||
// allocate and initialize memory on device0
|
||||
HIPCHECK (hipSetDevice(g_currentDevice));
|
||||
HIPCHECK (hipMalloc(&A_d0, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_d0, memsetval, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_d0, memsetval, Nbytes) );
|
||||
|
||||
// allocate and initialize memory on peer device
|
||||
HIPCHECK (hipSetDevice(g_peerDevice));
|
||||
HIPCHECK (hipMalloc(&A_d1, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) );
|
||||
|
||||
|
||||
|
||||
// Device0 push to device1, using P2P:
|
||||
// NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
|
||||
// NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
|
||||
// a host staging copy for the P2P access.
|
||||
HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
|
||||
HIPCHECK (myHipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault, 0/*stream*/, useAsyncCopy)); // This is P2P copy.
|
||||
@@ -177,7 +177,7 @@ void enablePeerFirst(bool useAsyncCopy)
|
||||
// Copy data back to host:
|
||||
// Have to wait for previous operation to finish, since we are switching to another one:
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
|
||||
HIPCHECK (hipSetDevice(g_peerDevice));
|
||||
HIPCHECK (myHipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost, 0/*stream*/, useAsyncCopy));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
@@ -215,12 +215,12 @@ void allocMemoryFirst(bool useAsyncCopy)
|
||||
// allocate and initialize memory on device0
|
||||
HIPCHECK (hipSetDevice(g_currentDevice));
|
||||
HIPCHECK (hipMalloc(&A_d0, Nbytes) );
|
||||
HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) );
|
||||
HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) );
|
||||
|
||||
// allocate and initialize memory on peer device
|
||||
HIPCHECK (hipSetDevice(g_peerDevice));
|
||||
HIPCHECK (hipMalloc(&A_d1, Nbytes) );
|
||||
HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) );
|
||||
HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) );
|
||||
|
||||
|
||||
//---
|
||||
@@ -268,7 +268,7 @@ void allocMemoryFirst(bool useAsyncCopy)
|
||||
// Test which tests peer H2D copy - ie: copy-engine=1, dst=1, src=0 (Host)
|
||||
// A_d0 is pinned host on dev0 (this)
|
||||
// A_d1 is device memory on dev1 (peer)
|
||||
//
|
||||
//
|
||||
void testPeerHostToDevice(bool useAsyncCopy)
|
||||
{
|
||||
printf ("\n==testing: %s useAsyncCopy=%d\n", __func__, useAsyncCopy);
|
||||
@@ -299,12 +299,12 @@ void testPeerHostToDevice(bool useAsyncCopy)
|
||||
// allocate and initialize memory on device0
|
||||
HIPCHECK (hipSetDevice(g_currentDevice));
|
||||
HIPCHECK (hipHostMalloc(&A_host_d0, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_host_d0, memsetval, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_host_d0, memsetval, Nbytes) );
|
||||
|
||||
// allocate and initialize memory on peer device
|
||||
HIPCHECK (hipSetDevice(g_peerDevice));
|
||||
HIPCHECK (hipMalloc(&A_d1, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) );
|
||||
HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) );
|
||||
|
||||
bool firstAsyncCopy = useAsyncCopy; /*TODO - should be useAsyncCopy*/
|
||||
|
||||
@@ -313,17 +313,17 @@ void testPeerHostToDevice(bool useAsyncCopy)
|
||||
|
||||
|
||||
// Device0 push to device1, using P2P:
|
||||
// NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
|
||||
// NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a
|
||||
// a host staging copy for the P2P access.
|
||||
if (p_memcpyWithPeer) {
|
||||
// p_memcpyWithPeer=1 case is HostToDevice.
|
||||
// if p_mirrorPeers = 1, this is accelerated copy over PCIe.
|
||||
// if p_mirrorPeers = 1, this is accelerated copy over PCIe.
|
||||
// if p_mirrorPeers = 0, this should fall back to host (because peer can't see A_host_d0)
|
||||
HIPCHECK (hipSetDevice(g_peerDevice));
|
||||
HIPCHECK (myHipMemcpy(A_d1, A_host_d0, Nbytes, hipMemcpyHostToDevice, 0/*stream*/, firstAsyncCopy)); // This is P2P copy.
|
||||
} else {
|
||||
// p_memcpyWithPeer=0 case is HostToDevice.
|
||||
// if p_mirrorPeers = 1, this is accelerated copy over PCIe.
|
||||
// if p_mirrorPeers = 1, this is accelerated copy over PCIe.
|
||||
// if p_mirrorPeers = 0, this should fall back to host (because device0 can't see A_d1)
|
||||
HIPCHECK (hipSetDevice(g_currentDevice));
|
||||
HIPCHECK (myHipMemcpy(A_d1, A_host_d0, Nbytes, hipMemcpyHostToDevice, 0/*stream*/, firstAsyncCopy)); // This is P2P copy.
|
||||
@@ -367,7 +367,7 @@ void simpleNegative()
|
||||
HIPASSERT( e == hipSuccess); // no error returned, it doesn't hurt to ask.
|
||||
HIPASSERT (canAccessPeer == 0); // but self is not a peer.
|
||||
|
||||
e = hipSuccess;
|
||||
e = hipSuccess;
|
||||
//---
|
||||
// Enable same device twice in a row:
|
||||
HIPCHECK(hipSetDevice(g_currentDevice));
|
||||
@@ -381,7 +381,7 @@ void simpleNegative()
|
||||
e =(hipDeviceDisablePeerAccess(g_peerDevice));
|
||||
HIPASSERT (e == hipErrorPeerAccessNotEnabled);
|
||||
|
||||
|
||||
|
||||
// More tests here:
|
||||
printf ("==done: %s\n\n", __func__);
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user