SWDEV-238517 - [dtest] Test for Multiple Target Architecture Code Object
Added a test to generate a code object for multiple target architectures (including for the current device), load and execute the kernel. SWDEV-238517 for enhancing hip unit tests Change-Id: I509d01124abdc0495cfc770ab5508738f108c91c
Этот коммит содержится в:
Rupam Chetia
коммит произвёл
Sumanth Tumbalam Gooty
Sumanth Tumbalam Gooty
родитель
72e7826adf
Коммит
7c20d3e4e4
@@ -20,84 +20,187 @@ THE SOFTWARE.
|
||||
/* HIT_START
|
||||
* BUILD_CMD: vcpy_kernel.code %hc --genco %S/vcpy_kernel.cpp -o vcpy_kernel.code
|
||||
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
||||
* TEST: %t
|
||||
* TEST: %t --tests 0x1
|
||||
* TEST: %t --tests 0x2 EXCLUDE_HIP_PLATFORM nvidia
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <thread>
|
||||
#include <chrono>
|
||||
|
||||
#ifdef __linux__
|
||||
#include <unistd.h>
|
||||
#endif
|
||||
#include "test_common.h"
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN << 2
|
||||
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define SIZE (LEN << 2)
|
||||
#define COMMAND_LEN 256
|
||||
#define CODE_OBJ_SINGLEARCH "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
#define CODE_OBJ_MULTIARCH "vcpy_kernel_multarch.code"
|
||||
|
||||
#define HIP_CHECK(status) \
|
||||
if (status != hipSuccess) { \
|
||||
std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \
|
||||
exit(0); \
|
||||
bool testCodeObjFile(const char *codeObjFile) {
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
A[i] = i * 1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
|
||||
HIPCHECK(hipMemcpyHtoD(Ad, A, SIZE));
|
||||
HIPCHECK(hipMemcpyHtoD(Bd, B, SIZE));
|
||||
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIPCHECK(hipModuleLoad(&Module, codeObjFile));
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
args._Ad = reinterpret_cast<void*>(Ad);
|
||||
args._Bd = reinterpret_cast<void*>(Bd);
|
||||
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, LEN, 1, 1, 0,
|
||||
stream, NULL,
|
||||
reinterpret_cast<void**>(&config)));
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
HIPCHECK(hipMemcpyDtoH(B, Bd, SIZE));
|
||||
|
||||
bool btestPassed = true;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (A[i] != B[i]) {
|
||||
btestPassed = false;
|
||||
break;
|
||||
}
|
||||
|
||||
int main() {
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
A[i] = i * 1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
HIPCHECK(hipInit(0));
|
||||
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
HIPCHECK(hipDeviceGet(&device, 0));
|
||||
HIPCHECK(hipCtxCreate(&context, 0, device));
|
||||
|
||||
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
|
||||
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
|
||||
HIPCHECK(hipMemcpyHtoD(Ad, A, SIZE));
|
||||
HIPCHECK(hipMemcpyHtoD(Bd, B, SIZE));
|
||||
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIPCHECK(hipModuleLoad(&Module, fileName));
|
||||
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
args._Ad = (void*) Ad;
|
||||
args._Bd = (void*) Bd;
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
|
||||
HIPCHECK(hipMemcpyDtoH(B, Bd, SIZE));
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
assert(A[i] == B[i]);
|
||||
}
|
||||
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
HIPCHECK(hipCtxDestroy(context));
|
||||
passed();
|
||||
}
|
||||
HIPCHECK(hipFree(reinterpret_cast<void*>(Bd)));
|
||||
HIPCHECK(hipFree(reinterpret_cast<void*>(Ad)));
|
||||
delete[] B;
|
||||
delete[] A;
|
||||
HIPCHECK(hipModuleUnload(Module));
|
||||
return btestPassed;
|
||||
}
|
||||
|
||||
#ifdef __linux__
|
||||
/**
|
||||
* Check if environment variable $ROCM_PATH is defined
|
||||
*
|
||||
*/
|
||||
bool isRocmPathSet() {
|
||||
FILE *fpipe;
|
||||
char const *command = "echo $ROCM_PATH";
|
||||
fpipe = popen(command, "r");
|
||||
|
||||
if (fpipe == nullptr) {
|
||||
printf("Unable to create command\n");
|
||||
return false;
|
||||
}
|
||||
char command_op[COMMAND_LEN];
|
||||
if (fgets(command_op, COMMAND_LEN, fpipe)) {
|
||||
size_t len = strlen(command_op);
|
||||
if (len > 1) { // This is because fgets always adds newline character
|
||||
pclose(fpipe);
|
||||
return true;
|
||||
}
|
||||
}
|
||||
pclose(fpipe);
|
||||
return false;
|
||||
}
|
||||
#endif
|
||||
|
||||
bool testMultiTargArchCodeObj() {
|
||||
bool btestPassed = true;
|
||||
#ifdef __linux__
|
||||
char command[COMMAND_LEN];
|
||||
hipDeviceProp_t props;
|
||||
hipGetDeviceProperties(&props, 0);
|
||||
// Hardcoding the codeobject lines in multiple string to avoid cpplint warning
|
||||
std::string CodeObjL1 = "#include \"hip/hip_runtime.h\"\n";
|
||||
std::string CodeObjL2 =
|
||||
"extern \"C\" __global__ void hello_world(float* a, float* b) {\n";
|
||||
std::string CodeObjL3 = " int tx = hipThreadIdx_x;\n";
|
||||
std::string CodeObjL4 = " b[tx] = a[tx];\n";
|
||||
std::string CodeObjL5 = "}";
|
||||
// Creating the full code object string
|
||||
static std::string CodeObj = CodeObjL1 + CodeObjL2 + CodeObjL3 +
|
||||
CodeObjL4 + CodeObjL5;
|
||||
std::ofstream ofs("/tmp/vcpy_kernel.cpp", std::ofstream::out);
|
||||
ofs << CodeObj;
|
||||
ofs.close();
|
||||
// Copy the file into current working location if not available
|
||||
if (access("/tmp/vcpy_kernel.cpp", F_OK) == -1) {
|
||||
printf("Code Object File: /tmp/vcpy_kernel.cpp not found \n");
|
||||
return true;
|
||||
}
|
||||
// Generate the command to generate multi architecture code object file
|
||||
const char* hipcc_path = nullptr;
|
||||
if (isRocmPathSet()) {
|
||||
hipcc_path = "$ROCM_PATH/bin/hipcc";
|
||||
} else {
|
||||
hipcc_path = "/opt/rocm/bin/hipcc";
|
||||
}
|
||||
/* Putting these command parameters into a variable to shorten the string
|
||||
literal length in order to avoid multiline string literal cpplint warning
|
||||
*/
|
||||
const char* genco_option = "--offload-arch";
|
||||
const char* input_codeobj = "/tmp/vcpy_kernel.cpp";
|
||||
snprintf(command, COMMAND_LEN,
|
||||
"%s --genco %s=gfx801,gfx802,gfx803,gfx900,gfx908,%s %s -o %s",
|
||||
hipcc_path, genco_option, props.gcnArchName, input_codeobj,
|
||||
CODE_OBJ_MULTIARCH);
|
||||
|
||||
printf("command = %s\n", command);
|
||||
system((const char*)command);
|
||||
// Check if the code object file is created
|
||||
snprintf(command, COMMAND_LEN, "./%s",
|
||||
CODE_OBJ_MULTIARCH);
|
||||
|
||||
if (access(command, F_OK) == -1) {
|
||||
printf("Code Object File not found \n");
|
||||
return true;
|
||||
}
|
||||
btestPassed = testCodeObjFile(CODE_OBJ_MULTIARCH);
|
||||
#else
|
||||
printf("This test is skipped due to non linux environment.\n");
|
||||
#endif
|
||||
return btestPassed;
|
||||
}
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
bool TestPassed = true;
|
||||
if (p_tests == 0x1) {
|
||||
/* In this test scenario a code object file for the current
|
||||
GPU architecture is generated, loaded and executed. */
|
||||
TestPassed = testCodeObjFile(CODE_OBJ_SINGLEARCH);
|
||||
} else if (p_tests == 0x2) {
|
||||
/* In this test scenario a code object file for the multiple
|
||||
GPU architectures (including the current) is generated, loaded
|
||||
and executed. */
|
||||
TestPassed = testMultiTargArchCodeObj();
|
||||
} else {
|
||||
printf("Invalid Test Case \n");
|
||||
exit(1);
|
||||
}
|
||||
if (TestPassed) {
|
||||
passed();
|
||||
} else {
|
||||
failed("Test Case %x Failed!", p_tests);
|
||||
}
|
||||
}
|
||||
|
||||
Ссылка в новой задаче
Block a user