[dtest] Enhanced tests to cover multiple hipmemset and hipmemsetAsync apis
SWDEV-238517 for enhancing hip unit tests
Change-Id: Iba2c419a4487955f34b4f19abe174ef427d289d8
[ROCm/clr commit: 7d77902198]
This commit is contained in:
Sudheer Kumar
committed by
Mohan Kumar Mithur
Mohan Kumar Mithur
parent
14e9205334
commit
31e62c5852
@@ -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");
|
||||
}
|
||||
}
|
||||
|
||||
+173
@@ -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();
|
||||
}
|
||||
Reference in New Issue
Block a user