Skip to content

Commit

Permalink
Clean up trailing whitespace so as to reduce noise in ROCm#246.
Browse files Browse the repository at this point in the history
  • Loading branch information
@AlexVlx
AlexVlx committed Nov 8, 2017
1 parent 9fd6c18 commit d8e323d
Show file tree
Hide file tree
Showing 8 changed files with 86 additions and 86 deletions.
4 changes: 2 additions & 2 deletions include/hip/hcc_detail/hip_runtime.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ THE SOFTWARE.
// define HIP_ENABLE_PRINTF to enable printf
#ifdef HIP_ENABLE_PRINTF
#define HCC_ENABLE_ACCELERATOR_PRINTF 1
#endif
#endif

//---
// Remainder of this file only compiles with HCC
Expand Down Expand Up @@ -481,7 +481,7 @@ do {\
type* var = \
(type*)__get_dynamicgroupbaseptr(); \

#define HIP_DYNAMIC_SHARED_ATTRIBUTE
#define HIP_DYNAMIC_SHARED_ATTRIBUTE



Expand Down
18 changes: 9 additions & 9 deletions src/hip_memory.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -69,7 +69,7 @@ int sharePtr(void *ptr, ihipCtx_t *ctx, bool shareWithAll, unsigned hipFlags)

if (shareWithAll) {
hsa_status_t s = hsa_amd_agents_allow_access(g_deviceCnt+1, g_allAgents, NULL, ptr);
tprintf (DB_MEM, " allow access to CPU + all %d GPUs (shareWithAll)\n", g_deviceCnt);
tprintf (DB_MEM, " allow access to CPU + all %d GPUs (shareWithAll)\n", g_deviceCnt);
if (s != HSA_STATUS_SUCCESS) {
ret = -1;
}
Expand Down Expand Up @@ -126,7 +126,7 @@ void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool
if (HIP_INIT_ALLOC != -1) {
// TODO , dont' call HIP API directly here:
hipMemset(ptr, HIP_INIT_ALLOC, sizeBytes);
}
}

if (ptr != nullptr) {
int r = sharePtr(ptr, ctx, shareWithAll, hipFlags);
Expand Down Expand Up @@ -255,7 +255,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
hip_status = hipErrorMemoryAllocation;
}

}
}


return ihipLogStatus(hip_status);
Expand Down Expand Up @@ -288,10 +288,10 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
}


const unsigned supportedFlags = hipHostMallocPortable
| hipHostMallocMapped
| hipHostMallocWriteCombined
| hipHostMallocCoherent
const unsigned supportedFlags = hipHostMallocPortable
| hipHostMallocMapped
| hipHostMallocWriteCombined
| hipHostMallocCoherent
| hipHostMallocNonCoherent;


