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

This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own.

Browse Source 
[ROCm/hip commit: c2482d1255]
This commit is contained in:
Alex Voicu
2017-11-01 15:09:59 +00:00
parent e6e90e9cfa
commit 70a41e7dac
27 changed files with 1457 additions and 1352 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/context/hipMemsetD8.cpp
+1 -2
View File
@@ -46,7 +46,6 @@ int main(int argc, char *argv[])
A_h = new char[Nbytes];
HIPCHECK ( hipMalloc((void **) &A_d, Nbytes) );
A_h = (char*)malloc(Nbytes);
printf ("Size=%zu memsetval=%2x \n", Nbytes, memsetval);
HIPCHECK ( hipMemsetD8(A_d, memsetval, Nbytes) );
@@ -61,7 +60,7 @@ int main(int argc, char *argv[])
}
hipFree((void *) A_d);
free(A_h);
delete [] A_h;
passed();
}
projects/hip/tests/src/deviceLib/hipTestDevice.cpp
+124 -17
View File
@@ -139,7 +139,14 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -174,7 +181,14 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -205,7 +219,13 @@ for(int i=0;i<512;i++){
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -234,7 +254,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -263,7 +288,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -291,7 +321,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -321,7 +356,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -350,7 +390,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -387,7 +432,16 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
if(passed == 1){
return true;
}
@@ -427,7 +481,18 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
delete [] E;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
hipFree(Ed);
if(passed == 1){
return true;
}
@@ -457,7 +522,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -489,7 +559,14 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -525,7 +602,16 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
if(passed == 1){
return true;
}
@@ -565,7 +651,18 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
delete [] E;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
hipFree(Ed);
if(passed == 1){
return true;
}
@@ -595,7 +692,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -622,7 +724,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -631,7 +738,7 @@ return false;
}
int main(){
if(run_sincosf() && run_sincospif() && run_fdividef() &&
if(run_sincosf() && run_sincospif() && run_fdividef() &&
run_llrintf() && run_norm3df() && run_norm4df() &&
run_normf() && run_rnorm3df() && run_rnorm4df() &&
run_rnormf() && run_lroundf() && run_llroundf() &&
projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp
+110 -14
View File
@@ -128,7 +128,14 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -163,7 +170,14 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -193,7 +207,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -221,7 +240,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -249,7 +273,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -278,7 +307,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -306,7 +340,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -343,7 +382,16 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
if(passed == 1){
return true;
}
@@ -383,7 +431,18 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
delete [] E;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
hipFree(Ed);
if(passed == 1){
return true;
}
@@ -416,7 +475,14 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
if(passed == 1){
return true;
}
@@ -452,7 +518,16 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
if(passed == 1){
return true;
}
@@ -492,7 +567,18 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
delete [] C;
delete [] D;
delete [] E;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
hipFree(Dd);
hipFree(Ed);
if(passed == 1){
return true;
}
@@ -522,7 +608,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
@@ -549,7 +640,12 @@ for(int i=0;i<512;i++){
passed = 1;
}
}
free(A);
delete [] A;
delete [] B;
hipFree(Ad);
hipFree(Bd);
if(passed == 1){
return true;
}
projects/hip/tests/src/deviceLib/hip_test_ldg.cpp
+30 -15
View File
@@ -159,11 +159,16 @@ bool dataTypesRun(){
HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernel(vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
deviceA ,deviceB ,WIDTH ,HEIGHT);
hipLaunchKernel(
vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0,
0,
deviceA,
static_cast<const T*>(deviceB),
WIDTH,
HEIGHT);
HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost));
@@ -221,11 +226,16 @@ bool dataTypesRun2(){
HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T)));
HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernel(vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
deviceA ,deviceB,WIDTH ,HEIGHT);
hipLaunchKernel(
vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0,
0,
deviceA,
static_cast<const T*>(deviceB),
WIDTH,
HEIGHT);
HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost));
@@ -281,11 +291,16 @@ bool dataTypesRun4(){
HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernel(vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
deviceA ,deviceB ,WIDTH ,HEIGHT);
hipLaunchKernel(
vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0,
0,
deviceA,
static_cast<const T*>(deviceB),
WIDTH,
HEIGHT);
HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost));
projects/hip/tests/src/experimental/xcompile/hipxxKer.cpp
+8 -2
View File
@@ -36,17 +36,23 @@ __global__ void Kern(hipLaunchParm lp, float *A)
int main()
{
float *A, *Ad;
float A[len];
float *Ad;
for(int i=0;i<len;i++)
{
A[i] = 1.0f;
}
Ad = (float*)mallocHip(size);
memcpyHipH2D(Ad, A, size);
hipLaunchKernel(HIP_KERNEL_NAME(Kern), dim3(len/1024), dim3(1024), 0, 0, A);
hipLaunchKernel(
HIP_KERNEL_NAME(Kern), dim3(len/1024), dim3(1024), 0, 0, Ad);
memcpyHipD2H(A, Ad, size);
for(int i=0;i<len;i++)
{
assert(A[i] == 2.0f);
}
hipFree(Ad);
}
projects/hip/tests/src/kernel/hipLanguageExtensions.cpp
+2 -2
View File
@@ -74,8 +74,8 @@ __global__ void MyKernel (const hipLaunchParm lp, const float *a, const float *b
void callMyKernel()
{
float *a, *b, *c;
unsigned N;
const unsigned blockSize = 256;
unsigned N = blockSize;
hipLaunchKernel(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N);
}
@@ -102,7 +102,7 @@ vectorADD(const hipLaunchParm lp,
int a = __shfl_up(x, 1);
#endif
float x;
float x = 1.0;
float z = sin(x);
#ifdef NOT_YET
float fastZ = __sin(x);
projects/hip/tests/src/kernel/hipTestMemKernel.cpp
+30 -15
View File
@@ -107,9 +107,12 @@ int main(){
assert(C[i] == 1);
}
delete A;
delete B;
delete C;
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
A = new uint8_t[LEN9];
B = new uint8_t[LEN9];
@@ -132,9 +135,12 @@ int main(){
assert(C[i] == 1);
}
delete A;
delete B;
delete C;
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
A = new uint8_t[LEN10];
B = new uint8_t[LEN10];
@@ -157,9 +163,12 @@ int main(){
assert(C[i] == 1);
}
delete A;
delete B;
delete C;
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
A = new uint8_t[LEN11];
B = new uint8_t[LEN11];
@@ -182,9 +191,12 @@ int main(){
assert(C[i] == 1);
}
delete A;
delete B;
delete C;
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
A = new uint8_t[LEN12];
B = new uint8_t[LEN12];
@@ -207,9 +219,12 @@ int main(){
assert(C[i] == 1);
}
delete A;
delete B;
delete C;
delete [] A;
delete [] B;
delete [] C;
hipFree(Ad);
hipFree(Bd);
hipFree(Cd);
passed();
}
projects/hip/tests/src/runtimeApi/event/hipEventRecord.cpp
+10 -1
View File
@@ -69,7 +69,16 @@ int main(int argc, char *argv[])
// Record the start event
HIPCHECK (hipEventRecord(start, NULL));
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, N);
hipLaunchKernel(
HipTest::vectorADD,
dim3(blocks),
dim3(threadsPerBlock),
0,
0,
static_cast<const float*>(A_d),
static_cast<const float*>(B_d),
C_d,
N);
HIPCHECK (hipEventRecord(stop, NULL));
projects/hip/tests/src/runtimeApi/event/record_event.cpp
+20 -11
View File
@@ -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);
@@ -77,7 +77,16 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, hipStream_
HIPCHECK(hipEventRecord(timingDisabled, stream));
// sandwhich a kernel:
HIPCHECK(hipEventRecord(start, stream));
hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, stream, C_d, C_h, numElements, count);
hipLaunchKernelGGL(
HipTest::addCountReverse,
dim3(blocks),
dim3(threadsPerBlock),
0,
stream,
static_cast<const int*>(C_d),
C_h,
numElements,
count);
HIPCHECK(hipEventRecord(stop, stream));
@@ -85,8 +94,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) {
@@ -97,12 +106,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;
@@ -111,25 +120,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));
projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp
+43 -34
View File
@@ -58,7 +58,7 @@ public:
void offset(int offset) { _offset = offset; };
int offset() const { return _offset; };
private:
T * _A_d;
T* _B_d;
@@ -72,7 +72,7 @@ private:
template<typename T>
DeviceMemory<T>::DeviceMemory(size_t numElements)
: _maxNumElements(numElements),
: _maxNumElements(numElements),
_offset(0)
{
T ** np = nullptr;
@@ -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;
};
@@ -125,7 +125,7 @@ public:
T * A_hh;
T* B_hh;
bool _usePinnedHost;
bool _usePinnedHost;
private:
size_t _maxNumElements;
@@ -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) {
@@ -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,
@@ -243,7 +243,16 @@ void memcpytest2(DeviceMemory<T> *dmem, HostMemory<T> *hmem, size_t numElements,
HIPCHECK ( hipMemcpy(dmem->B_d(), hmem->B_h(), sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
}
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, dmem->A_d(), dmem->B_d(), dmem->C_d(), numElements);
hipLaunchKernel(
HipTest::vectorADD,
dim3(blocks),
dim3(threadsPerBlock),
0,
0,
static_cast<const T*>(dmem->A_d()),
static_cast<const T*>(dmem->B_d()),
dmem->C_d(),
numElements);
if (useDeviceToDevice) {
// Do an extra device-to-device copy here to mix things up:
@@ -273,8 +282,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++) {
@@ -307,11 +316,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) {
@@ -336,11 +345,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;
@@ -380,16 +389,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();
@@ -427,21 +436,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);
}
@@ -464,7 +473,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);
projects/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp
+10 -1
View File
@@ -63,7 +63,16 @@ void simpleTest1()
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);
hipLaunchKernel(
HipTest::vectorADD,
dim3(blocks),
dim3(threadsPerBlock),
0,
0,
static_cast<const int*>(A_d),
static_cast<const int*>(B_d),
C_d,
N);
HIPCHECK ( memcopy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp
+12 -3
View File
@@ -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)
@@ -90,7 +90,16 @@ void simpleVectorAdd(size_t numElements, int iters, hipStream_t stream)
// This is the null stream?
//hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, numElements);
hipLaunchKernel(HipTest::vectorADDReverse, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, numElements);
hipLaunchKernel(
HipTest::vectorADDReverse,
dim3(blocks),
dim3(threadsPerBlock),
0,
0,
static_cast<const T*>(A_d),
static_cast<const T*>(B_d),
C_d,
numElements);
MemTraits<C>::Copy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream);
projects/hip/tests/src/runtimeApi/stream/hipNullStream.cpp
+38 -8
View File
@@ -119,7 +119,7 @@ void Streamer<T>::reset()
{
HipTest::setDefaultData(_numElements, _A_h, _B_h, _C_h);
H2D();
}
@@ -128,7 +128,17 @@ 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, _A_d, _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);
}
@@ -225,7 +235,17 @@ 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*/, lastStreamer->_C_d, lastStreamer->_C_d, nullStreamer->_C_d, numElements, 1/*repeat*/);
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) {
@@ -238,7 +258,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);
}
}
@@ -257,13 +277,23 @@ 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*/, lastStreamer->_C_d, lastStreamer->_C_d, nullStreamer->_C_d, numElements, 1/*repeat*/);
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();
HIPCHECK(hipDeviceSynchronize());
HipTest::checkTest(expected_H, nullStreamer->_C_h, numElements);
HipTest::checkTest(expected_H, nullStreamer->_C_h, numElements);
}
}
@@ -289,10 +319,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);
}
projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
+24 -15
View File
@@ -59,23 +59,23 @@ const char *syncModeString(int syncMode) {
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
// This test sends a long-running kernel to the null stream, then tests to see if the
// specified synchronization technique is effective.
//
// Some syncMode are not expected to correctly sync (for example "syncNone"). in these
// Some syncMode are not expected to correctly sync (for example "syncNone"). in these
// cases the test sets expectMismatch and the check logic below will attempt to ensure that
// the undesired synchronization did not occur - ie ensure the kernel is still running and did
// not yet update the stop event. This can be tricky since if the kernel runs fast enough it
// may complete before the check. To prevent this, the addCountReverse has a count parameter
// which causes it to loop repeatedly, and the results are checked in reverse order.
// may complete before the check. To prevent this, the addCountReverse has a count parameter
// which causes it to loop repeatedly, and the results are checked in reverse order.
//
// Tests with expectMismatch=true should ensure the kernel finishes correctly. This results
// are checked and we test to make sure stop event has completed.
if (!(testMask & p_tests)) {
return;
}
printf ("\ntest 0x%02x: syncMode=%s expectMismatch=%d\n",
printf ("\ntest 0x%02x: syncMode=%s expectMismatch=%d\n",
testMask, syncModeString(syncMode), expectMismatch);
size_t sizeBytes = numElements * sizeof(int);
@@ -97,8 +97,17 @@ 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*/, 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:
@@ -108,18 +117,18 @@ void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode s
break;
case syncOtherStream:
// Does this synchronize with the null stream?
HIPCHECK(hipStreamSynchronize(otherStream));
HIPCHECK(hipStreamSynchronize(otherStream));
break;
case syncMarkerThenOtherStream:
case syncMarkerThenOtherNonBlockingStream:
// this may wait for NULL stream depending hipStreamNonBlocking flag above
HIPCHECK(hipEventRecord(otherStreamEvent, otherStream));
HIPCHECK(hipStreamSynchronize(otherStream));
// this may wait for NULL stream depending hipStreamNonBlocking flag above
HIPCHECK(hipEventRecord(otherStreamEvent, otherStream));
HIPCHECK(hipStreamSynchronize(otherStream));
break;
case syncDevice:
HIPCHECK(hipDeviceSynchronize());
HIPCHECK(hipDeviceSynchronize());
break;
default:
assert(0);
@@ -197,7 +206,7 @@ void runTests(int64_t numElements)
int main(int argc, char *argv[])
{
// Can' destroy the default stream:// TODO - move to another test
HIPCHECK_API(hipStreamDestroy(0), hipErrorInvalidResourceHandle);
HIPCHECK_API(hipStreamDestroy(0), hipErrorInvalidResourceHandle);
HipTest::parseStandardArguments(argc, argv, true /*failOnUndefinedArg*/);
projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp
+32 -14
View File
@@ -88,7 +88,7 @@ private:
template <typename T>
Streamer<T>::Streamer(int deviceId, T * A_d, size_t numElements, int commandType) :
_preA_d(NULL),
_preA_d(NULL),
_A_d(A_d),
_deviceId(deviceId),
_numElements(numElements),
@@ -163,9 +163,27 @@ void Streamer<T>::runAsyncAfter(Streamer<T> *depStreamer, bool waitSameStream)
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
if (_commandType == COMMAND_ADD_REVERSE) {
hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, _stream, _A_d, _C_d, _numElements, 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, _A_d, _C_d, _numElements, 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));
} else {
@@ -239,7 +257,7 @@ size_t Streamer<T>::check(int streamerNum, T initValue, T expectedOffset, bool e
return _mismatchCount;
}
//---
//Parse arguments specific to this test.
@@ -300,7 +318,7 @@ void checkAll(int initValue, std::vector<IntStreamer *> &streamers, std::vector<
for (int i=0; i<streamers.size(); i++) {
expected += streamers[i]->expectedAdd();
mismatchCount += streamers[i]->check(i+1, initValue, expected, expectPass);
}
@@ -330,7 +348,7 @@ void checkAll(int initValue, std::vector<IntStreamer *> &streamers, std::vector<
void sync_none(void) {};
void sync_allDevices(int numDevices)
void sync_allDevices(int numDevices)
{
for (int d=0; d<numDevices; d++) {
HIPCHECK(hipSetDevice(d));
@@ -339,7 +357,7 @@ void sync_allDevices(int numDevices)
}
void sync_queryAllUntilComplete(std::vector<IntStreamer *> streamers)
void sync_queryAllUntilComplete(std::vector<IntStreamer *> streamers)
{
for (int i=streamers.size()-1; i>=0; i--) {
streamers[i]->queryUntilComplete();
@@ -347,7 +365,7 @@ void sync_queryAllUntilComplete(std::vector<IntStreamer *> streamers)
}
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));
@@ -389,7 +407,7 @@ int main(int argc, char *argv[])
initArray_h[i] = initValue;
}
HIPCHECK(hipMemcpy(initArray_d, initArray_h, sizeElements, hipMemcpyHostToDevice));
int numDevices;
HIPCHECK(hipGetDeviceCount(&numDevices));
@@ -414,7 +432,7 @@ int main(int argc, char *argv[])
// 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
// 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++) {
@@ -446,7 +464,7 @@ int main(int argc, char *argv[])
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);
@@ -463,7 +481,7 @@ int main(int argc, char *argv[])
HIPCHECK(hipEventRecord(e1, s1))
HIPCHECK(hipStreamWaitEvent(hipStream_t(0), e1, 0/*flags*/));
HIPCHECK(hipStreamDestroy(s1));
HIPCHECK(hipEventDestroy(e1));
}
@@ -476,11 +494,11 @@ int main(int argc, char *argv[])
HIPCHECK(hipEventRecord(e1, hipStream_t(0)))
HIPCHECK(hipStreamWaitEvent(s1, e1, 0/*flags*/));
HIPCHECK(hipStreamDestroy(s1));
HIPCHECK(hipEventDestroy(e1));
}
}
projects/hip/tests/src/stress/hipStressKernel.cpp
+3
View File
@@ -57,5 +57,8 @@ int main(){
}
std::cout<<std::endl;
hipDeviceSynchronize();
free(A);
hipFree(Ad);
}
}