Merge branch 'amd-master-next' into amd-npi-next
Change-Id: Id241c60d6c2ceb4049c3ec15d9fe06baf28bcb3a
This commit is contained in:
Vlad Sytchenko
@@ -0,0 +1,340 @@
|
||||
/*
|
||||
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#include "hip/hip_fp16.h"
|
||||
|
||||
#define test_passed(test_name) \
|
||||
printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM);
|
||||
|
||||
enum half2Op {
|
||||
HALF2_OP_HEQ2 = 0,
|
||||
HALF2_OP_HNE2,
|
||||
HALF2_OP_HLE2,
|
||||
HALF2_OP_HGE2,
|
||||
HALF2_OP_HLT2,
|
||||
HALF2_OP_HGT2,
|
||||
HALF2_OP_MAX
|
||||
};
|
||||
|
||||
enum half2Test {
|
||||
HALF2_TEST_FUNCTION = 0,
|
||||
HALF2_TEST_NAN,
|
||||
HALF2_TEST_MAX
|
||||
};
|
||||
|
||||
// Kernels for half2 comparision functions
|
||||
|
||||
__global__
|
||||
void __half2Compare(float* result_D, __half2 a, int n, int half2Op,
|
||||
int testType) {
|
||||
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = gputhread; i < n; i += stride) {
|
||||
switch (half2Op) {
|
||||
case HALF2_OP_HEQ2:
|
||||
if (testType == HALF2_TEST_FUNCTION) {
|
||||
result_D[i] = __high2float(__heq2(__hadd2(a, __half2{1, 1}),
|
||||
__half2{2, 2}));
|
||||
} else {
|
||||
result_D[i] = __high2float(__heq2(__h2div(a, __half2{0, 0}),
|
||||
__half2{0, 0}));
|
||||
}
|
||||
break;
|
||||
case HALF2_OP_HNE2:
|
||||
result_D[i] = __high2float(__hne2(__hadd2(a, __half2{1, 1}),
|
||||
__half2{2, 2}));
|
||||
break;
|
||||
case HALF2_OP_HLE2:
|
||||
if (testType == HALF2_TEST_FUNCTION) {
|
||||
result_D[i] = __high2float(__hle2(__hadd2(a, __half2{1, 1}),
|
||||
__half2{3, 3}));
|
||||
} else {
|
||||
result_D[i] = __high2float(__hle2(__h2div(a, __half2{0, 0}),
|
||||
__half2{0, 0}));
|
||||
}
|
||||
break;
|
||||
case HALF2_OP_HGE2:
|
||||
if (testType == HALF2_TEST_FUNCTION) {
|
||||
result_D[i] = __high2float(__hge2(__hadd2(a, __half2{1, 1}),
|
||||
__half2{2, 2}));
|
||||
} else {
|
||||
result_D[i] = __high2float(__hge2(__h2div(a, __half2{0, 0}),
|
||||
__half2{0, 0}));
|
||||
}
|
||||
break;
|
||||
case HALF2_OP_HLT2:
|
||||
if (testType == HALF2_TEST_FUNCTION) {
|
||||
result_D[i] = __high2float(__hlt2(__hadd2(a, __half2{1, 1}),
|
||||
__half2{3, 3}));
|
||||
} else {
|
||||
result_D[i] = __high2float(__hlt2(__h2div(a, __half2{0, 0}),
|
||||
__half2{0, 0}));
|
||||
}
|
||||
break;
|
||||
case HALF2_OP_HGT2:
|
||||
if (testType == HALF2_TEST_FUNCTION) {
|
||||
result_D[i] = __high2float(__hgt2(__hadd2(a, __half2{1, 1}),
|
||||
__half2{3, 3}));
|
||||
} else {
|
||||
result_D[i] = __high2float(__hgt2(__h2div(a, __half2{0, 0}),
|
||||
__half2{0, 0}));
|
||||
}
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static bool isFailed(float expectedValue, float *result_H, int size) {
|
||||
for (int index = 0; index < size; index++) {
|
||||
if (expectedValue != result_H[index]) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
int main() {
|
||||
const int n = 64;
|
||||
float* result_H = reinterpret_cast<float*>(malloc(n*sizeof(float)));
|
||||
float* result_D;
|
||||
bool bFunctionalTestFailed = false;
|
||||
bool bNanTestFailed = false;
|
||||
int index = 0;
|
||||
HIPCHECK(hipMalloc(&result_D, n*sizeof(float)));
|
||||
|
||||
// kernel launch and hipmemcpy operation to get return value for heq2
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{1, 1}, n, HALF2_OP_HEQ2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("heq2: failure when arguments are equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{2, 2}, n, HALF2_OP_HEQ2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("heq2: failure when arguments are not equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
// kernel launch and hipmemcpy operation to get return value for hne2
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{2, 2}, n, HALF2_OP_HNE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hne2: failure when arguments are not equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{1, 1}, n, HALF2_OP_HNE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hne2: failure when arguments are equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
// kernel launch and hipmemcpy operation to get return value for hle2
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{1, 1}, n, HALF2_OP_HLE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hle2: failure when argument is less than equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{2, 2}, n, HALF2_OP_HLE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hle2: failure when argument is equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{3, 3}, n, HALF2_OP_HLE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hle2: failure when argument is greater\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
// kernel launch and hipmemcpy operation to get return value for hge2
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{2, 2}, n, HALF2_OP_HGE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hge2: failure when argument is greater\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{1, 1}, n, HALF2_OP_HGE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hge2: failure when argument is equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{0, 0}, n, HALF2_OP_HGE2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hge2: failure when argument is less\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
// kernel launch and hipmemcpy operation to get return value for hlt2
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{1, 1}, n, HALF2_OP_HLT2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hlt2: failure when argument is less\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{2, 2}, n, HALF2_OP_HLT2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hlt2: failure when argument is equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{3, 3}, n, HALF2_OP_HLT2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hlt2: failure when argument is greater\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
// kernel launch and hipmemcpy operation to get return value for hgt2
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{3, 3}, n, HALF2_OP_HGT2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(1.0, result_H, n)) {
|
||||
printf("hgt2: failure when argument is greater\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{2, 2}, n, HALF2_OP_HGT2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hgt2: failure when argument is equal\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{1, 1}, n, HALF2_OP_HGT2,
|
||||
HALF2_TEST_FUNCTION);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("hgt2: failure when argument is less\n");
|
||||
bFunctionalTestFailed = true;
|
||||
}
|
||||
|
||||
for (int nanFunctionTest = HALF2_OP_HEQ2; nanFunctionTest < HALF2_OP_MAX;
|
||||
nanFunctionTest++) {
|
||||
// HNE2 will not have a NaN test
|
||||
if (nanFunctionTest != HALF2_OP_HNE2) {
|
||||
hipLaunchKernelGGL(__half2Compare, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
|
||||
result_D, __half2{0, 0}, n, nanFunctionTest,
|
||||
HALF2_TEST_NAN);
|
||||
hipDeviceSynchronize();
|
||||
HIPCHECK(hipMemcpy(result_H, result_D, n*sizeof(float),
|
||||
hipMemcpyDeviceToHost));
|
||||
if (isFailed(0.0, result_H, n)) {
|
||||
printf("NaN test failed for half function: %d\n", nanFunctionTest);
|
||||
bNanTestFailed = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
hipFree(result_D);
|
||||
free(result_H);
|
||||
|
||||
if ((false == bFunctionalTestFailed) && (false == bNanTestFailed)) {
|
||||
passed();
|
||||
} else {
|
||||
failed("Some Half2 tests failed");
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
@@ -186,6 +186,14 @@ void test_fp16() {
|
||||
CHECK_SIMPLE([]__device__(){ return min<__fp16>(1.0f, 2.0f); }, 1.0f);
|
||||
}
|
||||
|
||||
void test_pown() {
|
||||
CHECK_SIMPLE([]__device__(){ return powif(2.0f, 2); }, 4.0f);
|
||||
CHECK_SIMPLE([]__device__(){ return powi(2.0, 2); }, 4.0);
|
||||
CHECK_SIMPLE([]__device__(){ return pow(2.0f, 2); }, 4.0f);
|
||||
CHECK_SIMPLE([]__device__(){ return pow(2.0, 2); }, 4.0);
|
||||
CHECK_SIMPLE([]__device__(){ return pow(2.0f16, 2); }, 4.0f16);
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
@@ -195,5 +203,7 @@ int main(int argc, char* argv[]) {
|
||||
|
||||
test_fp16();
|
||||
|
||||
test_pown();
|
||||
|
||||
passed();
|
||||
}
|
||||
|
||||
@@ -0,0 +1,155 @@
|
||||
/*
|
||||
Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
|
||||
#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
|
||||
|
||||
static __global__
|
||||
void kernel_syncthreads_and(int *syncTestD,
|
||||
int *allThreadsZeroD,
|
||||
int *allThreadsOneD,
|
||||
int *oneThreadZeroD,
|
||||
int *allThreadsMinusOneD)
|
||||
{
|
||||
int blockSize = blockDim.x;
|
||||
int predicate = 10;
|
||||
// First block index starts with 0, and second block index starts
|
||||
// with blockSize
|
||||
int i = (blockIdx.x == 0) ? threadIdx.x : blockSize + threadIdx.x;
|
||||
|
||||
// At very first, we need to ensure work-group level syncronization
|
||||
// properly happened, don't bother about predicate testing for now.
|
||||
// Thread 0 and thread 1 writes to shared memory. After call to api,
|
||||
// every thread reads shared memory, and store product for verification
|
||||
__shared__ int sm[2];
|
||||
if (threadIdx.x == 0)
|
||||
sm[0] = 10;
|
||||
else if (threadIdx.x == 1)
|
||||
sm[1] = 20;
|
||||
__syncthreads_and(predicate);
|
||||
syncTestD[i] = sm[0] * sm[1];
|
||||
|
||||
// All threads pass 0 as predicate value, result should be 0
|
||||
predicate = 0;
|
||||
allThreadsZeroD[i] = __syncthreads_and(predicate);
|
||||
|
||||
// All threads pass 1 as predicate value, result should be 1
|
||||
predicate = 1;
|
||||
allThreadsOneD[i] = __syncthreads_and(predicate);
|
||||
|
||||
// Thread 0 pass 0, and all other threads 1 as predicate value,
|
||||
// result should be 0
|
||||
predicate = (threadIdx.x == 0) ? 0 : 1;
|
||||
oneThreadZeroD[i] = __syncthreads_and(predicate);
|
||||
|
||||
// All threads pass -1 as predicate value, result should be 1
|
||||
predicate = -1;
|
||||
allThreadsMinusOneD[i] = __syncthreads_and(predicate);
|
||||
}
|
||||
|
||||
static void test_syncthreads_and(int blockSize)
|
||||
{
|
||||
int nBytes = sizeof(int) * 2 * blockSize;
|
||||
int * syncTestD, *syncTestH;
|
||||
int *allThreadsZeroD, *allThreadsZeroH;
|
||||
int *allThreadsOneD, *allThreadsOneH;
|
||||
int *oneThreadZeroD, *oneThreadZeroH;
|
||||
int *allThreadsMinusOneD, *allThreadsMinusOneH;
|
||||
|
||||
// Allocate device memory
|
||||
ASSERT_EQUAL(hipMalloc((void**)&syncTestD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsZeroD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsOneD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&oneThreadZeroD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsMinusOneD, nBytes), hipSuccess);
|
||||
|
||||
// Allocate host memory
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&syncTestH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsZeroH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsOneH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&oneThreadZeroH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsMinusOneH, nBytes), hipSuccess);
|
||||
|
||||
// Launch Kernel
|
||||
hipLaunchKernelGGL(kernel_syncthreads_and,
|
||||
2,
|
||||
blockSize,
|
||||
0,
|
||||
0,
|
||||
syncTestD,
|
||||
allThreadsZeroD,
|
||||
allThreadsOneD,
|
||||
oneThreadZeroD,
|
||||
allThreadsMinusOneD);
|
||||
|
||||
// Copy result from device to host
|
||||
ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsZeroH, allThreadsZeroD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsOneH, allThreadsOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(oneThreadZeroH, oneThreadZeroD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsMinusOneH, allThreadsMinusOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
|
||||
// Validate results for both blocks together
|
||||
for (int i = 0; i < 2 * blockSize; ++i) {
|
||||
ASSERT_EQUAL(syncTestH[i], 200);
|
||||
ASSERT_EQUAL(allThreadsZeroH[i], 0);
|
||||
ASSERT_EQUAL(allThreadsOneH[i], 1);
|
||||
ASSERT_EQUAL(oneThreadZeroH[i], 0);
|
||||
ASSERT_EQUAL(allThreadsMinusOneH[i], 1);
|
||||
}
|
||||
|
||||
// Free device memory
|
||||
ASSERT_EQUAL(hipFree(syncTestD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsZeroD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsOneD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(oneThreadZeroD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsMinusOneD), hipSuccess);
|
||||
|
||||
//Free host memory
|
||||
ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsZeroH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsOneH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(oneThreadZeroH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsMinusOneH), hipSuccess);
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
int blockSizes[] = {10, 40, 70, 130, 240, 723, 32, 64, 128, 256, 512, 1024};
|
||||
for (int i = 0; i < (sizeof(blockSizes) / sizeof(blockSizes[0])); ++i)
|
||||
test_syncthreads_and(blockSizes[i]);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,169 @@
|
||||
/*
|
||||
Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
|
||||
#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
|
||||
|
||||
static __global__
|
||||
void kernel_syncthreads_count(int *syncTestD,
|
||||
int *allThreadsZeroD,
|
||||
int *allThreadsOneD,
|
||||
int *oddThreadsOneD,
|
||||
int *allThreadsMinusOneD,
|
||||
int *allThreadsIdD)
|
||||
{
|
||||
int blockSize = blockDim.x;
|
||||
int predicate = 10;
|
||||
// First block index starts with 0, and second block index starts
|
||||
// with blockSize
|
||||
int i = (blockIdx.x == 0) ? threadIdx.x : blockSize + threadIdx.x;
|
||||
|
||||
// At very first, we need to ensure work-group level syncronization
|
||||
// properly happened, don't bother about predicate testing for now.
|
||||
// Thread 0 and thread 1 writes to shared memory. After call to api,
|
||||
// every thread reads shared memory, and store sum for verification
|
||||
__shared__ int sm[2];
|
||||
if (threadIdx.x == 0)
|
||||
sm[0] = 10;
|
||||
else if (threadIdx.x == 1)
|
||||
sm[1] = 20;
|
||||
__syncthreads_count(predicate);
|
||||
syncTestD[i] = sm[0] + sm[1];
|
||||
|
||||
// All threads pass 0 as predicate value, result should be 0
|
||||
predicate = 0;
|
||||
allThreadsZeroD[i] = __syncthreads_count(predicate);
|
||||
|
||||
// All threads pass 1 as predicate value, result should be blockSize
|
||||
predicate = 1;
|
||||
allThreadsOneD[i] = __syncthreads_count(predicate);
|
||||
|
||||
// Odd numbered threads pass 1, and even numbered threads pass 0, as
|
||||
// predicate value, result should be blockSize / 2
|
||||
predicate = threadIdx.x % 2;
|
||||
oddThreadsOneD[i] = __syncthreads_count(predicate);
|
||||
|
||||
// All threads pass -1 as predicate value, result should blockSize
|
||||
predicate = -1;
|
||||
allThreadsMinusOneD[i] = __syncthreads_count(predicate);
|
||||
|
||||
// Each thread pass its ID as predicate value, result should be blockSize - 1
|
||||
predicate = threadIdx.x;
|
||||
allThreadsIdD[i] = __syncthreads_count(predicate);
|
||||
}
|
||||
|
||||
void test_syncthreads_count(int blockSize)
|
||||
{
|
||||
int nBytes = sizeof(int) * 2 * blockSize;
|
||||
int * syncTestD, *syncTestH;
|
||||
int *allThreadsZeroD, *allThreadsZeroH;
|
||||
int *allThreadsOneD, *allThreadsOneH;
|
||||
int *oddThreadsOneD, *oddThreadsOneH;
|
||||
int *allThreadsMinusOneD, *allThreadsMinusOneH;
|
||||
int *allThreadsIdD, *allThreadsIdH;
|
||||
|
||||
// Allocate device memory
|
||||
ASSERT_EQUAL(hipMalloc((void**)&syncTestD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsZeroD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsOneD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&oddThreadsOneD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsMinusOneD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsIdD, nBytes), hipSuccess);
|
||||
|
||||
// Allocate host memory
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&syncTestH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsZeroH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsOneH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&oddThreadsOneH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsMinusOneH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsIdH, nBytes), hipSuccess);
|
||||
|
||||
// Launch Kernel
|
||||
hipLaunchKernelGGL(kernel_syncthreads_count,
|
||||
2,
|
||||
blockSize,
|
||||
0,
|
||||
0,
|
||||
syncTestD,
|
||||
allThreadsZeroD,
|
||||
allThreadsOneD,
|
||||
oddThreadsOneD,
|
||||
allThreadsMinusOneD,
|
||||
allThreadsIdD);
|
||||
|
||||
// Copy result from device to host
|
||||
ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsZeroH, allThreadsZeroD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsOneH, allThreadsOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(oddThreadsOneH, oddThreadsOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsMinusOneH, allThreadsMinusOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsIdH, allThreadsIdD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
|
||||
// Validate results for both the blocks together
|
||||
for (int i = 0; i < 2 * blockSize; ++i) {
|
||||
ASSERT_EQUAL(syncTestH[i], 30);
|
||||
ASSERT_EQUAL(allThreadsZeroH[i], 0);
|
||||
ASSERT_EQUAL(allThreadsOneH[i], blockSize);
|
||||
ASSERT_EQUAL(oddThreadsOneH[i], blockSize / 2);
|
||||
ASSERT_EQUAL(allThreadsMinusOneH[i], blockSize);
|
||||
ASSERT_EQUAL(allThreadsIdH[i], (blockSize-1));
|
||||
}
|
||||
|
||||
// Free device memory
|
||||
ASSERT_EQUAL(hipFree(syncTestD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsZeroD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsOneD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(oddThreadsOneD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsMinusOneD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsIdD), hipSuccess);
|
||||
|
||||
//Free host memory
|
||||
ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsZeroH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsOneH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(oddThreadsOneH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsMinusOneH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsIdH), hipSuccess);
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
int blockSizes[] = {10, 40, 70, 130, 240, 723, 32, 64, 128, 256, 512, 1024};
|
||||
for (int i = 0; i < (sizeof(blockSizes) / sizeof(blockSizes[0])); ++i)
|
||||
test_syncthreads_count(blockSizes[i]);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,155 @@
|
||||
/*
|
||||
Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
|
||||
#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
|
||||
|
||||
static __global__
|
||||
void kernel_syncthreads_or(int *syncTestD,
|
||||
int *allThreadsZeroD,
|
||||
int *allThreadsOneD,
|
||||
int *oneThreadOneD,
|
||||
int *allThreadsMinusOneD)
|
||||
{
|
||||
int blockSize = blockDim.x;
|
||||
int predicate = 10;
|
||||
// First block index starts with 0, and second block index starts
|
||||
// with blockSize
|
||||
int i = (blockIdx.x == 0) ? threadIdx.x : blockSize + threadIdx.x;
|
||||
|
||||
// At very first, we need to ensure work-group level syncronization
|
||||
// properly happened, don't bother about predicate testing for now.
|
||||
// Thread 0 and thread 1 writes to shared memory. After call to api,
|
||||
// every thread reads shared memory, and store subtraction for verification
|
||||
__shared__ int sm[2];
|
||||
if (threadIdx.x == 0)
|
||||
sm[0] = 10;
|
||||
else if (threadIdx.x == 1)
|
||||
sm[1] = 20;
|
||||
__syncthreads_or(predicate);
|
||||
syncTestD[i] = sm[1] - sm[0];
|
||||
|
||||
// All threads pass 0 as predicate value, result should be 0
|
||||
predicate = 0;
|
||||
allThreadsZeroD[i] = __syncthreads_or(predicate);
|
||||
|
||||
// All threads pass 1 as predicate value, result should be 1
|
||||
predicate = 1;
|
||||
allThreadsOneD[i] = __syncthreads_or(predicate);
|
||||
|
||||
// Thread 0 pass 1, and all other threads 0 as predicate value,
|
||||
// result should be 1
|
||||
predicate = (threadIdx.x == 0) ? 1 : 0;
|
||||
oneThreadOneD[i] = __syncthreads_or(predicate);
|
||||
|
||||
// All threads pass -1 as predicate value, result should be 1
|
||||
predicate = -1;
|
||||
allThreadsMinusOneD[i] = __syncthreads_or(predicate);
|
||||
}
|
||||
|
||||
static void test_syncthreads_or(int blockSize)
|
||||
{
|
||||
int nBytes = sizeof(int) * 2 * blockSize;
|
||||
int * syncTestD, *syncTestH;
|
||||
int *allThreadsZeroD, *allThreadsZeroH;
|
||||
int *allThreadsOneD, *allThreadsOneH;
|
||||
int *oneThreadOneD, *oneThreadOneH;
|
||||
int *allThreadsMinusOneD, *allThreadsMinusOneH;
|
||||
|
||||
// Allocate device memory
|
||||
ASSERT_EQUAL(hipMalloc((void**)&syncTestD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsZeroD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsOneD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&oneThreadOneD, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipMalloc((void**)&allThreadsMinusOneD, nBytes), hipSuccess);
|
||||
|
||||
// Allocate host memory
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&syncTestH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsZeroH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsOneH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&oneThreadOneH, nBytes), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostMalloc((void**)&allThreadsMinusOneH, nBytes), hipSuccess);
|
||||
|
||||
// Launch Kernel
|
||||
hipLaunchKernelGGL(kernel_syncthreads_or,
|
||||
2,
|
||||
blockSize,
|
||||
0,
|
||||
0,
|
||||
syncTestD,
|
||||
allThreadsZeroD,
|
||||
allThreadsOneD,
|
||||
oneThreadOneD,
|
||||
allThreadsMinusOneD);
|
||||
|
||||
// Copy result from device to host
|
||||
ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsZeroH, allThreadsZeroD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsOneH, allThreadsOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(oneThreadOneH, oneThreadOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
ASSERT_EQUAL(hipMemcpy(allThreadsMinusOneH, allThreadsMinusOneD, nBytes, hipMemcpyDeviceToHost),
|
||||
hipSuccess);
|
||||
|
||||
// Validate results for both blocks together
|
||||
for (int i = 0; i < 2 * blockSize; ++i) {
|
||||
ASSERT_EQUAL(syncTestH[i], 10);
|
||||
ASSERT_EQUAL(allThreadsZeroH[i], 0);
|
||||
ASSERT_EQUAL(allThreadsOneH[i], 1);
|
||||
ASSERT_EQUAL(oneThreadOneH[i], 1);
|
||||
ASSERT_EQUAL(allThreadsMinusOneH[i], 1);
|
||||
}
|
||||
|
||||
// Free device memory
|
||||
ASSERT_EQUAL(hipFree(syncTestD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsZeroD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsOneD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(oneThreadOneD), hipSuccess);
|
||||
ASSERT_EQUAL(hipFree(allThreadsMinusOneD), hipSuccess);
|
||||
|
||||
//Free host memory
|
||||
ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsZeroH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsOneH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(oneThreadOneH), hipSuccess);
|
||||
ASSERT_EQUAL(hipHostFree(allThreadsMinusOneH), hipSuccess);
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
int blockSizes[] = {10, 40, 70, 130, 240, 723, 32, 64, 128, 256, 512, 1024};
|
||||
for (int i = 0; i < (sizeof(blockSizes) / sizeof(blockSizes[0])); ++i)
|
||||
test_syncthreads_or(blockSizes[i]);
|
||||
passed();
|
||||
}
|
||||
Executable
+90
@@ -0,0 +1,90 @@
|
||||
/*
|
||||
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
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
#define N 1024
|
||||
#define OFFSET 128
|
||||
|
||||
void single_process() {
|
||||
int* ipc_dptr = nullptr;
|
||||
int* ipc_hptr = nullptr;
|
||||
int* ipc_out_dptr = nullptr;
|
||||
int* ipc_out_hptr = nullptr;
|
||||
|
||||
int* ipc_offset_dptr = nullptr;
|
||||
|
||||
hipIpcMemHandle_t ipc_handle;
|
||||
hipIpcMemHandle_t ipc_offset_handle;
|
||||
|
||||
HIPCHECK(hipMalloc((void**)&ipc_dptr, N * sizeof(int)));
|
||||
|
||||
// Negative, Make sure we return error when an offset of original ptr is passed
|
||||
ipc_offset_dptr = ipc_dptr + (OFFSET * sizeof(int));
|
||||
assert(hipErrorInvalidDevicePointer == hipIpcGetMemHandle(&ipc_offset_handle, ipc_offset_dptr));
|
||||
|
||||
// Get handle for the device_ptr
|
||||
HIPCHECK(hipIpcGetMemHandle(&ipc_handle, ipc_dptr));
|
||||
|
||||
// Set Values @ Host Ptr
|
||||
ipc_hptr = new int[N];
|
||||
for (size_t idx = 0; idx < N; ++idx) {
|
||||
ipc_hptr[idx] = idx;
|
||||
}
|
||||
|
||||
// Copy values to Device ptr
|
||||
HIPCHECK(hipMemset(ipc_dptr, 0x00, (N * sizeof(int))));
|
||||
HIPCHECK(hipMemcpy(ipc_dptr, ipc_hptr, (N * sizeof(int)), hipMemcpyHostToDevice));
|
||||
|
||||
// Open handle to get dev_ptr
|
||||
ipc_out_hptr = new int[N];
|
||||
memset(ipc_out_hptr, 0x00, (N * sizeof(int)));
|
||||
HIPCHECK(hipIpcOpenMemHandle((void**)&ipc_out_dptr, ipc_handle, 0));
|
||||
|
||||
// Copy Values from Device to Host and Check for correctness
|
||||
HIPCHECK(hipMemcpy(ipc_out_hptr, ipc_out_dptr, (N * sizeof(int)), hipMemcpyDeviceToHost));
|
||||
for (size_t idx = 0; idx < N; ++idx) {
|
||||
if(ipc_out_hptr[idx] != idx) {
|
||||
std::cout<<"Failing @ idx: "<<idx<<std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
//Close All Mem Handle
|
||||
HIPCHECK(hipIpcCloseMemHandle(ipc_out_dptr));
|
||||
HIPCHECK(hipFree(ipc_dptr));
|
||||
|
||||
delete[] ipc_hptr;
|
||||
delete[] ipc_out_hptr;
|
||||
}
|
||||
|
||||
void multi_process() {
|
||||
//To create and open IPC handle via multiple process
|
||||
}
|
||||
|
||||
int main() {
|
||||
single_process();
|
||||
multi_process();
|
||||
passed();
|
||||
}
|
||||
@@ -571,6 +571,10 @@ __global__ void hipLaunchKernelStructFunc21(
|
||||
|
||||
__global__ void vAdd(float* a) {}
|
||||
|
||||
template<class T1, class T2>
|
||||
__global__ void myKernel(T1 a, T2 b) {}
|
||||
|
||||
|
||||
//---
|
||||
// Some wrapper macro for testing:
|
||||
#define WRAP(...) __VA_ARGS__
|
||||
@@ -913,6 +917,18 @@ int main() {
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad);
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad);
|
||||
|
||||
// Test: Passing macro to hipLaunchKernelGGL
|
||||
#define KERNEL_CONFIG dim3(1024), dim3(1), 0, 0
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), KERNEL_CONFIG, Ad);
|
||||
|
||||
// Test: Same thing with templates:
|
||||
int a;
|
||||
float b;
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(myKernel<int, float>), KERNEL_CONFIG, a, b);
|
||||
|
||||
#define TYPE_PARAM_CONFIG int, float
|
||||
hipLaunchKernelGGL(HIP_KERNEL_NAME(myKernel<TYPE_PARAM_CONFIG>), KERNEL_CONFIG, a, b);
|
||||
|
||||
// Test: Passing hipLaunchKernelGGL inside another macro:
|
||||
float e0;
|
||||
MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024),
|
||||
|
||||
@@ -1,51 +1,168 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
/*
|
||||
* Test to compare
|
||||
* 1.pciBusID from hipDeviceGetPCIBusId and hipDeviceGetAttribute **
|
||||
* 2.{pciDomainID, pciBusID, pciDeviceID} values hipDeviceGetPCIBusId vs lspci **
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* TEST: %t
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST_NAMED: %t hipDeviceGetPCIBusId-vs-hipDeviceGetAttribute --tests 0x1
|
||||
* TEST_NAMED: %t hipDeviceGetPCIBusId-vs-lspci --tests 0x2
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <stdio.h>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
#define MAX_DEVICE_LENGTH 20
|
||||
|
||||
int main(void) {
|
||||
char pciBusId[13];
|
||||
int deviceCount = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&deviceCount));
|
||||
HIPASSERT(deviceCount != 0);
|
||||
for (int i = 0; i < deviceCount; i++) {
|
||||
int pciBusID = -1;
|
||||
int pciDeviceID = -1;
|
||||
int pciDomainID = -1;
|
||||
int tempPciBusId = -1;
|
||||
HIPCHECK(hipDeviceGetPCIBusId(&pciBusId[0], 13, i));
|
||||
sscanf(pciBusId, "%04x:%02x:%02x", &pciDomainID, &pciBusID, &pciDeviceID);
|
||||
HIPCHECK(hipDeviceGetAttribute(&tempPciBusId, hipDeviceAttributePciBusId, i));
|
||||
if (pciBusID != tempPciBusId) {
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
passed();
|
||||
static bool getPciBusId(int deviceCount, char hipDeviceList[][MAX_DEVICE_LENGTH]) {
|
||||
for (int i = 0; i < deviceCount; i++) {
|
||||
HIPCHECK(hipDeviceGetPCIBusId(hipDeviceList[i], MAX_DEVICE_LENGTH, i));
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
bool comparePciBusIDWithHipDeviceGetAttribute() {
|
||||
bool testResult = true;
|
||||
int deviceCount = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&deviceCount));
|
||||
HIPASSERT(deviceCount != 0);
|
||||
printf("No.of gpus in the system: %d\n", deviceCount);
|
||||
char hipDeviceList[deviceCount][MAX_DEVICE_LENGTH];
|
||||
char pciDeviceList[deviceCount][MAX_DEVICE_LENGTH];
|
||||
|
||||
getPciBusId(deviceCount, hipDeviceList);
|
||||
|
||||
for (int i = 0; i < deviceCount; i++) {
|
||||
int pciBusID = -1;
|
||||
int pciDeviceID = -1;
|
||||
int pciDomainID = -1;
|
||||
int tempPciBusId = -1;
|
||||
sscanf(hipDeviceList[i], "%04x:%02x:%02x", &pciDomainID, &pciBusID,
|
||||
&pciDeviceID);
|
||||
HIPCHECK(hipDeviceGetAttribute(&tempPciBusId, hipDeviceAttributePciBusId, i));
|
||||
if (pciBusID != tempPciBusId) {
|
||||
testResult = false;
|
||||
printf("pciBusID from hipDeviceGetPCIBusId mismatched to that from "
|
||||
"hipDeviceGetAttribute for gpu %d\n", i);
|
||||
}
|
||||
}
|
||||
|
||||
printf("pciBusID output of both hipDeviceGetPCIBusId and"
|
||||
" hipDeviceGetAttribute matched for all gpus\n");
|
||||
return testResult;
|
||||
}
|
||||
|
||||
bool compareHipDeviceGetPCIBusIdWithLspci() {
|
||||
FILE *fpipe;
|
||||
bool testResult = false;
|
||||
|
||||
{
|
||||
// Check if lspci is installed, if not, don't proceed
|
||||
char const *cmd = "lspci --version";
|
||||
char *lspciCheck;
|
||||
char temp[20];
|
||||
fpipe = popen(cmd, "r");
|
||||
|
||||
if (fpipe == nullptr) {
|
||||
printf("Unable to create command file\n");
|
||||
return testResult;
|
||||
}
|
||||
|
||||
lspciCheck = fgets(temp, 20, fpipe);
|
||||
pclose(fpipe);
|
||||
|
||||
if (!lspciCheck) {
|
||||
printf("lspci not found. Skipping the test\n");
|
||||
return true;
|
||||
}
|
||||
}
|
||||
|
||||
int deviceCount = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&deviceCount));
|
||||
HIPASSERT(deviceCount != 0);
|
||||
printf("No.of gpus in the system: %d\n", deviceCount);
|
||||
char hipDeviceList[deviceCount][MAX_DEVICE_LENGTH];
|
||||
char pciDeviceList[deviceCount][MAX_DEVICE_LENGTH];
|
||||
|
||||
getPciBusId(deviceCount, hipDeviceList);
|
||||
|
||||
// Get lspci device list and compare with hip device list
|
||||
char const *command = "lspci -D | grep controller | grep AMD/ATI | "
|
||||
"cut -d ' ' -f 1";
|
||||
fpipe = popen(command, "r");
|
||||
|
||||
if (fpipe == nullptr) {
|
||||
printf("Unable to create command file\n");
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int index = 0;
|
||||
int deviceMatchCount = 0;
|
||||
|
||||
while (fgets(pciDeviceList[index], sizeof(pciDeviceList[index]), fpipe)) {
|
||||
bool bMatchFound = false;
|
||||
for (int deviceNo = 0; deviceNo < deviceCount; deviceNo++) {
|
||||
if (!strncmp(pciDeviceList[index], hipDeviceList[deviceNo], 10)) {
|
||||
deviceMatchCount++;
|
||||
bMatchFound = true;
|
||||
}
|
||||
}
|
||||
if (bMatchFound == false) {
|
||||
printf("PCI device: %s is not reported by HIP\n", pciDeviceList[index]);
|
||||
}
|
||||
index++;
|
||||
}
|
||||
|
||||
pclose(fpipe);
|
||||
|
||||
if (deviceMatchCount == deviceCount) {
|
||||
printf("hip and lspci output for {pciDomainID, pciBusID, pciDeviceID} "
|
||||
"matched for all gpus\n");
|
||||
testResult = true;
|
||||
} else {
|
||||
printf("Mismatch in number GPUs reported by HIP with lscpi\n");
|
||||
}
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
bool testResult = true;
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
if (p_tests & 0x1) {
|
||||
testResult &= comparePciBusIDWithHipDeviceGetAttribute();
|
||||
}
|
||||
|
||||
if (p_tests & 0x2) {
|
||||
#ifdef __unix__
|
||||
testResult &= compareHipDeviceGetPCIBusIdWithLspci();
|
||||
#else
|
||||
printf("Detected non-linux OS. Skipping the test\n");
|
||||
#endif
|
||||
}
|
||||
|
||||
if (testResult) {
|
||||
passed();
|
||||
} else {
|
||||
failed("one or more tests failed\n");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -145,5 +145,6 @@ int main(int argc, char* argv[]) {
|
||||
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeKernelExecTimeout, props.kernelExecTimeoutEnabled));
|
||||
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeCanMapHostMemory, props.canMapHostMemory));
|
||||
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeEccEnabled, props.ECCEnabled));
|
||||
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeAsicRevision, props.asicRevision));
|
||||
passed();
|
||||
};
|
||||
|
||||
Regular → Executable
+8
-4
@@ -39,6 +39,9 @@ int main() {
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
if (numDevices > 1) {
|
||||
int canAccessPeer = 0;
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
|
||||
if (canAccessPeer) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
@@ -47,23 +50,21 @@ int main() {
|
||||
HIPCHECK(hipMalloc(&Y_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&Z_d, Nbytes));
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(1));
|
||||
HIPCHECK(hipMemcpyDtoD((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, Nbytes));
|
||||
HIPCHECK(hipMemcpyDtoD((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, Nbytes));
|
||||
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
|
||||
static_cast<const int*>(X_d), static_cast<const int*>(Y_d), Z_d, N);
|
||||
static_cast<const int*>(X_d), static_cast<const int*>(Y_d), Z_d, N);
|
||||
HIPCHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
@@ -72,6 +73,9 @@ int main() {
|
||||
HIPCHECK(hipFree(X_d));
|
||||
HIPCHECK(hipFree(Y_d));
|
||||
HIPCHECK(hipFree(Z_d));
|
||||
} else {
|
||||
std::cout<<"Machine does not seem to have P2P Capabilities, Empty Pass"<<std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
Regular → Executable
+7
@@ -40,6 +40,10 @@ int main() {
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
if (numDevices > 1) {
|
||||
|
||||
int canAccessPeer = 0;
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
|
||||
if (canAccessPeer) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
@@ -75,6 +79,9 @@ int main() {
|
||||
HIPCHECK(hipFree(X_d));
|
||||
HIPCHECK(hipFree(Y_d));
|
||||
HIPCHECK(hipFree(Z_d));
|
||||
} else {
|
||||
std::cout<<"Machine does not seem to have P2P Capabilities, Empty Pass"<<std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
Regular → Executable
+6
@@ -39,6 +39,9 @@ int main() {
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
if (numDevices > 1) {
|
||||
int canAccessPeer = 0;
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
|
||||
if (canAccessPeer) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
@@ -72,6 +75,9 @@ int main() {
|
||||
HIPCHECK(hipFree(X_d));
|
||||
HIPCHECK(hipFree(Y_d));
|
||||
HIPCHECK(hipFree(Z_d));
|
||||
} else {
|
||||
std::cout<<"Machine does not seem to have P2P Capabilities, Empty Pass"<<std::endl;
|
||||
}
|
||||
}
|
||||
passed();
|
||||
}
|
||||
|
||||
Regular → Executable
+7
@@ -42,6 +42,10 @@ int main() {
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
if (numDevices > 1) {
|
||||
|
||||
int canAccessPeer = 0;
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
|
||||
if (canAccessPeer) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
@@ -77,6 +81,9 @@ int main() {
|
||||
HIPCHECK(hipFree(X_d));
|
||||
HIPCHECK(hipFree(Y_d));
|
||||
HIPCHECK(hipFree(Z_d));
|
||||
} else {
|
||||
std::cout<<"Machine does not seem to have P2P Capabilities, Empty Pass"<<std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
Regular → Executable
+476
-22
@@ -18,9 +18,9 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Conformance test for checking functionality of
|
||||
* hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t
|
||||
* sizeBytes);
|
||||
* Different test for checking functionality of
|
||||
* hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,hipMemcpyKind kind,
|
||||
* hipStream_t stream);
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
@@ -31,27 +31,481 @@ THE SOFTWARE.
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
int main() {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM);
|
||||
#define test_failed(test_name) printf("%s %s FAILED!%s\n", KRED, #test_name, KNRM);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice,stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
class HipMemcpyWithStreamtests {
|
||||
public:
|
||||
// Test hipMemcpyWithStream with one streams and launch kernel in
|
||||
// that stream, verify the data
|
||||
void TestwithOnestream(void);
|
||||
// Test hipMemcpyWithStream with two streams and launch kernels in
|
||||
// two streams, verify the data
|
||||
void TestwithTwoStream(void);
|
||||
// Test hipMemcpyWithStream with one stream for each gpu and launch
|
||||
// kernels in each, verify the data
|
||||
void TestOnMultiGPUwithOneStream(void);
|
||||
// Test hipMemcpyWithStream to copy data from device to host (hipMemcpyDeviceToHost)
|
||||
void TestkindDtoH(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU
|
||||
void TestkindDtoD(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyHostToHost
|
||||
void TestkindHtoH(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault
|
||||
void TestkindDefault(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault for device to device transfer case
|
||||
void TestkindDefaultForDtoD(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device
|
||||
void TestDtoDonSameDevice(void);
|
||||
};
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
passed();
|
||||
void HipMemcpyWithStreamtests::TestwithOnestream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestwithTwoStream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int noOfstreams = 2;
|
||||
int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams];
|
||||
int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams];
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i], N, false);
|
||||
}
|
||||
|
||||
hipStream_t stream[noOfstreams];
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestDtoDonSameDevice(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int noOfstreams = 2;
|
||||
int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams];
|
||||
int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams];
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
|
||||
hipStream_t stream[noOfstreams];
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMalloc(&A_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[1], Nbytes));
|
||||
C_h[1] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[1] != NULL);
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
|
||||
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
|
||||
if (A_d[1]) {
|
||||
HIPCHECK(hipFree(A_d[1]));
|
||||
}
|
||||
if (B_d[1]) {
|
||||
HIPCHECK(hipFree(B_d[1]));
|
||||
}
|
||||
if (C_d[1]) {
|
||||
HIPCHECK(hipFree(C_d[1]));
|
||||
}
|
||||
if (C_h[1]) {
|
||||
free(C_h[1]);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestOnMultiGPUwithOneStream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i], N, false);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDtoH(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDtoD(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
|
||||
int canAccessPeer = 0;
|
||||
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
|
||||
if (!canAccessPeer) {
|
||||
std::cout<<"Machine does not seem to have P2P Capabilities"<<std::endl;
|
||||
return;
|
||||
}
|
||||
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
// Initialize and create the host and device elements for first device
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i))
|
||||
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
|
||||
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[i] != NULL);
|
||||
}
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDefault(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDefaultForDtoD(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
// Initialize and create the host and device elements for first device
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
|
||||
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[i] != NULL);
|
||||
}
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDefault, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindHtoH(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_h, *B_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
// Allocate memory to A_h and B_h
|
||||
A_h = static_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(A_h != NULL);
|
||||
B_h = static_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(B_h != NULL);
|
||||
|
||||
for (size_t i = 0; i < N; ++i) {
|
||||
if (A_h) (A_h)[i] = 3.146f + i; // Pi
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
HIPASSERT(A_h[i] == B_h[i]);
|
||||
}
|
||||
|
||||
if (A_h) {
|
||||
free(A_h);
|
||||
}
|
||||
if (B_h) {
|
||||
free(B_h);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
int main() {
|
||||
HipMemcpyWithStreamtests tests;
|
||||
tests.TestwithOnestream();
|
||||
test_passed(TestwithOnestream);
|
||||
tests.TestwithTwoStream();
|
||||
test_passed(TestwithTwoStream);
|
||||
tests.TestkindDtoH();
|
||||
test_passed(TestkindsDtoH);
|
||||
tests.TestkindDefault();
|
||||
test_passed(TestkindDefault);
|
||||
tests.TestDtoDonSameDevice();
|
||||
test_passed(TestDtoDonSameDevice);
|
||||
tests.TestOnMultiGPUwithOneStream();
|
||||
test_passed(TestOnMultiGPUwithOneStream);
|
||||
tests.TestkindDtoD();
|
||||
test_passed(TestkindDtoD);
|
||||
tests.TestkindDefaultForDtoD();
|
||||
test_passed(TestkindDefaultForDtoD);
|
||||
tests.TestkindHtoH();
|
||||
test_passed(TestkindsHtoH);
|
||||
}
|
||||
|
||||
@@ -0,0 +1,659 @@
|
||||
/*
|
||||
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Different test for checking functionality of
|
||||
* hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
|
||||
* hipMemcpyKind kind, hipStream_t stream);
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <vector>
|
||||
#include <thread>
|
||||
#include <chrono>
|
||||
#include "test_common.h"
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN << 2
|
||||
#define THREADS 2
|
||||
#define MAX_THREADS 16
|
||||
|
||||
|
||||
#define test_passed(test_name) printf("%s %s PASSED!%s\n", \
|
||||
KGRN, #test_name, KNRM);
|
||||
#define test_failed(test_name) printf("%s %s FAILED!%s\n", \
|
||||
KRED, #test_name, KNRM);
|
||||
|
||||
enum class ops
|
||||
{ TestwithOnestream,
|
||||
TestwithTwoStream,
|
||||
TestOnMultiGPUwithOneStream,
|
||||
TestkindDtoH,
|
||||
TestkindDtoD,
|
||||
TestkindHtoH,
|
||||
TestkindDefault,
|
||||
TestkindDefaultForDtoD,
|
||||
TestDtoDonSameDevice,
|
||||
END_OF_LIST
|
||||
};
|
||||
|
||||
|
||||
class HipMemcpyWithStreamMultiThreadtests {
|
||||
// Test hipMemcpyWithStream with one streams and launch kernel in
|
||||
// that stream, verify the data.
|
||||
void TestwithOnestream(void);
|
||||
// Test hipMemcpyWithStream with two streams and launch kernels in
|
||||
// two streams, verify the data.
|
||||
void TestwithTwoStream(void);
|
||||
// Test hipMemcpyWithStream with one stream for each gpu and launch
|
||||
// kernels in each, verify the data
|
||||
void TestOnMultiGPUwithOneStream(void);
|
||||
// Test hipMemcpyWithStream to copy data from
|
||||
// device to host (hipMemcpyDeviceToHost).
|
||||
void TestkindDtoH(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU.
|
||||
void TestkindDtoD(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyHostToHost.
|
||||
void TestkindHtoH(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault.
|
||||
void TestkindDefault(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault for
|
||||
// device to device transfer case.
|
||||
void TestkindDefaultForDtoD(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device.
|
||||
void TestDtoDonSameDevice(void);
|
||||
|
||||
public:
|
||||
// run all the tests on multithreaded.
|
||||
void TestwithMultiThreaded(ops op);
|
||||
};
|
||||
|
||||
struct joinable_thread : std::thread {
|
||||
template <class... Xs>
|
||||
explicit joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...)
|
||||
{} // NOLINT
|
||||
|
||||
joinable_thread& operator=(joinable_thread&& other) = default;
|
||||
joinable_thread(joinable_thread&& other) = default;
|
||||
|
||||
~joinable_thread() {
|
||||
if (this->joinable())
|
||||
this->join();
|
||||
}
|
||||
};
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes,
|
||||
hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes,
|
||||
hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int noOfstreams = 2;
|
||||
int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams];
|
||||
int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams];
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i],
|
||||
&A_h[i], &B_h[i], &C_h[i], N, false);
|
||||
}
|
||||
|
||||
hipStream_t stream[noOfstreams];
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream[i], static_cast<const int*>(A_d[i]),
|
||||
static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int noOfstreams = 2;
|
||||
int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams];
|
||||
int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams];
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0],
|
||||
&A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
|
||||
hipStream_t stream[noOfstreams];
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMalloc(&A_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[1], Nbytes));
|
||||
C_h[1] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[1] != NULL);
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes,
|
||||
hipMemcpyDeviceToDevice, stream[1]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes,
|
||||
hipMemcpyDeviceToDevice, stream[1]));
|
||||
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream[i], static_cast<const int*>(A_d[i]),
|
||||
static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
|
||||
if (A_d[1]) {
|
||||
HIPCHECK(hipFree(A_d[1]));
|
||||
}
|
||||
if (B_d[1]) {
|
||||
HIPCHECK(hipFree(B_d[1]));
|
||||
}
|
||||
if (C_d[1]) {
|
||||
HIPCHECK(hipFree(C_d[1]));
|
||||
}
|
||||
if (C_h[1]) {
|
||||
free(C_h[1]);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestOnMultiGPUwithOneStream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i],
|
||||
&A_h[i], &B_h[i], &C_h[i], N, false);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream[i], static_cast<const int*>(A_d[i]),
|
||||
static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes,
|
||||
hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes,
|
||||
hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes,
|
||||
hipMemcpyDeviceToHost, stream));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
// Initialize and create the host and device elements for first device
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0],
|
||||
&A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i))
|
||||
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
|
||||
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[i] != NULL);
|
||||
}
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
// Copying device data from 1st GPU to the rest of the the GPUs that is
|
||||
// numDevices in the setup. 1st GPU start numbering from 0,1,2..n etc.
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes,
|
||||
hipMemcpyDeviceToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes,
|
||||
hipMemcpyDeviceToDevice, stream[i]));
|
||||
}
|
||||
|
||||
|
||||
// Launching the kernel including the 1st GPU to the no of GPUs present
|
||||
// in the setup. 1st GPU start numbering from 0,1,2..n etc.
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream[i], static_cast<const int*>(A_d[i]),
|
||||
static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindDefault(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream, static_cast<const int*>(A_d),
|
||||
static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// Test case will not run on single GPU setup.
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
// Initialize and create the host and device elements for first device
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0],
|
||||
&A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
|
||||
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[i] != NULL);
|
||||
}
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
|
||||
hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
// Copying device data from 1st GPU to the rest of the the GPUs
|
||||
// using hipMemcpyDefault kind that is numDevices in the setup.
|
||||
// 1st GPU start numbering from 0,1,2..n etc.
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes,
|
||||
hipMemcpyDefault, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes,
|
||||
hipMemcpyDefault, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, stream[i], static_cast<const int*>(A_d[i]),
|
||||
static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
// Output of each GPU is getting validated with input of 1st GPU.
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_h, *B_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
// Allocate memory to A_h and B_h
|
||||
A_h = static_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(A_h != NULL);
|
||||
B_h = static_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(B_h != NULL);
|
||||
|
||||
for (size_t i = 0; i < N; ++i) {
|
||||
if (A_h) {
|
||||
(A_h)[i] = 3.146f + i; // Pi
|
||||
}
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
HIPASSERT(A_h[i] == B_h[i]);
|
||||
}
|
||||
|
||||
if (A_h) {
|
||||
free(A_h);
|
||||
}
|
||||
if (B_h) {
|
||||
free(B_h);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestwithMultiThreaded(ops op) {
|
||||
int n = min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS);
|
||||
std::vector<joinable_thread> threads;
|
||||
|
||||
for (uint32_t i = 0; i < n; i++) {
|
||||
threads.emplace_back(std::thread{[&] {
|
||||
switch ( op ) {
|
||||
case ops::TestwithOnestream:
|
||||
TestwithOnestream();
|
||||
break;
|
||||
case ops::TestwithTwoStream:
|
||||
TestwithTwoStream();
|
||||
break;
|
||||
case ops::TestkindDtoH:
|
||||
TestkindDtoH();
|
||||
break;
|
||||
case ops::TestkindHtoH:
|
||||
TestkindHtoH();
|
||||
break;
|
||||
case ops::TestkindDtoD:
|
||||
TestkindDtoD();
|
||||
break;
|
||||
case ops::TestOnMultiGPUwithOneStream:
|
||||
TestOnMultiGPUwithOneStream();
|
||||
break;
|
||||
case ops::TestkindDefault:
|
||||
TestkindDefault();
|
||||
break;
|
||||
case ops::TestkindDefaultForDtoD:
|
||||
TestkindDefaultForDtoD();
|
||||
break;
|
||||
case ops::TestDtoDonSameDevice:
|
||||
TestDtoDonSameDevice();
|
||||
break;
|
||||
default:{}
|
||||
}
|
||||
}});
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main() {
|
||||
HipMemcpyWithStreamMultiThreadtests tests;
|
||||
for (int op = static_cast<int>(ops::TestwithOnestream);
|
||||
op < static_cast<int>(ops::END_OF_LIST); ++op) {
|
||||
tests.TestwithMultiThreaded(static_cast<ops>(op));
|
||||
switch ( static_cast<ops>(op) ) {
|
||||
case ops::TestwithOnestream:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestwithOnestream);
|
||||
break;
|
||||
case ops::TestwithTwoStream:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestwithTwoStream);
|
||||
break;
|
||||
case ops::TestkindDtoH:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestkindDtoH);
|
||||
break;
|
||||
case ops::TestkindHtoH:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestkindHtoH);
|
||||
break;
|
||||
case ops::TestkindDtoD:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestkindDtoD);
|
||||
break;
|
||||
case ops::TestOnMultiGPUwithOneStream:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestOnMultiGPUwithOneStream);
|
||||
break;
|
||||
case ops::TestkindDefault:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestkindDefault);
|
||||
break;
|
||||
case ops::TestkindDefaultForDtoD:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestkindDefaultForDtoD);
|
||||
break;
|
||||
case ops::TestDtoDonSameDevice:
|
||||
test_passed(HipMemcpyWithStreamMultiThreadtests
|
||||
::TestDtoDonSameDevice);
|
||||
break;
|
||||
default: { test_failed("No Operation to done with API"); }
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -1,119 +1,234 @@
|
||||
/*
|
||||
Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
// Simple test for memset.
|
||||
// Also serves as a template for other tests.
|
||||
// Test for hipMemset2D functionality for different width and height values
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* TEST: %t
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST_NAMED: %t hipMemset2D-basic
|
||||
* TEST_NAMED: %t hipMemset2D-dim1 --width2D 10 --height2D 10 --memsetWidth 4 --memsetHeight 4
|
||||
* TEST_NAMED: %t hipMemset2D-dim2 --width2D 100 --height2D 100 --memsetWidth 20 --memsetHeight 40
|
||||
* TEST_NAMED: %t hipMemset2D-dim3 --width2D 256 --height2D 256 --memsetWidth 39 --memsetHeight 19
|
||||
* TEST_NAMED: %t hipMemset2D-zeroH --width2D 100 --height2D 100 --memsetWidth 20 --memsetHeight 0
|
||||
* TEST_NAMED: %t hipMemset2D-zeroW --width2D 100 --height2D 100 --memsetWidth 0 --memsetHeight 20
|
||||
* TEST_NAMED: %t hipMemset2D-zeroW*H --width2D 100 --height2D 100 --memsetWidth 0 --memsetHeight 0
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
bool testhipMemset2D(int memsetval,int p_gpuDevice)
|
||||
{
|
||||
size_t numH = 256;
|
||||
size_t numW = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW* numH;
|
||||
// Check hipMemset2D functionality
|
||||
bool testhipMemset2D(int memsetval, int p_gpuDevice) {
|
||||
bool testResult = true;
|
||||
size_t numH = 256;
|
||||
size_t numW = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW* numH;
|
||||
printf("testhipMemset2D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width ,
|
||||
numH));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
HIPASSERT(A_h != NULL);
|
||||
|
||||
printf ("testhipMemset2D memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
bool testResult = true;
|
||||
HIPCHECK (hipMallocPitch((void**)&A_d, &pitch_A, width , numH));
|
||||
A_h = (char*)malloc(sizeElements);
|
||||
HIPASSERT(A_h != NULL);
|
||||
for (size_t i=0; i<elements; i++) {
|
||||
A_h[i] = 1;
|
||||
for (size_t i=0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
|
||||
HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i=0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
printf("testhipMemset2D mismatch at index:%d computed:%02x, memsetval:"
|
||||
"%02x\n", i, static_cast<int>(A_h[i]), static_cast<int>(memsetval));
|
||||
break;
|
||||
}
|
||||
HIPCHECK ( hipMemset2D(A_d, pitch_A, memsetval, numW, numH) );
|
||||
HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
for (int i=0; i<elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
printf("testhipMemset2D mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
|
||||
break;
|
||||
}
|
||||
}
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
return testResult;
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
bool testhipMemset2DAsync(int memsetval,int p_gpuDevice)
|
||||
{
|
||||
size_t numH = 256;
|
||||
size_t numW = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW* numH;
|
||||
// Check hipMemset2DAsync functionality
|
||||
bool testhipMemset2DAsync(int memsetval, int p_gpuDevice) {
|
||||
size_t numH = 256;
|
||||
size_t numW = 256;
|
||||
size_t pitch_A;
|
||||
size_t width = numW * sizeof(char);
|
||||
size_t sizeElements = width * numH;
|
||||
size_t elements = numW * numH;
|
||||
printf("testhipMemset2DAsync memsetval=%2x device=%d\n", memsetval,
|
||||
p_gpuDevice);
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
bool testResult = true;
|
||||
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width , numH));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
HIPASSERT(A_h != NULL);
|
||||
|
||||
printf ("testhipMemset2DAsync memsetval=%2x device=%d\n", memsetval, p_gpuDevice);
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
bool testResult = true;
|
||||
for (size_t i=0; i < elements; i++) {
|
||||
A_h[i] = 1;
|
||||
}
|
||||
|
||||
HIPCHECK (hipMallocPitch((void**)&A_d, &pitch_A, width , numH));
|
||||
A_h = (char*)malloc(sizeElements);
|
||||
HIPASSERT(A_h != NULL);
|
||||
for (size_t i=0; i<elements; i++) {
|
||||
A_h[i] = 1;
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i=0; i < elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
printf("testhipMemset2DAsync mismatch at index:%d computed:%02x, memsetval:"
|
||||
"%02x\n", i, static_cast<int>(A_h[i]), static_cast<int>(memsetval));
|
||||
break;
|
||||
}
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
for (int i=0; i<elements; i++) {
|
||||
if (A_h[i] != memsetval) {
|
||||
testResult = false;
|
||||
printf("testhipMemset2DAsync mismatch at index:%d computed:%02x, memsetval:%02x\n", i, (int)A_h[i], (int)memsetval);
|
||||
break;
|
||||
}
|
||||
}
|
||||
hipFree(A_d);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
return testResult;
|
||||
hipFree(A_d);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
bool testResult = true;
|
||||
int width2D = 20;
|
||||
int height2D = 20;
|
||||
int memsetWidth = 20;
|
||||
int memsetHeight = 20;
|
||||
|
||||
int parseExtraArguments(int argc, char* argv[]) {
|
||||
int i = 0;
|
||||
for (i = 1; i < argc; i++) {
|
||||
const char* arg = argv[i];
|
||||
if (!strcmp(arg, " ")) {
|
||||
// skip NULL args.
|
||||
} else if (!strcmp(arg, "--width2D")) {
|
||||
if (++i >= argc || !HipTest::parseInt(argv[i], &width2D)) {
|
||||
failed("Bad width2D argument");
|
||||
}
|
||||
} else if (!strcmp(arg, "--height2D")) {
|
||||
if (++i >= argc || !HipTest::parseInt(argv[i], &height2D)) {
|
||||
failed("Bad height2D argument");
|
||||
}
|
||||
} else if (!strcmp(arg, "--memsetWidth")) {
|
||||
if (++i >= argc || !HipTest::parseInt(argv[i], &memsetWidth)) {
|
||||
failed("Bad memsetWidth argument");
|
||||
}
|
||||
} else if (!strcmp(arg, "--memsetHeight")) {
|
||||
if (++i >= argc || !HipTest::parseInt(argv[i], &memsetHeight)) {
|
||||
failed("Bad memsetHeight argument");
|
||||
}
|
||||
} else {
|
||||
failed("Bad argument");
|
||||
}
|
||||
}
|
||||
return i;
|
||||
}
|
||||
|
||||
// Memset random dimensions
|
||||
bool testMemset2DPartial(int memsetval, int p_gpuDevice) {
|
||||
bool testResult = true;
|
||||
size_t NUM_H = height2D;
|
||||
size_t NUM_W = width2D;
|
||||
size_t Nbytes = N*sizeof(char);
|
||||
size_t pitch_A;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
size_t sizeElements = width * NUM_H;
|
||||
size_t elements = NUM_W * NUM_H;
|
||||
char *A_d;
|
||||
char *A_h;
|
||||
printf("testhipMemset2DPartial memsetval=%2x device=%d\n", memsetval,
|
||||
p_gpuDevice);
|
||||
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
|
||||
width, NUM_H));
|
||||
hipError_t e;
|
||||
int index;
|
||||
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
HIPASSERT(A_h != NULL);
|
||||
|
||||
for (index = 0; index < sizeElements; index++) {
|
||||
A_h[0] = 'c';
|
||||
}
|
||||
|
||||
printf("2D Dimension: %zuX%zu, MemsetWidth:%d, memsetHeight:%d\n",
|
||||
NUM_W, NUM_H, memsetWidth, memsetHeight);
|
||||
e = hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight);
|
||||
HIPASSERT(e == hipSuccess);
|
||||
|
||||
HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, NUM_W, NUM_H,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int row = 0; row < memsetHeight; row++) {
|
||||
for (int column = 0; column < memsetWidth; column++) {
|
||||
if (A_h[(row * width) + column] != memsetval) {
|
||||
printf("A_h[%d][%d] did not match %d", row, column, memsetval);
|
||||
testResult = false;
|
||||
}
|
||||
}
|
||||
}
|
||||
hipFree(A_d);
|
||||
free(A_h);
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
int extraArgs = 0;
|
||||
bool testResult = true;
|
||||
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
extraArgs = HipTest::parseStandardArguments(argc, argv, false);
|
||||
parseExtraArguments(extraArgs, argv);
|
||||
|
||||
if (extraArgs == 1) {
|
||||
testResult &= testhipMemset2D(memsetval, p_gpuDevice);
|
||||
testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice);
|
||||
if(testResult){
|
||||
passed();
|
||||
if (!(testResult)) {
|
||||
printf("hipMemset2D failed\n");
|
||||
}
|
||||
testResult &= testhipMemset2DAsync(memsetval, p_gpuDevice);
|
||||
if (!(testResult)) {
|
||||
printf("hipMemset2DAsync failed\n");
|
||||
}
|
||||
} else if (extraArgs == 9) {
|
||||
testResult &= testMemset2DPartial(memsetval, p_gpuDevice);
|
||||
if (!(testResult)) {
|
||||
printf("hipMemset2D at random dimensions failed\n");
|
||||
}
|
||||
} else {
|
||||
failed("Wrong Arguments for test\n");
|
||||
}
|
||||
|
||||
if (testResult) {
|
||||
passed();
|
||||
} else {
|
||||
failed("one or more hipMemset2D tests failed");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,173 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
//
|
||||
// Test to verify
|
||||
// a) Order of execution of device kernel and hipMemset2DAsync api
|
||||
// b) hipMemSet2DAsync execution in multiple threads
|
||||
//
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#define NUM_THREADS 1000
|
||||
#define ITER 100
|
||||
#define NUM_H 256
|
||||
#define NUM_W 256
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
hipStream_t stream;
|
||||
|
||||
bool testResult = true;
|
||||
char *A_d, *A_h, *B_d, *B_h, *C_d;
|
||||
int validateCount;
|
||||
|
||||
size_t pitch_A, pitch_B, pitch_C;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
size_t sizeElements = width * NUM_H;
|
||||
size_t elements = NUM_W * NUM_H;
|
||||
|
||||
/*
|
||||
* Square each element in the array B and write to array C.
|
||||
*/
|
||||
|
||||
__global__ void
|
||||
vector_square(char* B_d, char* C_d, size_t elements) {
|
||||
for (int i=0 ; i < elements ; i++) {
|
||||
C_d[i] = B_d[i] * B_d[i];
|
||||
}
|
||||
}
|
||||
|
||||
void memAllocate() {
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, NUM_H));
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d), &pitch_B, width, NUM_H));
|
||||
A_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
HIPASSERT(A_h != NULL);
|
||||
B_h = reinterpret_cast<char*>(malloc(sizeElements));
|
||||
HIPASSERT(B_h != NULL);
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&C_d), &pitch_C, width, NUM_H));
|
||||
|
||||
for (int i = 0 ; i < elements ; i++) {
|
||||
B_h[i] = i;
|
||||
}
|
||||
HIPCHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H,
|
||||
hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
}
|
||||
|
||||
void memDeallocate() {
|
||||
HIPCHECK(hipFree(A_d)); HIPCHECK(hipFree(B_d)); HIPCHECK(hipFree(C_d));
|
||||
free(A_h); free(B_h);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch,
|
||||
size_t width) {
|
||||
HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream));
|
||||
HIPCHECK(hipMemcpy2DAsync(A_h, width, A_d, pitch, NUM_W, NUM_H,
|
||||
hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
bool testhipMemset2DAsyncWithKernel() {
|
||||
validateCount = 0;
|
||||
memAllocate();
|
||||
printf("info: Launching vector_square kernel and hipMemset2DAsync "
|
||||
"simultaneously\n");
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
stream, B_d, C_d, elements);
|
||||
|
||||
HIPCHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (int p = 0 ; p < elements ; p++) {
|
||||
if (A_h[p] == memsetval) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
testResult = (validateCount == (ITER * elements)) ? true : false;
|
||||
memDeallocate();
|
||||
return testResult;
|
||||
}
|
||||
|
||||
bool testhipMemset2DAsyncMultiThread() {
|
||||
validateCount = 0;
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
memAllocate();
|
||||
|
||||
printf("info: Queueing up hipMemset2DAsync jobs over multiple threads\n");
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, A_h, pitch_A,
|
||||
width);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, B_h, pitch_A,
|
||||
width);
|
||||
}
|
||||
}
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
for (int k = 0 ; k < elements ; k++) {
|
||||
if ((A_h[k] == memsetval) && (B_h[k] == memsetval)) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
memDeallocate();
|
||||
testResult = (validateCount == (ITER * elements)) ? true : false;
|
||||
return testResult;
|
||||
}
|
||||
|
||||
int main() {
|
||||
bool testResult = true;
|
||||
|
||||
testResult &= testhipMemset2DAsyncWithKernel();
|
||||
if (testResult) {
|
||||
printf("Kernel and hipMemset2DAsync executed in correct order!\n");
|
||||
} else {
|
||||
printf("Kernel and hipMemset2DAsync order of execution failed\n");
|
||||
}
|
||||
|
||||
testResult &= testhipMemset2DAsyncMultiThread();
|
||||
if (testResult) {
|
||||
printf("hipMemset2DAsync jobs on all threads finished successfully!\n");
|
||||
passed();
|
||||
} else {
|
||||
printf("hipMemset2DAsync failed in multi thread scenario\n");
|
||||
}
|
||||
|
||||
if (testResult) {
|
||||
passed();
|
||||
} else {
|
||||
failed("One or more tests failed\n");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,191 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Test for checking order of execution of device kernel and
|
||||
* hipMemsetAsync apis on all gpus
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#define ITER 10
|
||||
#define N 1024 * 1024
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
template <typename T>
|
||||
__global__ void vector_square(T* B_d, T* C_d, size_t M) {
|
||||
for (int i=0 ; i < M ; i++) {
|
||||
C_d[i] = B_d[i] * B_d[i];
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
class MemSetTest {
|
||||
public:
|
||||
T *A_h, *B_d, *B_h, *C_d;
|
||||
T memSetVal;
|
||||
size_t Nbytes;
|
||||
bool testResult = true;
|
||||
int validateCount = 0;
|
||||
hipStream_t stream;
|
||||
|
||||
void memAllocate(T memSetValue) {
|
||||
memSetVal = memSetValue;
|
||||
Nbytes = N * sizeof(T);
|
||||
|
||||
A_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIPASSERT(A_h != NULL);
|
||||
HIPCHECK(hipMalloc(&B_d , Nbytes));
|
||||
B_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIPASSERT(B_h != NULL);
|
||||
HIPCHECK(hipMalloc(&C_d , Nbytes));
|
||||
|
||||
for (int i = 0 ; i < N ; i++) {
|
||||
B_h[i] = i;
|
||||
}
|
||||
HIPCHECK(hipMemcpy(B_d , B_h , Nbytes , hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
}
|
||||
|
||||
void memDeallocate() {
|
||||
HIPCHECK(hipFree(B_d)); HIPCHECK(hipFree(C_d));
|
||||
free(B_h); free(A_h);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void validateExecutionOrder() {
|
||||
for (int p = 0 ; p < N ; p++) {
|
||||
if (A_h[p] == memSetVal) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
bool resultAfterAllIterations() {
|
||||
testResult = (validateCount == (ITER * N)) ? true : false;
|
||||
memDeallocate();
|
||||
return testResult;
|
||||
}
|
||||
};
|
||||
|
||||
bool testhipMemsetAsyncWithKernel() {
|
||||
MemSetTest <char> obj;
|
||||
obj.memAllocate(memsetval);
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIPCHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream));
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
HIPCHECK(hipMemcpy(obj.A_h , obj.C_d , obj.Nbytes , hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD32AsyncWithKernel() {
|
||||
MemSetTest <int32_t> obj;
|
||||
obj.memAllocate(memsetD32val);
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIPCHECK(hipMemsetD32Async(obj.C_d , obj.memSetVal , N , obj.stream));
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
HIPCHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD16AsyncWithKernel() {
|
||||
MemSetTest <int16_t> obj;
|
||||
obj.memAllocate(memsetD16val);
|
||||
for (int k = 0 ; k < ITER ; k++) {
|
||||
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIPCHECK(hipMemsetD16Async(obj.C_d , obj.memSetVal , N , obj.stream));
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
HIPCHECK(hipMemcpy(obj.A_h , obj.C_d, obj.Nbytes , hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD8AsyncWithKernel() {
|
||||
MemSetTest <char> obj;
|
||||
obj.memAllocate(memsetD8val);
|
||||
for (int k = 0; k < ITER; k++) {
|
||||
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
obj.stream, obj.B_d, obj.C_d, N);
|
||||
HIPCHECK(hipMemsetD8Async(obj.C_d, obj.memSetVal, N, obj.stream));
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
HIPCHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
obj.validateExecutionOrder();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
int main() {
|
||||
bool testResult = true;
|
||||
int numDevices = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
printf("total number of gpus in the system: %d\n", numDevices);
|
||||
|
||||
for (int i = 0; i < numDevices; i++) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
printf("test running on gpu %d\n", i);
|
||||
|
||||
testResult &= testhipMemsetAsyncWithKernel();
|
||||
if (!(testResult)) {
|
||||
printf("Mismatch in order of execution of hipMemsetAsync and kernel\n");
|
||||
}
|
||||
|
||||
testResult &= testhipMemsetD32AsyncWithKernel();
|
||||
if (!(testResult)) {
|
||||
printf("Mismatch in order of execution of hipMemsetD32Async and kernel\n");
|
||||
}
|
||||
|
||||
testResult &= testhipMemsetD16AsyncWithKernel();
|
||||
if (!(testResult)) {
|
||||
printf("Mismatch in order of execution of hipMemsetD16Async and kernel\n");
|
||||
}
|
||||
|
||||
testResult &= testhipMemsetD8AsyncWithKernel();
|
||||
if (!(testResult)) {
|
||||
printf("Mismatch in order of execution of hipMemsetD8Async and kernel\n");
|
||||
}
|
||||
}
|
||||
|
||||
if (testResult) {
|
||||
printf("Execution order of Kernel and hipMemsetAsync apis on "
|
||||
"all gpus is correct!\n");
|
||||
passed();
|
||||
} else {
|
||||
failed("One or more hipMemsetAsync tests failed\n");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,247 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Test that validates functionality of hipmemsetAsync apis over multi threads
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#define NUM_THREADS 50
|
||||
#define ITER 50
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
template <typename T>
|
||||
class MemSetTest {
|
||||
public:
|
||||
T *A_h, *A_d, *B_h;
|
||||
T memSetVal;
|
||||
size_t Nbytes;
|
||||
bool testResult = true;
|
||||
int validateCount = 0;
|
||||
hipStream_t stream;
|
||||
|
||||
void memAllocate(T memSetValue) {
|
||||
memSetVal = memSetValue;
|
||||
Nbytes = N * sizeof(T);
|
||||
|
||||
A_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIPASSERT(A_h != NULL);
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
B_h = reinterpret_cast<T*>(malloc(Nbytes));
|
||||
HIPASSERT(B_h != NULL);
|
||||
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
}
|
||||
|
||||
void threadCompleteStatus() {
|
||||
for (int k = 0 ; k < N ; k++) {
|
||||
if ((A_h[k] == memSetVal) && (B_h[k] == memSetVal)) {
|
||||
validateCount+= 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
bool resultAfterAllIterations() {
|
||||
memDeallocate();
|
||||
testResult = (validateCount == (ITER * N)) ? true: false;
|
||||
return testResult;
|
||||
}
|
||||
|
||||
void memDeallocate() {
|
||||
HIPCHECK(hipFree(A_d));
|
||||
free(A_h);
|
||||
free(B_h);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetAsync(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetAsync(A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetD32Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetD32Async(A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetD16Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetD16Async(A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void queueJobsForhipMemsetD8Async(T* A_d, T* A_h, T memSetVal, size_t Nbytes,
|
||||
hipStream_t stream) {
|
||||
HIPCHECK(hipMemsetD8Async(A_d, memSetVal, N, stream));
|
||||
HIPCHECK(hipMemcpyAsync(A_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
/* Queue hipMemsetAsync jobs on multiple threads and verify they all
|
||||
* finished on all threads successfully
|
||||
*/
|
||||
|
||||
bool testhipMemsetAsyncWithMultiThread() {
|
||||
MemSetTest <char> obj;
|
||||
obj.memAllocate(memsetval);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetAsync<char>, obj.A_d, obj.A_h,
|
||||
obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetAsync<char>, obj.A_d, obj.B_h,
|
||||
obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD32AsyncWithMultiThread() {
|
||||
MemSetTest <int32_t> obj;
|
||||
obj.memAllocate(memsetD32val);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD32Async<int32_t>, obj.A_d,
|
||||
obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD32Async<int32_t>, obj.A_d,
|
||||
obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD16AsyncWithMultiThread() {
|
||||
MemSetTest <int16_t> obj;
|
||||
obj.memAllocate(memsetD16val);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD16Async<int16_t>, obj.A_d,
|
||||
obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD16Async<int16_t>, obj.A_d,
|
||||
obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
bool testhipMemsetD8AsyncWithMultiThread() {
|
||||
MemSetTest <char> obj;
|
||||
obj.memAllocate(memsetD8val);
|
||||
std::thread t[NUM_THREADS];
|
||||
|
||||
for (int i = 0 ; i < ITER ; i++) {
|
||||
for (int k = 0 ; k < NUM_THREADS ; k++) {
|
||||
if (k%2) {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD8Async<char>, obj.A_d,
|
||||
obj.A_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
} else {
|
||||
t[k] = std::thread(queueJobsForhipMemsetD8Async<char>, obj.A_d,
|
||||
obj.B_h, obj.memSetVal, obj.Nbytes, obj.stream);
|
||||
}
|
||||
}
|
||||
for (int j = 0 ; j < NUM_THREADS ; j++) {
|
||||
t[j].join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(obj.stream));
|
||||
obj.threadCompleteStatus();
|
||||
}
|
||||
return obj.resultAfterAllIterations();
|
||||
}
|
||||
|
||||
int main() {
|
||||
bool testResult = true;
|
||||
printf("Queueing up hipMemSetAsync jobs on multiple threads"
|
||||
"and checking results\n");
|
||||
|
||||
testResult &= testhipMemsetAsyncWithMultiThread();
|
||||
if (!(testResult)) {
|
||||
printf("Thread execution did not complete for hipMemsetAsync\n");
|
||||
}
|
||||
|
||||
testResult &= testhipMemsetD32AsyncWithMultiThread();
|
||||
if (!(testResult)) {
|
||||
printf("Thread execution did not complete for hipMemsetD32Async\n");
|
||||
}
|
||||
|
||||
testResult &= testhipMemsetD16AsyncWithMultiThread();
|
||||
if (!(testResult)) {
|
||||
printf("Thread execution did not complete for hipMemsetD16Async\n");
|
||||
}
|
||||
testResult &= testhipMemsetD8AsyncWithMultiThread();
|
||||
if (!(testResult)) {
|
||||
printf("Thread execution did not complete for hipMemsetD8Async\n");
|
||||
}
|
||||
|
||||
if (testResult) {
|
||||
printf("All threads ran successfully for all hipMemsetAsync apis\n");
|
||||
passed();
|
||||
} else {
|
||||
failed("One or more tests failed\n");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,97 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// * To test invalid pointer to hipMemset* apis
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
#define N 50
|
||||
#define MEMSETVAL 0x42
|
||||
#define NUM_H 256
|
||||
#define NUM_W 256
|
||||
|
||||
int main() {
|
||||
size_t Nbytes = N*sizeof(char);
|
||||
size_t pitch_A;
|
||||
size_t width = NUM_W * sizeof(char);
|
||||
size_t sizeElements = width * NUM_H;
|
||||
size_t elements = NUM_W * NUM_H;
|
||||
char *A_d;
|
||||
|
||||
HIPCHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width , NUM_H));
|
||||
|
||||
hipError_t e;
|
||||
|
||||
e = hipMemset(NULL , MEMSETVAL , Nbytes);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetD32(NULL , MEMSETVAL , Nbytes);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetD16(NULL , MEMSETVAL , Nbytes);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetD8(NULL , MEMSETVAL , Nbytes);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetAsync(NULL , MEMSETVAL , Nbytes , 0);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetD32Async(NULL , MEMSETVAL , Nbytes, 0);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetD16Async(NULL , MEMSETVAL , Nbytes, 0);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemsetD8Async(NULL , MEMSETVAL , Nbytes, 0);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
/* Passing host pointer to hipMemset.Ticket SWDEV-243206 is open for this.
|
||||
* Disabling this test until the ticket is closed
|
||||
*
|
||||
char *A_h;
|
||||
A_h = (char*)malloc(Nbytes);
|
||||
e = hipMemset(A_h, MEMSETVAL , Nbytes);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
*/
|
||||
|
||||
/* Passing invalid pitch to hipMemset2D.Ticket SWDEV-243104 is open for this.
|
||||
* Disabling this test until the ticket is closed
|
||||
*
|
||||
e = hipMemset2D(A_d, 0, MEMSETVAL, NUM_W, NUM_H);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
|
||||
e = hipMemset2DAsync(A_d, 0, MEMSETVAL, NUM_W, NUM_H,0);
|
||||
HIPASSERT(e == hipErrorInvalidValue);
|
||||
*/
|
||||
|
||||
hipFree(A_d);
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,187 @@
|
||||
/*
|
||||
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
|
||||
// Testcase Description: This test case is used to verify if the callback
|
||||
// function called through hipStreamAddCallback() api completes the execution
|
||||
// in order as hipStreamAddCallback() api queued in their respective streams
|
||||
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <vector>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
|
||||
hipStream_t mystream1, mystream2;
|
||||
size_t Num = 4096;
|
||||
std::vector<int> Stream1_Order, Stream2_Order;
|
||||
|
||||
|
||||
__global__ void vector_square(float* C_d, float* A_d, size_t Num) {
|
||||
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = gputhread; i < Num; i += stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
|
||||
// Delay thread 1 only in the GPU
|
||||
if (gputhread == 1) {
|
||||
unsigned long long int wait_t = 3200000000, start = clock64(), cur;
|
||||
do {
|
||||
cur = clock64() - start;
|
||||
} while (cur < wait_t);
|
||||
}
|
||||
}
|
||||
|
||||
float *A_h, *C_h, *A_h1, *C_h1;
|
||||
|
||||
static void HIPRT_CB Callback_Stream1(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
std::cout << "Data mismatch in stream1 at: " << i << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
// Storing the int passed into this callback into Stream1_Order
|
||||
// this will help verify the order in which this Callback function
|
||||
// is called.
|
||||
Stream1_Order.push_back(*(reinterpret_cast<int*>(userData)));
|
||||
delete reinterpret_cast<int*>(userData);
|
||||
}
|
||||
|
||||
static void HIPRT_CB Callback_Stream2(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
if (C_h1[i] != A_h1[i] * A_h1[i]) {
|
||||
std::cout << "Data mismatch in stream2 at: " << i << std::endl;
|
||||
}
|
||||
}
|
||||
// Storing the int passed into this callback into Stream2_Order
|
||||
// this will help verify the order in which this Callback function
|
||||
// is called.
|
||||
Stream2_Order.push_back(*(reinterpret_cast<int*>(userData)));
|
||||
delete reinterpret_cast<int*>(userData);
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
float *A_d, *C_d;
|
||||
size_t Nbytes = Num * sizeof(float);
|
||||
|
||||
A_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
A_h1 = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h1 = reinterpret_cast<float*>(malloc(Nbytes));
|
||||
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
|
||||
// Fill with Phi + i
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
A_h1[i] = 1.618f + i;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream1, hipStreamNonBlocking));
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream2, hipStreamNonBlocking));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream1));
|
||||
|
||||
const unsigned threadsPerBlock = 256;
|
||||
const unsigned blocks = (Num + 255)/threadsPerBlock;
|
||||
int *ptr = NULL;
|
||||
int *ptr1 = NULL;
|
||||
// Queing jobs in both mystream1/2 followed by hipStreamAddCallback
|
||||
for (int i = 1; i < 5; ++i) {
|
||||
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock),
|
||||
0, mystream1, C_d, A_d, Num);
|
||||
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost,
|
||||
mystream1));
|
||||
ptr = new int;
|
||||
*ptr = i;
|
||||
HIPCHECK(hipStreamAddCallback(mystream1, Callback_Stream1,
|
||||
reinterpret_cast<void*>(ptr), 0));
|
||||
|
||||
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock),
|
||||
0, mystream2, C_d, A_d, Num);
|
||||
HIPCHECK(hipMemcpyAsync(C_h1, C_d, Nbytes,
|
||||
hipMemcpyDeviceToHost, mystream2));
|
||||
ptr1 = new int;
|
||||
*ptr1 = i;
|
||||
HIPCHECK(hipStreamAddCallback(mystream2, Callback_Stream2,
|
||||
reinterpret_cast<void*>(ptr1), 0));
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(mystream1));
|
||||
HIPCHECK(hipStreamSynchronize(mystream2));
|
||||
|
||||
HIPCHECK(hipStreamDestroy(mystream1));
|
||||
HIPCHECK(hipStreamDestroy(mystream2));
|
||||
|
||||
HIPCHECK(hipFree(A_d));
|
||||
HIPCHECK(hipFree(C_d));
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
free(A_h1);
|
||||
free(C_h1);
|
||||
|
||||
// Checking if Stream1_Order has ints in sequencial order or not
|
||||
int i = 1;
|
||||
for (auto itr=Stream1_Order.begin(); itr != Stream1_Order.end(); ++itr) {
|
||||
if (*itr != i) {
|
||||
printf("hipStreamAddCallBack() did not execute in sequence");
|
||||
printf(" in first stream\n");
|
||||
failed("Unexpected behavior!");
|
||||
}
|
||||
++i;
|
||||
}
|
||||
|
||||
// Checking if Stream2_Order has ints in sequencial order or not
|
||||
i = 1;
|
||||
for (auto itr=Stream2_Order.begin(); itr != Stream2_Order.end(); ++itr) {
|
||||
if (*itr != i) {
|
||||
printf("hipStreamAddCallBack() did not execute in sequence");
|
||||
printf(" in second stream\n");
|
||||
failed("Unexpected behavior!");
|
||||
}
|
||||
++i;
|
||||
}
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,180 @@
|
||||
/*
|
||||
Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// Testcase Description: Streams are launched in individual GPUs with different
|
||||
// kernel. Verify that all the kernels queued are executed before the callback.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <stdio.h>
|
||||
#include <unistd.h>
|
||||
#include <thread>
|
||||
#include <chrono>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
|
||||
size_t N_ELMTS = 4096;
|
||||
|
||||
// Data structure for holding and validating data
|
||||
struct gpu_data {
|
||||
int *int_ptr = NULL;
|
||||
int gpu;
|
||||
int acknowledge;
|
||||
};
|
||||
|
||||
enum {
|
||||
SUCCESS = 0,
|
||||
KERNEL_EXECUTION_MISMATCH,
|
||||
KERNEL_COMPUTATION_MISMATCH
|
||||
};
|
||||
|
||||
__global__ void Add_Data(int* A_d, size_t N_ELMTS) {
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = offset; i < N_ELMTS; i += stride) {
|
||||
// Increment the value of A_d[i] by 1
|
||||
A_d[i] = A_d[i] + 1;
|
||||
}
|
||||
}
|
||||
|
||||
// below kernel is just to load the gpu with multiple jobs
|
||||
__global__ void Square_plus_one(int* A_d, int* C_d, size_t N_ELMTS) {
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = offset; i < N_ELMTS; i += stride) {
|
||||
C_d[i] = A_d[i]*A_d[i] + 1;
|
||||
}
|
||||
}
|
||||
|
||||
static void HIPRT_CB Stream_Callback(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
gpu_data *ptr = reinterpret_cast<gpu_data *>(userData);
|
||||
|
||||
// int_ptr in the passed userData will contain the data copied from device to
|
||||
// host. Expected data in this field is the gpu ordinal.
|
||||
if (*((*ptr).int_ptr) != (*ptr).gpu + 1) {
|
||||
(*ptr).acknowledge = 100; // Assign unexpected value to indicate fail
|
||||
} else {
|
||||
(*ptr).acknowledge = (*ptr).gpu; // Assign the gpu ordinal received
|
||||
}
|
||||
}
|
||||
|
||||
void launch_gpu(int gpu_ordinal) {
|
||||
HIPCHECK(hipSetDevice(gpu_ordinal));
|
||||
int *A_d, *A_h, *C_h, *C_d;
|
||||
size_t Nbytes = N_ELMTS * sizeof(int), Data_mismatch = 0;
|
||||
bool cb = false;
|
||||
A_h = (int *)malloc(Nbytes);
|
||||
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (int *)malloc(Nbytes);
|
||||
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
|
||||
// Fill with 0
|
||||
for (size_t i = 0; i < N_ELMTS; i++) {
|
||||
A_h[i] = 0;
|
||||
}
|
||||
|
||||
// setting gpu value in the struct object
|
||||
gpu_data *ptr = new gpu_data;
|
||||
ptr->int_ptr = C_h;
|
||||
ptr->gpu = gpu_ordinal;
|
||||
ptr->acknowledge = 100;
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
hipStream_t mystream;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
|
||||
|
||||
const unsigned threadsPerBlock = 256;
|
||||
const unsigned blocks = (N_ELMTS + 255)/threadsPerBlock;
|
||||
|
||||
// A_d is initialized to 0. Add_Data kernel does A_d = A_d + 1
|
||||
// The Add_data kernel is called 1 time for gpu0, 2 times for gpu1 etc.
|
||||
// At the end of the loop, A_d should have the gpu_ordinal number
|
||||
for (int i = 0; i < gpu_ordinal + 1; i++) {
|
||||
hipLaunchKernelGGL(Add_Data, dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
mystream, A_d, N_ELMTS);
|
||||
hipLaunchKernelGGL(Square_plus_one, 1, 1, 0, mystream, A_d, C_d, N_ELMTS);
|
||||
}
|
||||
HIPCHECK(hipMemcpyAsync(C_h, A_d, Nbytes, hipMemcpyDeviceToHost, mystream));
|
||||
|
||||
// Pass the ptr as user data which contains the gpu_ordinal, default value
|
||||
// for ack and the data that is copied to host
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Stream_Callback,
|
||||
reinterpret_cast<void *>(ptr), 0));
|
||||
HIPCHECK(hipStreamSynchronize(mystream));
|
||||
|
||||
HIPCHECK(hipFree(A_d));
|
||||
HIPCHECK(hipFree(C_d));
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
|
||||
int result = SUCCESS;
|
||||
if (C_h[0] != gpu_ordinal + 1) {
|
||||
result = KERNEL_EXECUTION_MISMATCH;
|
||||
}
|
||||
|
||||
if (ptr->gpu != ptr->acknowledge) {
|
||||
result = KERNEL_COMPUTATION_MISMATCH;
|
||||
}
|
||||
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
free(ptr);
|
||||
|
||||
if (result == KERNEL_EXECUTION_MISMATCH) {
|
||||
failed("Number of kernels expected to be executed does not match");
|
||||
} else if (result == KERNEL_COMPUTATION_MISMATCH) {
|
||||
failed("Mismatch found in the result of the computation!");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main() {
|
||||
int gpu_cnt = 0;
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&gpu_cnt));
|
||||
if (gpu_cnt < 2) {
|
||||
printf("Minimum of 2 gpus are needed for this test, skipping the test\n");
|
||||
passed();
|
||||
}
|
||||
|
||||
std::thread T[gpu_cnt];
|
||||
|
||||
// Launching threads for each GPU
|
||||
for (int i = 0; i < gpu_cnt; i++) {
|
||||
T[i] = std::thread(launch_gpu, i);
|
||||
}
|
||||
|
||||
for (int i=0; i < gpu_cnt; i++) {
|
||||
T[i].join();
|
||||
}
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,130 @@
|
||||
/*
|
||||
* Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
* */
|
||||
|
||||
// Testcase Description:: This test case is used to check if the runtime is ok
|
||||
// when hipStreamAddCallback() is called back to back multiple calls
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <unistd.h>
|
||||
#include <mutex>
|
||||
#include <atomic>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
#define NUM_CALLS 1000
|
||||
|
||||
hipStream_t mystream;
|
||||
size_t Num = 4096;
|
||||
std::atomic<size_t>Cb_count{0}, Data_mismatch{0};
|
||||
float *A_h, *C_h;
|
||||
|
||||
__global__ void vector_square(float* C_d, float* A_d, size_t Num) {
|
||||
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = gputhread; i < Num; i += stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
|
||||
// Delay thread 1 only in the GPU
|
||||
if (gputhread == 1) {
|
||||
unsigned long long int wait_t = 3200000000, start = clock64(), cur;
|
||||
do {
|
||||
cur = clock64() - start;
|
||||
} while (cur < wait_t);
|
||||
}
|
||||
}
|
||||
|
||||
static void HIPRT_CB Stream_Callback(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
// Validate the data and update Data_mismatch
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
Data_mismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
// Increment the Cb_count to indicate that the callback is processed.
|
||||
++Cb_count;
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
float *A_d, *C_d;
|
||||
size_t Nbytes = Num * sizeof(float);
|
||||
|
||||
A_h = (float*)malloc(Nbytes);
|
||||
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (float*)malloc(Nbytes);
|
||||
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
|
||||
// Fill with Phi + i
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
|
||||
|
||||
const unsigned threadsPerBlock = 256;
|
||||
const unsigned blocks = (Num+255)/threadsPerBlock;
|
||||
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
mystream, C_d, A_d, Num);
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
|
||||
|
||||
// Add multiple callbacks to the stream
|
||||
for (int i = 0; i< NUM_CALLS; i++) {
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Stream_Callback, NULL, 0));
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(mystream));
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
|
||||
HIPCHECK(hipFree(A_d));
|
||||
HIPCHECK(hipFree(C_d));
|
||||
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
|
||||
// Each callback would have validated the data and if any mismatch is found,
|
||||
// Data_mismatch will not have proper data. Validate the same.
|
||||
// Cb_count should match the number of callbacks added.
|
||||
if (Data_mismatch.load() != 0) {
|
||||
failed("Mismatch found in the result of the computation!");
|
||||
} else if (Cb_count.load() != NUM_CALLS) {
|
||||
failed("All callbacks for stream did not get called!");
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,165 @@
|
||||
/*
|
||||
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// Testcase Description: This test case is used to check the behaviour of HIP
|
||||
// when multiple hipStreaAddCallback() are called over multiple Threads
|
||||
// This test case is disabled currently.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM all
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <thread>
|
||||
#include <chrono>
|
||||
#include <atomic>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
#define NUM_THREADS 2000
|
||||
|
||||
size_t Num = 4096;
|
||||
std::atomic<size_t>Cb_count{0}, Data_mismatch{0};
|
||||
hipStream_t mystream;
|
||||
float *A_h, *C_h;
|
||||
|
||||
__global__ void vector_square(float* C_d, float* A_d, size_t Num) {
|
||||
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = gputhread; i < Num; i += stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
|
||||
// Delay thread 1 only in the GPU
|
||||
if (gputhread == 1) {
|
||||
unsigned long long int wait_t = 3200000000, start = clock64(), cur;
|
||||
do {
|
||||
cur = clock64() - start;
|
||||
} while (cur < wait_t);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
static void HIPRT_CB Thread1_Callback(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
// Validate the data and update Data_mismatch
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
Data_mismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
// Increment the Cb_count to indicate that the callback is processed.
|
||||
++Cb_count;
|
||||
}
|
||||
|
||||
static void HIPRT_CB Thread2_Callback(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
// Validate the data and update Data_mismatch
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
Data_mismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
// Increment the Cb_count to indicate that the callback is processed.
|
||||
++Cb_count;
|
||||
}
|
||||
|
||||
void Thread1_func() {
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Thread1_Callback, NULL, 0));
|
||||
}
|
||||
|
||||
void Thread2_func() {
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Thread2_Callback, NULL, 0));
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
float *A_d, *C_d;
|
||||
size_t Nbytes = Num * sizeof(float);
|
||||
|
||||
A_h = (float*)malloc(Nbytes);
|
||||
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (float*)malloc(Nbytes);
|
||||
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
|
||||
// Fill with Phi + i
|
||||
for (size_t i = 0; i < Num; i++) {
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
|
||||
|
||||
const unsigned threadsPerBlock = 256;
|
||||
const unsigned blocks = (Num+255)/threadsPerBlock;
|
||||
|
||||
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
mystream, C_d, A_d, Num);
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
|
||||
|
||||
std::thread T[NUM_THREADS];
|
||||
for (int i = 0; i < NUM_THREADS; i++) {
|
||||
// Use different callback for every even thread
|
||||
// The callbacks will be added to same stream from different threads
|
||||
if ((i%2) == 0)
|
||||
T[i] = std::thread(Thread1_func);
|
||||
else
|
||||
T[i] = std::thread(Thread2_func);
|
||||
}
|
||||
|
||||
// Wait until all the threads finish their execution
|
||||
for (int i = 0; i < NUM_THREADS; i++) {
|
||||
T[i].join();
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(mystream));
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
|
||||
HIPCHECK(hipFree(A_d));
|
||||
HIPCHECK(hipFree(C_d));
|
||||
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
|
||||
// Cb_count should match total number of callbacks added from both threads
|
||||
// Data_mismatch will be updated if there is problem in data validation
|
||||
if (Cb_count.load() != NUM_THREADS) {
|
||||
failed("All callbacks for stream did not get called!");
|
||||
} else if (Data_mismatch.load() != 0) {
|
||||
failed("Mismatch found in the result of the computation!");
|
||||
}
|
||||
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,147 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// Testcase Description: This test case checks whether hipStreamSynchronize()
|
||||
// is taking less time than the time taken by Callback() function launched
|
||||
// by hipStreamAddCallback() api.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <stdio.h>
|
||||
#include <unistd.h>
|
||||
#include <chrono>
|
||||
#include <atomic>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
#define SECONDS_TO_WAIT 5
|
||||
#define TO_MICROSECONDS 1000000
|
||||
|
||||
hipStream_t mystream;
|
||||
size_t N_elmts = 4096;
|
||||
bool Init_callback = false;
|
||||
std::atomic<int> Data_mismatch{0};
|
||||
|
||||
__global__ void vector_square(float* C_d, float* A_d, size_t N_elmts) {
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = offset; i < N_elmts; i += stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
|
||||
// Delay the thread 1
|
||||
if (offset == 1) {
|
||||
unsigned long long int wait_t = 3200000000, start = clock64(), cur;
|
||||
do {
|
||||
cur = clock64() - start;
|
||||
} while (cur < wait_t);
|
||||
}
|
||||
}
|
||||
|
||||
float *A_h, *C_h;
|
||||
|
||||
static void HIPRT_CB Callback1(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
// Mark that the callback is entered. This is checked in main thread.
|
||||
Init_callback = true;
|
||||
|
||||
// Validate the data
|
||||
for (size_t i = 0; i < N_elmts; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
Data_mismatch++;
|
||||
}
|
||||
}
|
||||
|
||||
// Delay the callback completion
|
||||
sleep(SECONDS_TO_WAIT);
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
float *A_d, *C_d;
|
||||
size_t Nbytes = N_elmts * sizeof(float);
|
||||
float tElapsed = 1.0f;
|
||||
|
||||
A_h = (float*)malloc(Nbytes);
|
||||
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (float*)malloc(Nbytes);
|
||||
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
|
||||
// Fill with Phi + i
|
||||
for (size_t i = 0; i < N_elmts; i++) {
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
|
||||
|
||||
const unsigned threadsPerBlock = 256;
|
||||
const unsigned blocks = (N_elmts + 255)/threadsPerBlock;
|
||||
|
||||
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0,
|
||||
mystream, C_d, A_d, N_elmts);
|
||||
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Callback1, NULL, 0));
|
||||
|
||||
// Wait untill Callback() function changes the Init_callback value to true
|
||||
while (!Init_callback) {}
|
||||
|
||||
// Since the callback is supposed to be called only after an implicit stream
|
||||
// synchronization, hipStreamSynchronize call shoud not take much time.
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
HIPCHECK(hipStreamSynchronize(mystream));
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
auto duration = std::chrono::duration_cast<std::chrono::microseconds>(stop - start);
|
||||
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
HIPCHECK(hipFree(C_d));
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
|
||||
if (Data_mismatch.load() != 0) {
|
||||
failed("Output from kernel execution is not as expected");
|
||||
}
|
||||
|
||||
// There is a delay of 5000000 microseconds in the Callback() function, the
|
||||
// duration.count() value is expected to less than 5000000 microseconds
|
||||
// because it is expected that stream synchronization completed the moment
|
||||
// Callback function starts the execution and not untill Callback function
|
||||
// completes the execution. Therefore the hipStreamSynchronize() in the
|
||||
// main thread should hardly take any time to complete.
|
||||
|
||||
if (duration.count() < SECONDS_TO_WAIT * TO_MICROSECONDS) {
|
||||
passed();
|
||||
} else {
|
||||
failed("hipStreamSynchronize is waiting untill Callback() completes.");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,63 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
* */
|
||||
|
||||
// Testcase Description: This test case tests if Host thread continues with
|
||||
// next command after hipStreamAddCallback() api or wait for callback() call to
|
||||
// finish. Ideally Host thread should not wait for callback to finish.
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <unistd.h>
|
||||
#include <stdio.h>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
bool Callback_Completed = false;
|
||||
|
||||
void HIPRT_CB Callback1(hipStream_t stream, hipError_t status, void* userData) {
|
||||
sleep(5);
|
||||
Callback_Completed = true;
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
hipStream_t mystream;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Callback1, NULL, 0));
|
||||
sleep(1);
|
||||
|
||||
// Callback_Completed is initialized to false. The same is set to true at
|
||||
// the end of callback and callback sleeps for 5 seconds.
|
||||
// So, in case Callback_Completed is true here, it means the main thread
|
||||
// has waited till callback is complete and is a fail case.
|
||||
if (Callback_Completed == false) {
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
passed();
|
||||
} else {
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
failed("Unexpected: Host thread is waiting for callback to finish");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,81 @@
|
||||
/*
|
||||
* Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
* Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
* of this software and associated documentation files (the "Software"), to deal
|
||||
* in the Software without restriction, including without limitation the rights
|
||||
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
* copies of the Software, and to permit persons to whom the Software is
|
||||
* furnished to do so, subject to the following conditions:
|
||||
* The above copyright notice and this permission notice shall be included in
|
||||
* all copies or substantial portions of the Software.
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
* IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
* THE SOFTWARE.
|
||||
* */
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
// Checks the callback execution in the same order it was added
|
||||
// Also, it checks if the number of callbacks executed are same as the number
|
||||
// of callbacks added
|
||||
|
||||
#include <stdio.h>
|
||||
#include <atomic>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
#define HIPRT_CB
|
||||
#endif
|
||||
|
||||
#define NUM_CALLS 10
|
||||
hipStream_t mystream;
|
||||
bool Callback_SequenceMismatch = false;
|
||||
std::atomic<int> Cb_ordinal{0};
|
||||
|
||||
void HIPRT_CB Stream_Callback(hipStream_t stream, hipError_t status,
|
||||
void* userData) {
|
||||
// Userdata has the order of the callback. It should match with
|
||||
// the callback counter Cb_ordinal as the sequence of callback
|
||||
// should match the sequence of callback addition
|
||||
if (*(reinterpret_cast<int*>(userData)) == Cb_ordinal) {
|
||||
// Increment the Cb_ordinal to prepare for next sequence
|
||||
Cb_ordinal++;
|
||||
} else {
|
||||
Callback_SequenceMismatch = true;
|
||||
}
|
||||
|
||||
delete reinterpret_cast<int*>(userData);
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
int *ptr;
|
||||
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
|
||||
for (int i = 0; i< NUM_CALLS; i++) {
|
||||
ptr = new int;
|
||||
*ptr = i;
|
||||
// Pass the userdata with the order of the callback addition
|
||||
HIPCHECK(hipStreamAddCallback(mystream, Stream_Callback,
|
||||
reinterpret_cast<void*>(ptr), 0));
|
||||
}
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(mystream));
|
||||
HIPCHECK(hipStreamDestroy(mystream));
|
||||
|
||||
if (!(Cb_ordinal == (NUM_CALLS))) {
|
||||
failed("All callbacks for stream did not get called!");
|
||||
}
|
||||
|
||||
if (Callback_SequenceMismatch == false) {
|
||||
passed();
|
||||
} else {
|
||||
failed("hipStreamAddCallback() calls did not execute in sequence!");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,92 @@
|
||||
/*
|
||||
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
int numDevices;
|
||||
hipGetDeviceCount(&numDevices);
|
||||
for (int i = 0; i < numDevices; i++) {
|
||||
hipStream_t stream;
|
||||
int priority;
|
||||
int priority_normal;
|
||||
int priority_low;
|
||||
int priority_high;
|
||||
int priority_check;
|
||||
|
||||
// Test is to get the Stream Priority Range
|
||||
HIPCHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
printf("Priority range is %d for low and %d for high \n", priority_low, priority_high);
|
||||
priority_normal = priority_low + priority_high;
|
||||
|
||||
// Check if priorities are indeed supported
|
||||
if ((priority_low + priority_high) != 0) {
|
||||
failed("Priorities are not supported");
|
||||
}
|
||||
|
||||
// Checking Priority of default stream
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIPCHECK(hipStreamGetPriority(stream, &priority));
|
||||
if (priority_normal != priority) {
|
||||
failed("Unable to set Normal Priority for the stream");
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
// Creating Stream with Priorities
|
||||
HIPCHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
|
||||
HIPCHECK(hipStreamGetPriority(stream, &priority_check));
|
||||
if (priority_check != priority_high) {
|
||||
failed("Unable to set high priority for the stream");
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
HIPCHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
|
||||
HIPCHECK(hipStreamGetPriority(stream, &priority_check));
|
||||
if (priority_check != priority_low) {
|
||||
failed("Unable to set low priority for the stream");
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
// creating a stream with boundry cases
|
||||
HIPCHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low+1));
|
||||
HIPCHECK(hipStreamGetPriority(stream, &priority_check));
|
||||
if (priority_check != priority_low) {
|
||||
failed("setting priority failed ");
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
HIPCHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high-1));
|
||||
HIPCHECK(hipStreamGetPriority(stream, &priority_check));
|
||||
if (priority_check != priority_high) {
|
||||
failed("setting priority failed ");
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
passed();
|
||||
return 0;
|
||||
}
|
||||
@@ -57,7 +57,7 @@ void texture2Dtest()
|
||||
|
||||
// Use the texture object
|
||||
hipResourceDesc texRes;
|
||||
hipMemset(&texRes, 0, sizeof(texRes));
|
||||
memset(&texRes, 0, sizeof(texRes));
|
||||
texRes.resType = hipResourceTypePitch2D;
|
||||
texRes.res.pitch2D.devPtr = devPtrA;
|
||||
texRes.res.pitch2D.height = SIZE_H;
|
||||
@@ -66,7 +66,7 @@ void texture2Dtest()
|
||||
texRes.res.pitch2D.desc = hipCreateChannelDesc<TYPE_t>();
|
||||
|
||||
hipTextureDesc texDescr;
|
||||
hipMemset(&texDescr, 0, sizeof(texDescr));
|
||||
memset(&texDescr, 0, sizeof(texDescr));
|
||||
texDescr.normalizedCoords = false;
|
||||
texDescr.filterMode = hipFilterModePoint;
|
||||
texDescr.mipmapFilterMode = hipFilterModePoint;
|
||||
|
||||
Reference in New Issue
Block a user