From 480049fb09c4e6011888fa4ade2ed95a98ff7c07 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Sun, 16 Feb 2020 17:58:04 -0500 Subject: [PATCH] Fix bug of hip/samples/0_Intro/hcc_dialects/vadd_hip For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need insert hmod->kernargs[name_str] which is empty. Change-Id: I95f818d0525da84452e66c5778f0648a643843c7 [ROCm/hip commit: 30da92e2afcfe5cd1b34e3389c0a92974fcd3b8c] --- projects/hip/src/hip_module.cpp | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index b8952f8a68..9c1e2998c7 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -1040,11 +1040,15 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t if (kernel.handle == 0u) return hipErrorNotFound; + //For hipModuleLoad(), hmod->kernargs must contain a args with key + //name_str or namekd_str. + //For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need + //insert hmod->kernargs[name_str] 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; + if (it != hmod->kernargs.end()) { + name_str = namekd_str; } } @@ -1052,7 +1056,8 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t // below, due to hipFunction_t being a pointer to ihipModuleSymbol_t. func[0][0] = *static_cast( - Kernel_descriptor{kernel_object(kernel), name_str, it->second}); + Kernel_descriptor{kernel_object(kernel), name_str, + it != hmod->kernargs.end() ? it->second : hmod->kernargs[name_str]}); return hipSuccess; }