From 9a005eacf3d30a50db23a47778807d49965cf37f Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Sat, 15 Feb 2020 18:41:15 -0500 Subject: [PATCH] Support defaultDriver test with Clang+Hcc Runtime Fix the following issues: 1.Ignore hidden arguments of kernel functions. 2.Look up both origial function name and function name with .kd postfix when argments are retrived from module. 3.Addition, fix compiling issue of LaunchKernel test app. Change-Id: I9400943f2f02433cb4409b19c0cac3626c2bc454 --- src/hip_module.cpp | 16 +++++++++++++--- src/program_state.inl | 21 +++++++++++++++++++++ tests/src/gcc/LaunchKernel.c | 12 ++++++------ 3 files changed, 40 insertions(+), 9 deletions(-) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 16606e8016..b8952f8a68 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -1030,19 +1030,29 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t if (!*func) return hipErrorInvalidValue; std::string name_str(name); + std::string namekd_str(name_str + ".kd"); + auto kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent); if (kernel.handle == 0u) { - name_str.append(".kd"); - kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent); + kernel = find_kernel_by_name(hmod->executable, namekd_str.c_str(), agent); } if (kernel.handle == 0u) return hipErrorNotFound; + auto it = hmod->kernargs.find(name_str); //Look up args from the original name + if (it == hmod->kernargs.end()) { + it = hmod->kernargs.find(namekd_str); //Look up args from .kd name + if (it == hmod->kernargs.end()) { + return hipErrorNotFound; + } + } + // TODO: refactor the whole ihipThisThat, which is a mess and yields the // below, due to hipFunction_t being a pointer to ihipModuleSymbol_t. + func[0][0] = *static_cast( - Kernel_descriptor{kernel_object(kernel), name_str, hmod->kernargs[name_str]}); + Kernel_descriptor{kernel_object(kernel), name_str, it->second}); return hipSuccess; } diff --git a/src/program_state.inl b/src/program_state.inl index 272addd053..2e959a9928 100644 --- a/src/program_state.inl +++ b/src/program_state.inl @@ -714,6 +714,27 @@ public: != AMD_COMGR_STATUS_SUCCESS) return; + //Look up “.value_kind” to decide whether to ignore it + //See http://llvm.org/docs/AMDGPUUsage.html#code-object-v3-metadata-mattr-code-object-v3 + amd_comgr_metadata_node_t arg_value_kind_md; + if (amd_comgr_metadata_lookup(arg_md, ".value_kind", &arg_value_kind_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + + std::string arg_value_kind{ metadata_to_string(arg_value_kind_md) }; + + if (amd_comgr_destroy_metadata(arg_value_kind_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + + if (arg_value_kind.find("hidden_") == 0) { + if (amd_comgr_destroy_metadata(arg_md) + != AMD_COMGR_STATUS_SUCCESS) + return; + + continue; //Ignore hidden arg + } + amd_comgr_metadata_node_t arg_size_md; if (amd_comgr_metadata_lookup(arg_md, ".size", &arg_size_md) != AMD_COMGR_STATUS_SUCCESS) diff --git a/tests/src/gcc/LaunchKernel.c b/tests/src/gcc/LaunchKernel.c index 042089a71a..d2fc854510 100644 --- a/tests/src/gcc/LaunchKernel.c +++ b/tests/src/gcc/LaunchKernel.c @@ -20,7 +20,7 @@ /* HIT_START * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvcc vdi - * BUILD_CMD: launchkernel.o %cc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc vdi + * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc vdi * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvcc vdi * TEST: %t EXCLUDE_HIP_PLATFORM nvcc vdi * HIT_END @@ -36,7 +36,7 @@ bool LaunchKernelArg() dim3 blocks = {1,1,1}; dim3 threads = {1,1,1}; - HIPCHECK(hipLaunchKernel(kernel, blocks, threads, NULL, 0, 0)); + HIPCHECK(hipLaunchKernel((const void *)kernel, blocks, threads, NULL, 0, 0)); return true; } @@ -52,7 +52,7 @@ bool LaunchKernelArg1() HIPCHECK(hipMalloc((void**)&A_d, sizeof(int))); void* Args[]={&A_d}; - HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args, 0, 0)); + HIPCHECK(hipLaunchKernel((const void *)kernel1, blocks, threads, Args, 0, 0)); // Get the result back to host memory HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); @@ -84,7 +84,7 @@ bool LaunchKernelArg2() HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice)); void* Args[]={&A_d, &B_d}; - HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0)); + HIPCHECK(hipLaunchKernel((const void *)kernel2, blocks, threads, Args,0,0)); // Get the result back to host memory HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost)); @@ -123,7 +123,7 @@ bool LaunchKernelArg3() HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice)); void* Args[]={&A_d, &B_d, &C_d}; - HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0)); + HIPCHECK(hipLaunchKernel((const void *)kernel3, blocks, threads, Args,0,0)); // Get the result back to host memory HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost)); @@ -154,7 +154,7 @@ bool LaunchKernelArg4() struct things t = {2,20,200}; void* Args[]={&A_d, &c, &s, &i, &t}; - HIPCHECK(hipLaunchKernel(kernel4, blocks, threads, Args, 0, 0)); + HIPCHECK(hipLaunchKernel((const void *)kernel4, blocks, threads, Args, 0, 0)); // Get the result back to host memory HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));