Expand All @@ -304,7 +304,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
hip_status = hipErrorInvalidValue;
} else {
auto device = ctx->getWriteableDevice();

unsigned amFlags = 0;
if (flags & hipHostMallocCoherent) {
amFlags = amHostCoherent;
Expand Down Expand Up @@ -585,7 +585,7 @@ hipError_t hipMalloc3DArray(hipArray_t *array,
hsa_ext_image_data_info_t imageInfo;
hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;

*ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0, alignment);

if (size && (*ptr == NULL)) {
Expand Down
20 changes: 10 additions & 10 deletions tests/src/runtimeApi/event/record_event.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, hipStream_
if (!(testMask & p_tests)) {
return;
}
printf ("\ntest 0x%3x: stream=%p waitStart=%d syncMode=%s\n",
printf ("\ntest 0x%3x: stream=%p waitStart=%d syncMode=%s\n",
testMask, stream, waitStart, syncModeString(syncMode));

size_t sizeBytes = numElements * sizeof(int);
Expand Down Expand Up @@ -85,8 +85,8 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, hipStream_
HIPCHECK(hipEventSynchronize(start));
}

hipError_t expectedStopError = hipSuccess;

hipError_t expectedStopError = hipSuccess;

// How to wait for the events to finish:
switch (syncMode) {
Expand All @@ -97,12 +97,12 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, hipStream_
HIPCHECK(hipStreamSynchronize(stream)); // wait for recording to finish...
break;
case syncStopEvent:
HIPCHECK(hipEventSynchronize(stop));
HIPCHECK(hipEventSynchronize(stop));
break;
default:
assert(0);
};


float t;

Expand All @@ -111,25 +111,25 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, hipStream_
failed ("start event not in expected state, was %d=%s\n", e, hipGetErrorName(e));
}

if (e == hipSuccess)
if (e == hipSuccess)
assert (t==0.0f);


// stop usually ready unless we skipped the synchronization (syncNone)
HIPCHECK_API(hipEventElapsedTime(&t, stop, stop), expectedStopError);
if (e == hipSuccess)
if (e == hipSuccess)
assert (t==0.0f);


e = hipEventElapsedTime(&t, start, stop);
HIPCHECK_API(e, expectedStopError);
if (expectedStopError == hipSuccess)
if (expectedStopError == hipSuccess)
assert (t>0.0f);
printf ("time=%6.2f error=%s\n", t, hipGetErrorName(e));

e = hipEventElapsedTime(&t, stop, start);
HIPCHECK_API(e, expectedStopError);
if (expectedStopError == hipSuccess)
if (expectedStopError == hipSuccess)
assert (t<0.0f);
printf ("negtime=%6.2f error=%s\n", t, hipGetErrorName(e));

Expand Down
66 changes: 33 additions & 33 deletions tests/src/runtimeApi/memory/hipMemcpy.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ class DeviceMemory

void offset(int offset) { _offset = offset; };
int offset() const { return _offset; };

private:
T * _A_d;
T* _B_d;
Expand All @@ -72,7 +72,7 @@ class DeviceMemory

template<typename T>
DeviceMemory<T>::DeviceMemory(size_t numElements)
: _maxNumElements(numElements),
: _maxNumElements(numElements),
_offset(0)
{
T ** np = nullptr;
Expand All @@ -93,7 +93,7 @@ DeviceMemory<T>::~DeviceMemory ()
HipTest::freeArrays (_A_d, _B_d, _C_d, np, np, np, 0);

HIPCHECK (hipFree(_C_dd));

_C_dd = NULL;
};

Expand Down Expand Up @@ -125,7 +125,7 @@ class HostMemory
T * A_hh;
T* B_hh;

bool _usePinnedHost;
bool _usePinnedHost;
private:
size_t _maxNumElements;

Expand Down Expand Up @@ -165,11 +165,11 @@ HostMemory<T>::HostMemory(size_t numElements, bool usePinnedHost)

template<typename T>
void
HostMemory<T>::reset(size_t numElements, bool full)
HostMemory<T>::reset(size_t numElements, bool full)
{
// Initialize the host data:
for (size_t i=0; i<numElements; i++) {
(A_hh)[i] = 1097.0 + i;
(A_hh)[i] = 1097.0 + i;
(B_hh)[i] = 1492.0 + i; // Phi

if (full) {
Expand Down Expand Up @@ -213,8 +213,8 @@ template <typename T>
void memcpytest2(DeviceMemory<T> *dmem, HostMemory<T> *hmem, size_t numElements, 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, offsets:dev:%+d host:+%d\n",
__func__,
printf ("test: %s<%s> size=%lu (%6.2fMB) usePinnedHost:%d, useHostToHost:%d, useDeviceToDevice:%d, useMemkindDefault:%d, offsets:dev:%+d host:+%d\n",
__func__,
TYPENAME(T),
sizeElements, sizeElements/1024.0/1024.0,
hmem->_usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault,
Expand Down Expand Up @@ -273,8 +273,8 @@ void memcpytest2_for_type(size_t numElements)
{
printSep();

DeviceMemory<T> memD(numElements);
HostMemory<T> memU(numElements, 0/*usePinnedHost*/);
DeviceMemory<T> memD(numElements);
HostMemory<T> memU(numElements, 0/*usePinnedHost*/);
HostMemory<T> memP(numElements, 1/*usePinnedHost*/);

for (int usePinnedHost =0; usePinnedHost<=1; usePinnedHost++) {
Expand Down Expand Up @@ -307,11 +307,11 @@ void memcpytest2_sizes(size_t maxElem=0)
maxElem = free/sizeof(T)/8;
}

printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n",
printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n",
deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0);
HIPCHECK ( hipDeviceReset() );
DeviceMemory<T> memD(maxElem);
HostMemory<T> memU(maxElem, 0/*usePinnedHost*/);
DeviceMemory<T> memD(maxElem);
HostMemory<T> memU(maxElem, 0/*usePinnedHost*/);
HostMemory<T> memP(maxElem, 1/*usePinnedHost*/);

for (size_t elem=1; elem<=maxElem; elem*=2) {
Expand All @@ -336,11 +336,11 @@ void memcpytest2_offsets(size_t maxElem, bool devOffsets, bool hostOffsets)
HIPCHECK(hipMemGetInfo(&free, &total));


printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n",
printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n",
deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0);
HIPCHECK ( hipDeviceReset() );
DeviceMemory<T> memD(maxElem);
HostMemory<T> memU(maxElem, 0/*usePinnedHost*/);
DeviceMemory<T> memD(maxElem);
HostMemory<T> memU(maxElem, 0/*usePinnedHost*/);
HostMemory<T> memP(maxElem, 1/*usePinnedHost*/);

size_t elem = maxElem / 2;
Expand Down Expand Up @@ -380,16 +380,16 @@ void multiThread_1(bool serialize, bool usePinnedHost)
{
printSep();
printf ("test: %s<%s> serialize=%d usePinnedHost=%d\n", __func__, TYPENAME(T), serialize, usePinnedHost);
DeviceMemory<T> memD(N);
HostMemory<T> mem1(N, usePinnedHost);
HostMemory<T> mem2(N, usePinnedHost);
DeviceMemory<T> memD(N);
HostMemory<T> mem1(N, usePinnedHost);
HostMemory<T> mem2(N, usePinnedHost);

std::thread t1 (memcpytest2<T>, &memD, &mem1, N, 0,0,0);
if (serialize) {
t1.join();
}


std::thread t2 (memcpytest2<T>,&memD, &mem2, N, 0,0,0);
if (serialize) {
t2.join();
Expand Down Expand Up @@ -427,21 +427,21 @@ int main(int argc, char *argv[])
// Some tests around the 64KB boundary which have historically shown issues:
printf ("\n\n=== tests&0x2 (64KB boundary)\n");
size_t maxElem = 32*1024*1024;
DeviceMemory<float> memD(maxElem);
HostMemory<float> memU(maxElem, 0/*usePinnedHost*/);
HostMemory<float> memP(maxElem, 0/*usePinnedHost*/);
DeviceMemory<float> memD(maxElem);
HostMemory<float> memU(maxElem, 0/*usePinnedHost*/);
HostMemory<float> memP(maxElem, 0/*usePinnedHost*/);
// These all pass:
memcpytest2<float>(&memD, &memP, 15*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 16*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 16*1024*1024+16*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 15*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 16*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 16*1024*1024+16*1024, 0, 0, 0);

// Just over 64MB:
memcpytest2<float>(&memD, &memP, 16*1024*1024+512*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 17*1024*1024+1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 32*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memU, 32*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 32*1024*1024, 1, 1, 0);
memcpytest2<float>(&memD, &memP, 32*1024*1024, 1, 1, 0);
memcpytest2<float>(&memD, &memP, 16*1024*1024+512*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 17*1024*1024+1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 32*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memU, 32*1024*1024, 0, 0, 0);
memcpytest2<float>(&memD, &memP, 32*1024*1024, 1, 1, 0);
memcpytest2<float>(&memD, &memP, 32*1024*1024, 1, 1, 0);


}
Expand All @@ -464,7 +464,7 @@ int main(int argc, char *argv[])

// 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);
multiThread_1<float>(true, true);

// Serialize, but use unpinned memory to stress the unpinned memory xfer path.
multiThread_1<float>(true, false);
Expand Down
4 changes: 2 additions & 2 deletions tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -41,8 +41,8 @@ void printSep()
// Designed to stress a small number of simple smoke tests

template<
typename T=float,
class P=HipTest::Unpinned,
typename T=float,
class P=HipTest::Unpinned,
class C=HipTest::Memcpy
>
void simpleVectorAdd(size_t numElements, int iters, hipStream_t stream)
Expand Down
10 changes: 5 additions & 5 deletions tests/src/runtimeApi/stream/hipNullStream.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -119,7 +119,7 @@ void Streamer<T>::reset()
{
HipTest::setDefaultData(_numElements, _A_h, _B_h, _C_h);
H2D();

}


Expand Down Expand Up @@ -238,7 +238,7 @@ int main(int argc, char *argv[])
nullStreamer->D2H();
HIPCHECK(hipDeviceSynchronize());

HipTest::checkTest(expected_H, nullStreamer->_C_h, numElements);
HipTest::checkTest(expected_H, nullStreamer->_C_h, numElements);
}
}

Expand All @@ -263,7 +263,7 @@ int main(int argc, char *argv[])

HIPCHECK(hipDeviceSynchronize());

HipTest::checkTest(expected_H, nullStreamer->_C_h, numElements);
HipTest::checkTest(expected_H, nullStreamer->_C_h, numElements);
}
}

Expand All @@ -289,10 +289,10 @@ int main(int argc, char *argv[])
// 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());

HipTest::checkTest(expected_H, streamers[0]->_C_h, numElements);
HipTest::checkTest(expected_H, streamers[0]->_C_h, numElements);
}


Expand Down
Loading

0 comments on commit d8e323d

Please sign in to comment.