[dtest] hipModuleLoad and related API tests
APIs: hipModuleLoad, hipModuleLoadData and hipModuleLoadDataEx,
hipModuleGetGlobal(), hipModuleGetFunction
Functional and negative tests
Repeated call to ModuleLoadXX/ModuleUnloadXX
Few tests are disabled now, will be enabled when functional
SWDEV-238517 for enhancing hip unit tests
Change-Id: I65c12027e32db80213468fdee1c5cc1aa3e60bfd
[ROCm/clr commit: 4f48154f46]
Bu işleme şunda yer alıyor:
Rupam Chetia
işlemeyi yapan:
Mohan Kumar Mithur
Mohan Kumar Mithur
ebeveyn
427cb5bf4c
işleme
07c5bbbf6c
+312
@@ -0,0 +1,312 @@
|
||||
/*
|
||||
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_CMD: kernel_composite_test.code %hc --genco %S/kernel_composite_test.cpp -o kernel_composite_test.code
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t --tests 0x1
|
||||
* TEST: %t --tests 0x2
|
||||
* TEST: %t --tests 0x3
|
||||
* HIT_END
|
||||
*/
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#ifdef __linux__
|
||||
#include <unistd.h>
|
||||
#include <sys/wait.h>
|
||||
#endif
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <cstddef>
|
||||
#include <vector>
|
||||
#include "test_common.h"
|
||||
|
||||
#define TEST_ITERATIONS 1000
|
||||
#define CODEOBJ_FILE "kernel_composite_test.code"
|
||||
#define CODEOBJ_GLOB_KERNEL1 "testWeightedCopy"
|
||||
#define CODEOBJ_GLOB_KERNEL2 "getAvg"
|
||||
#define BLOCKSPERCULDULD 6
|
||||
#define THREADSPERBLOCKLDULD 256
|
||||
|
||||
unsigned int globTestID = 0;
|
||||
|
||||
/**
|
||||
* Fetches Gpu device count
|
||||
*/
|
||||
void getDeviceCount(int *pdevCnt) {
|
||||
#ifdef __linux__
|
||||
int fd[2], val = 0;
|
||||
pid_t childpid;
|
||||
|
||||
// create pipe descriptors
|
||||
pipe(fd);
|
||||
|
||||
// disable visible_devices env from shell
|
||||
unsetenv("ROCR_VISIBLE_DEVICES");
|
||||
unsetenv("HIP_VISIBLE_DEVICES");
|
||||
|
||||
childpid = fork();
|
||||
|
||||
if (childpid > 0) { // Parent
|
||||
close(fd[1]);
|
||||
// parent will wait to read the device cnt
|
||||
read(fd[0], &val, sizeof(val));
|
||||
|
||||
// close the read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
// wait for child exit
|
||||
wait(NULL);
|
||||
|
||||
*pdevCnt = val;
|
||||
} else if (!childpid) { // Child
|
||||
int devCnt = 1;
|
||||
// writing only, no need for read-descriptor
|
||||
close(fd[0]);
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&devCnt));
|
||||
// send the value on the write-descriptor:
|
||||
write(fd[1], &devCnt, sizeof(devCnt));
|
||||
|
||||
// close the write descriptor:
|
||||
close(fd[1]);
|
||||
exit(0);
|
||||
} else { // failure
|
||||
*pdevCnt = 1;
|
||||
return;
|
||||
}
|
||||
|
||||
#else
|
||||
HIPCHECK(hipGetDeviceCount(pdevCnt));
|
||||
#endif
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates hipModuleLoadUnload if globTestID = 1
|
||||
* Validates hipModuleLoadDataUnload if globTestID = 2
|
||||
* Validates hipModuleLoadDataExUnload if globTestID = 3
|
||||
*/
|
||||
bool testhipModuleLoadUnloadFunc(const std::vector<char>& buffer) {
|
||||
size_t N = 16*16;
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int *A_d, *B_d;
|
||||
int *A_h, *B_h;
|
||||
unsigned blocks = HipTest::setNumBlocks(BLOCKSPERCULDULD,
|
||||
THREADSPERBLOCKLDULD, N);
|
||||
int deviceid;
|
||||
hipGetDevice(&deviceid);
|
||||
printf("pid = %u deviceid = %d\n", getpid(), deviceid);
|
||||
// allocate host and device buffer
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d, Nbytes));
|
||||
|
||||
A_h = reinterpret_cast<int *>(malloc(Nbytes));
|
||||
if (NULL == A_h) {
|
||||
failed("Failed to allocate using malloc");
|
||||
}
|
||||
B_h = reinterpret_cast<int *>(malloc(Nbytes));
|
||||
if (NULL == B_h) {
|
||||
failed("Failed to allocate using malloc");
|
||||
}
|
||||
// set host buffers
|
||||
for (int idx = 0; idx < N; idx++) {
|
||||
A_h[idx] = deviceid;
|
||||
}
|
||||
// Copy buffer from host to device
|
||||
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
if (1 == globTestID) {
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
} else if (2 == globTestID) {
|
||||
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
|
||||
} else if (3 == globTestID) {
|
||||
HIPCHECK(hipModuleLoadDataEx(&Module,
|
||||
&buffer[0], 0, nullptr, nullptr));
|
||||
}
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module,
|
||||
CODEOBJ_GLOB_KERNEL1));
|
||||
float deviceGlobalFloatH = 3.14;
|
||||
int deviceGlobalInt1H = 100*deviceid;
|
||||
int deviceGlobalInt2H = 50*deviceid;
|
||||
short deviceGlobalShortH = 25*deviceid;
|
||||
char deviceGlobalCharH = 13*deviceid;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize,
|
||||
Module, "deviceGlobalFloat"));
|
||||
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
|
||||
&deviceGlobalFloatH,
|
||||
deviceGlobalSize));
|
||||
HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize,
|
||||
Module, "deviceGlobalInt1"));
|
||||
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
|
||||
&deviceGlobalInt1H,
|
||||
deviceGlobalSize));
|
||||
HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize,
|
||||
Module,
|
||||
"deviceGlobalInt2"));
|
||||
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
|
||||
&deviceGlobalInt2H, deviceGlobalSize));
|
||||
HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize,
|
||||
Module, "deviceGlobalShort"));
|
||||
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
|
||||
&deviceGlobalShortH, deviceGlobalSize));
|
||||
HIPCHECK(hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize, Module, "deviceGlobalChar"));
|
||||
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal),
|
||||
&deviceGlobalCharH, deviceGlobalSize));
|
||||
// Launch Function kernel function
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
args._Ad = reinterpret_cast<void*>(A_d);
|
||||
args._Bd = reinterpret_cast<void*>(B_d);
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1,
|
||||
N, 1, 1, 0, stream, NULL,
|
||||
reinterpret_cast<void**>(&config)));
|
||||
// Copy buffer from decice to host
|
||||
HIPCHECK(hipMemcpyAsync(B_h, B_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
// Check the results
|
||||
for (int idx = 0; idx < N; idx++) {
|
||||
if (B_h[idx] != (deviceGlobalInt1H*A_h[idx]
|
||||
+ deviceGlobalInt2H
|
||||
+ static_cast<int>(deviceGlobalShortH) +
|
||||
+ static_cast<int>(deviceGlobalCharH)
|
||||
+ static_cast<int>(deviceGlobalFloatH*deviceGlobalFloatH))) {
|
||||
printf("Matrix Addition Failed\n");
|
||||
// exit the current process with failure
|
||||
return false;
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
// free memory
|
||||
HIPCHECK(hipFree(B_d));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
free(B_h);
|
||||
free(A_h);
|
||||
printf("pid:%u PASSED\n", getpid());
|
||||
return true;
|
||||
}
|
||||
|
||||
/**
|
||||
* Spawn 1 Process for each device
|
||||
*
|
||||
*/
|
||||
void spawnProc(int deviceCount, const std::vector<char>& buffer) {
|
||||
int numDevices = deviceCount;
|
||||
bool TestPassed = true;
|
||||
#ifdef __linux__
|
||||
pid_t pid = 0;
|
||||
// spawn a process for each device
|
||||
for (int deviceNo = 0; deviceNo < numDevices; deviceNo++) {
|
||||
if ((pid = fork()) < 0) {
|
||||
printf("Child_Concurrency_MultiGpu : fork() returned error %d\n",
|
||||
pid);
|
||||
failed("Test Failed!");
|
||||
} else if (!pid) { // Child process
|
||||
bool TestPassedChild = true;
|
||||
// set the device id for the current process
|
||||
HIPCHECK(hipSetDevice(deviceNo));
|
||||
TestPassedChild = testhipModuleLoadUnloadFunc(buffer);
|
||||
|
||||
if (TestPassedChild) {
|
||||
exit(0); // child exit with success status
|
||||
} else {
|
||||
printf("Child_Concurrency_MultiGpu : childpid %d failed\n",
|
||||
getpid());
|
||||
exit(1); // child exit with failure status
|
||||
}
|
||||
}
|
||||
}
|
||||
int cumStatus = 0;
|
||||
// Parent shall wait for child to complete
|
||||
for (int i = 0; i < numDevices; i++) {
|
||||
int pidwait = 0, exitStatus;
|
||||
pidwait = wait(&exitStatus);
|
||||
cumStatus |= WEXITSTATUS(exitStatus);
|
||||
}
|
||||
if (cumStatus) {
|
||||
TestPassed &= false;
|
||||
}
|
||||
#else
|
||||
for (int deviceNo = 0; deviceNo < numDevices; deviceNo++) {
|
||||
// set the device id for the current process
|
||||
HIPCHECK(hipSetDevice(deviceNo));
|
||||
TestPassed &= testhipModuleLoadUnloadFunc(buffer);
|
||||
}
|
||||
#endif
|
||||
if (TestPassed) {
|
||||
passed();
|
||||
} else {
|
||||
failed("hipMallocChild_Concurrency_MultiGpu Failed!");
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
int numDevices = 0;
|
||||
getDeviceCount(&numDevices);
|
||||
if (1 == numDevices) {
|
||||
printf("Testing on Single GPU machine.\n");
|
||||
}
|
||||
std::ifstream file(CODEOBJ_FILE,
|
||||
std::ios::binary | std::ios::ate);
|
||||
std::streamsize fsize = file.tellg();
|
||||
file.seekg(0, std::ios::beg);
|
||||
|
||||
std::vector<char> buffer(fsize);
|
||||
if (!file.read(buffer.data(), fsize)) {
|
||||
failed("could not open code object '%s'\n", CODEOBJ_FILE);
|
||||
}
|
||||
file.close();
|
||||
if (p_tests == 0x1) {
|
||||
globTestID = 1;
|
||||
spawnProc(numDevices, buffer);
|
||||
} else if (p_tests == 0x2) {
|
||||
globTestID = 2;
|
||||
spawnProc(numDevices, buffer);
|
||||
} else if (p_tests == 0x3) {
|
||||
globTestID = 3;
|
||||
spawnProc(numDevices, buffer);
|
||||
}
|
||||
}
|
||||
+119
@@ -0,0 +1,119 @@
|
||||
/*
|
||||
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 --tests 0x1
|
||||
* TEST: %t --tests 0x2
|
||||
* TEST: %t --tests 0x3
|
||||
* HIT_END
|
||||
*/
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <cstddef>
|
||||
#include <vector>
|
||||
#include "test_common.h"
|
||||
|
||||
#define TEST_ITERATIONS 1000
|
||||
#define CODEOBJ_FILE "kernel_composite_test.code"
|
||||
/**
|
||||
* Run Valgrind tool with these test cases to validate memory leakage.
|
||||
* E.g. valgrind --leak-check=yes ./a.out --tests 0x1
|
||||
*/
|
||||
|
||||
/**
|
||||
* Internal Function
|
||||
*/
|
||||
std::vector<char> load_file() {
|
||||
std::ifstream file(CODEOBJ_FILE, std::ios::binary | std::ios::ate);
|
||||
std::streamsize fsize = file.tellg();
|
||||
file.seekg(0, std::ios::beg);
|
||||
std::vector<char> buffer(fsize);
|
||||
if (!file.read(buffer.data(), fsize)) {
|
||||
failed("could not open code object '%s'\n", CODEOBJ_FILE);
|
||||
}
|
||||
file.close();
|
||||
return buffer;
|
||||
}
|
||||
/**
|
||||
* Validates no memory leakage for hipModuleLoad
|
||||
*/
|
||||
void testhipModuleLoadUnloadStress() {
|
||||
for (int count = 0; count < TEST_ITERATIONS; count++) {
|
||||
hipModule_t Module;
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
hipFunction_t Function;
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy"));
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
}
|
||||
}
|
||||
/**
|
||||
* Validates no memory leakage for hipModuleLoadData
|
||||
*/
|
||||
void testhipModuleLoadDataUnloadStress() {
|
||||
auto buffer = load_file();
|
||||
for (int count = 0; count < TEST_ITERATIONS; count++) {
|
||||
hipModule_t Module;
|
||||
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
|
||||
hipFunction_t Function;
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy"));
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
}
|
||||
}
|
||||
/**
|
||||
* Validates no memory leakage for hipModuleLoadDataEx
|
||||
*/
|
||||
void testhipModuleLoadDataExUnloadStress() {
|
||||
auto buffer = load_file();
|
||||
for (int count = 0; count < TEST_ITERATIONS; count++) {
|
||||
hipModule_t Module;
|
||||
HIPCHECK(hipModuleLoadDataEx(&Module, &buffer[0], 0,
|
||||
nullptr, nullptr));
|
||||
hipFunction_t Function;
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module, "testWeightedCopy"));
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipInit(0));
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
HIPCHECK(hipDeviceGet(&device, 0));
|
||||
HIPCHECK(hipCtxCreate(&context, 0, device));
|
||||
#endif
|
||||
if (p_tests == 0x1) {
|
||||
testhipModuleLoadUnloadStress();
|
||||
} else if (p_tests == 0x2) {
|
||||
testhipModuleLoadDataUnloadStress();
|
||||
} else if (p_tests == 0x3) {
|
||||
testhipModuleLoadDataExUnloadStress();
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,866 @@
|
||||
/*
|
||||
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 --tests 0x10 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x11 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x12 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x13
|
||||
* TEST: %t --tests 0x14
|
||||
* TEST: %t --tests 0x15
|
||||
* TEST: %t --tests 0x20 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x21 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x22
|
||||
* TEST: %t --tests 0x30 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x31 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x32
|
||||
* TEST: %t --tests 0x40 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x41
|
||||
* TEST: %t --tests 0x42
|
||||
* TEST: %t --tests 0x43 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x44 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t --tests 0x45
|
||||
* TEST: %t --tests 0x50 EXCLUDE_HIP_PLATFORM hcc rocclr nvcc
|
||||
* TEST: %t --tests 0x51 EXCLUDE_HIP_PLATFORM hcc rocclr nvcc
|
||||
* TEST: %t --tests 0x52 EXCLUDE_HIP_PLATFORM hcc rocclr
|
||||
* TEST: %t --tests 0x53
|
||||
* TEST: %t --tests 0x54
|
||||
* TEST: %t --tests 0x55 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t --tests 0x56
|
||||
* TEST: %t --tests 0x60 EXCLUDE_HIP_PLATFORM nvcc
|
||||
* HIT_END
|
||||
*/
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <ctime>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <cstddef>
|
||||
#include <vector>
|
||||
#include "test_common.h"
|
||||
|
||||
#define FILENAME_NONEXST "sample_nonexst.code"
|
||||
#define FILENAME_EMPTY "emptyfile.code"
|
||||
#define FILENAME_RAND "rand_file.code"
|
||||
#define RANDOMFILE_LEN 2048
|
||||
#define CODEOBJ_FILE "vcpy_kernel.code"
|
||||
#define KERNEL_NAME "hello_world"
|
||||
#define KERNEL_NAME_NONEXST "xyz"
|
||||
#define CODEOBJ_GLOBAL "global_kernel.code"
|
||||
#define DEVGLOB_VAR_NONEXIST "xyz"
|
||||
#define DEVGLOB_VAR "myDeviceGlobal"
|
||||
/**
|
||||
* Internal Function
|
||||
*/
|
||||
std::vector<char> load_file(const char* filename) {
|
||||
std::ifstream file(filename, std::ios::binary | std::ios::ate);
|
||||
std::streamsize fsize = file.tellg();
|
||||
file.seekg(0, std::ios::beg);
|
||||
std::vector<char> buffer(fsize);
|
||||
if (!file.read(buffer.data(), fsize)) {
|
||||
failed("could not open code object '%s'\n", filename);
|
||||
}
|
||||
file.close();
|
||||
return buffer;
|
||||
}
|
||||
|
||||
/**
|
||||
* Internal Function
|
||||
*/
|
||||
void createRandomFile(const char* filename) {
|
||||
std::ofstream outfile(filename, std::ios::binary);
|
||||
char buf[RANDOMFILE_LEN];
|
||||
unsigned int seed = 1;
|
||||
for (int i = 0; i < RANDOMFILE_LEN; i++) {
|
||||
buf[i] = rand_r(&seed) % 256;
|
||||
}
|
||||
outfile.write(buf, RANDOMFILE_LEN);
|
||||
outfile.close();
|
||||
}
|
||||
|
||||
/**
|
||||
* Internal Function
|
||||
*/
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
void initHipCtx(hipCtx_t *pcontext) {
|
||||
HIPCHECK(hipInit(0));
|
||||
hipDevice_t device;
|
||||
HIPCHECK(hipDeviceGet(&device, 0));
|
||||
HIPCHECK(hipCtxCreate(pcontext, 0, device));
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoad
|
||||
* module = nullptr
|
||||
*/
|
||||
bool testhipModuleLoadNeg10() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoad(nullptr, CODEOBJ_FILE))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoad
|
||||
* fname = nullptr
|
||||
*/
|
||||
bool testhipModuleLoadNeg11() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoad(&Module, nullptr))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoad
|
||||
* fname = empty file
|
||||
*/
|
||||
bool testhipModuleLoadNeg12() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
// Create an empty
|
||||
std::fstream fs;
|
||||
fs.open(FILENAME_EMPTY, std::ios::out);
|
||||
fs.close();
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoad(&Module, FILENAME_EMPTY))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
remove(FILENAME_EMPTY);
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoad
|
||||
* fname = ramdom file
|
||||
*/
|
||||
bool testhipModuleLoadNeg13() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
// Create a binary file with random numbers
|
||||
createRandomFile(FILENAME_RAND);
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoad(&Module, FILENAME_RAND))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
remove(FILENAME_RAND);
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoad
|
||||
* fname = non existent file
|
||||
*/
|
||||
bool testhipModuleLoadNeg14() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoad(&Module, FILENAME_NONEXST))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoad
|
||||
* fname = empty string ""
|
||||
*/
|
||||
bool testhipModuleLoadNeg15() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoad(&Module, "")) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoadData
|
||||
* module = nullptr
|
||||
*/
|
||||
bool testhipModuleLoadDataNeg20() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
auto buffer = load_file(CODEOBJ_FILE);
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoadData(nullptr, &buffer[0]))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoadData
|
||||
* image = nullptr
|
||||
*/
|
||||
bool testhipModuleLoadDataNeg21() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoadData(&Module, nullptr))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoadData
|
||||
* image = ramdom file
|
||||
*/
|
||||
bool testhipModuleLoadDataNeg22() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
// Create a binary file with random numbers
|
||||
createRandomFile(FILENAME_RAND);
|
||||
// Open the code object file and copy it in a buffer
|
||||
auto buffer = load_file(FILENAME_RAND);
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoadData(&Module, &buffer[0]))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
remove(FILENAME_RAND);
|
||||
return TestPassed;
|
||||
}
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoadDataEx
|
||||
* module = nullptr
|
||||
*/
|
||||
bool testhipModuleLoadDataExNeg30() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
// Open the code object file and copy it in a buffer
|
||||
auto buffer = load_file(CODEOBJ_FILE);
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoadDataEx(nullptr, &buffer[0], 0, nullptr, nullptr))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoadDataEx
|
||||
* image = nullptr
|
||||
*/
|
||||
bool testhipModuleLoadDataExNeg31() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoadDataEx(&Module, nullptr, 0, nullptr, nullptr))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleLoadDataEx
|
||||
* image = ramdom file
|
||||
*/
|
||||
bool testhipModuleLoadDataExNeg32() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
// Create a binary file with random numbers
|
||||
createRandomFile(FILENAME_RAND);
|
||||
// Open the code object file and copy it in a buffer
|
||||
auto buffer = load_file(FILENAME_RAND);
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleLoadDataEx(&Module, &buffer[0], 0, nullptr, nullptr))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
remove(FILENAME_RAND);
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetFunction
|
||||
* Function = nullptr
|
||||
*/
|
||||
bool testhipModuleGetFunctionNeg40() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
if ((ret = hipModuleGetFunction(nullptr, Module, KERNEL_NAME))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetFunction
|
||||
* Module is uninitialized
|
||||
*/
|
||||
bool testhipModuleGetFunctionNeg41() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetFunction
|
||||
* kname = non existing function
|
||||
*/
|
||||
bool testhipModuleGetFunctionNeg42() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
if ((ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME_NONEXST))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetFunction
|
||||
* kname = nullptr
|
||||
*/
|
||||
bool testhipModuleGetFunctionNeg43() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
if ((ret = hipModuleGetFunction(&Function, Module, nullptr))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetFunction
|
||||
* Module = Unloaded Module
|
||||
*/
|
||||
bool testhipModuleGetFunctionNeg44() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
if ((ret = hipModuleGetFunction(&Function, Module, KERNEL_NAME))
|
||||
!= hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetFunction
|
||||
* kname = Empty String ""
|
||||
*/
|
||||
bool testhipModuleGetFunctionNeg45() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
if ((ret = hipModuleGetFunction(&Function,
|
||||
Module, "")) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* dptr = nullptr
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg50() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
size_t deviceGlobalSize;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
|
||||
if ((ret = hipModuleGetGlobal(nullptr,
|
||||
&deviceGlobalSize, Module, DEVGLOB_VAR)) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* bytes = nullptr
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg51() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
|
||||
if ((ret = hipModuleGetGlobal(&deviceGlobal, nullptr,
|
||||
Module, DEVGLOB_VAR)) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* name = nullptr
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg52() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
|
||||
if ((ret = hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize, Module, nullptr)) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* name = wrong name
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg53() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
|
||||
if ((ret = hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize,
|
||||
Module, DEVGLOB_VAR_NONEXIST)) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* name = Empty String ""
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg54() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
|
||||
if ((ret = hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize, Module, "")) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* Module = Unloaded Module
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg55() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_GLOBAL));
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
if ((ret = hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize, Module, DEVGLOB_VAR)) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleGetGlobal
|
||||
* Module = Uninitialized Module
|
||||
*/
|
||||
bool testhipModuleGetGlobalNeg56() {
|
||||
bool TestPassed = false;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
hipDeviceptr_t deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
if ((ret = hipModuleGetGlobal(&deviceGlobal,
|
||||
&deviceGlobalSize, Module, DEVGLOB_VAR)) != hipSuccess) {
|
||||
TestPassed = true;
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
/**
|
||||
* Validates negative scenarios for hipModuleUnload
|
||||
* 1. Unload an uninitialized module
|
||||
* 2. Unload an unloaded module
|
||||
*/
|
||||
bool testhipModuleLoadNeg60() {
|
||||
bool TestPassed = true;
|
||||
hipError_t ret;
|
||||
hipModule_t Module;
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
hipCtx_t context;
|
||||
initHipCtx(&context);
|
||||
#endif
|
||||
// test case 1
|
||||
if ((ret = hipModuleUnload(Module)) != hipSuccess) {
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
} else {
|
||||
TestPassed &= false;
|
||||
}
|
||||
// test case 2
|
||||
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
if ((ret = hipModuleUnload(Module)) != hipSuccess) {
|
||||
printf("Test Passed: Error Code Returned: '%s'(%d)\n",
|
||||
hipGetErrorString(ret), ret);
|
||||
} else {
|
||||
TestPassed &= false;
|
||||
}
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
#endif
|
||||
return TestPassed;
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
bool TestPassed = true;
|
||||
if (p_tests == 0x10) {
|
||||
TestPassed = testhipModuleLoadNeg10();
|
||||
} else if (p_tests == 0x11) {
|
||||
TestPassed = testhipModuleLoadNeg11();
|
||||
} else if (p_tests == 0x12) {
|
||||
TestPassed = testhipModuleLoadNeg12();
|
||||
} else if (p_tests == 0x13) {
|
||||
TestPassed = testhipModuleLoadNeg13();
|
||||
} else if (p_tests == 0x14) {
|
||||
TestPassed = testhipModuleLoadNeg14();
|
||||
} else if (p_tests == 0x15) {
|
||||
TestPassed = testhipModuleLoadNeg15();
|
||||
} else if (p_tests == 0x20) {
|
||||
TestPassed = testhipModuleLoadDataNeg20();
|
||||
} else if (p_tests == 0x21) {
|
||||
TestPassed = testhipModuleLoadDataNeg21();
|
||||
} else if (p_tests == 0x22) {
|
||||
TestPassed = testhipModuleLoadDataNeg22();
|
||||
} else if (p_tests == 0x30) {
|
||||
TestPassed = testhipModuleLoadDataExNeg30();
|
||||
} else if (p_tests == 0x31) {
|
||||
TestPassed = testhipModuleLoadDataExNeg31();
|
||||
} else if (p_tests == 0x32) {
|
||||
TestPassed = testhipModuleLoadDataExNeg32();
|
||||
} else if (p_tests == 0x40) {
|
||||
TestPassed = testhipModuleGetFunctionNeg40();
|
||||
} else if (p_tests == 0x41) {
|
||||
TestPassed = testhipModuleGetFunctionNeg41();
|
||||
} else if (p_tests == 0x42) {
|
||||
TestPassed = testhipModuleGetFunctionNeg42();
|
||||
} else if (p_tests == 0x43) {
|
||||
TestPassed = testhipModuleGetFunctionNeg43();
|
||||
} else if (p_tests == 0x44) {
|
||||
TestPassed = testhipModuleGetFunctionNeg44();
|
||||
} else if (p_tests == 0x45) {
|
||||
TestPassed = testhipModuleGetFunctionNeg45();
|
||||
} else if (p_tests == 0x50) {
|
||||
TestPassed = testhipModuleGetGlobalNeg50();
|
||||
} else if (p_tests == 0x51) {
|
||||
TestPassed = testhipModuleGetGlobalNeg51();
|
||||
} else if (p_tests == 0x52) {
|
||||
TestPassed = testhipModuleGetGlobalNeg52();
|
||||
} else if (p_tests == 0x53) {
|
||||
TestPassed = testhipModuleGetGlobalNeg53();
|
||||
} else if (p_tests == 0x54) {
|
||||
TestPassed = testhipModuleGetGlobalNeg54();
|
||||
} else if (p_tests == 0x55) {
|
||||
TestPassed = testhipModuleGetGlobalNeg55();
|
||||
} else if (p_tests == 0x56) {
|
||||
TestPassed = testhipModuleGetGlobalNeg56();
|
||||
} else if (p_tests == 0x60) {
|
||||
TestPassed = testhipModuleLoadNeg60();
|
||||
} else {
|
||||
printf("Invalid Test Case \n");
|
||||
exit(1);
|
||||
}
|
||||
if (TestPassed) {
|
||||
passed();
|
||||
} else {
|
||||
failed("Test Case %x Failed!", p_tests);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,41 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#define GLOBAL_BUF_SIZE 2048
|
||||
|
||||
__device__ float deviceGlobalFloat;
|
||||
__device__ int deviceGlobalInt1;
|
||||
__device__ int deviceGlobalInt2;
|
||||
__device__ short deviceGlobalShort;
|
||||
__device__ char deviceGlobalChar;
|
||||
|
||||
__device__ int getSquareOfGlobalFloat() {
|
||||
return static_cast<int>(deviceGlobalFloat*deviceGlobalFloat);
|
||||
}
|
||||
|
||||
extern "C" __global__ void testWeightedCopy(int* a, int* b) {
|
||||
int tx = hipThreadIdx_x;
|
||||
b[tx] = deviceGlobalInt1*a[tx] + deviceGlobalInt2 +
|
||||
static_cast<int>(deviceGlobalShort) + static_cast<int>(deviceGlobalChar)
|
||||
+ getSquareOfGlobalFloat();
|
||||
}
|
||||
Yeni konuda referans
Bir kullanıcı engelle