From 7c20d3e4e4ffadc2bfcd83d49afec7ac11b94989 Mon Sep 17 00:00:00 2001 From: Rupam Chetia Date: Thu, 26 Nov 2020 17:54:14 +0530 Subject: [PATCH] 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 --- tests/src/runtimeApi/module/hipModule.cpp | 241 +++++++++++++++------- 1 file changed, 172 insertions(+), 69 deletions(-) diff --git a/tests/src/runtimeApi/module/hipModule.cpp b/tests/src/runtimeApi/module/hipModule.cpp index 9ed5a72415..8e1044804c 100755 --- a/tests/src/runtimeApi/module/hipModule.cpp +++ b/tests/src/runtimeApi/module/hipModule.cpp @@ -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 #include -#include -#include -#include - +#ifdef __linux__ +#include +#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(&Ad), SIZE)); + HIPCHECK(hipMalloc(reinterpret_cast(&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(Ad); + args._Bd = reinterpret_cast(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(&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(Bd))); + HIPCHECK(hipFree(reinterpret_cast(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); + } }