From 59ce739c08d64d7b2a45756015fc54816f54fea1 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 20 Jul 2017 14:41:30 +0530 Subject: [PATCH 01/10] Renable frexp(f) device math function Change-Id: I53c022b8ddf38cd17ddb42eba457b9020db66395 [ROCm/clr commit: cb5cf20c21e60ca9cfb5175bf3d03ab666a03636] --- projects/clr/hipamd/include/hip/hcc_detail/math_functions.h | 4 ++-- .../tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp | 2 +- .../tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index 9faff2743a..79bacf274b 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h @@ -60,7 +60,7 @@ __device__ float fmaf(float x, float y, float z); __device__ float fmaxf(float x, float y); __device__ float fminf(float x, float y); __device__ float fmodf(float x, float y); -//__device__ float frexpf(float x, int* nptr); +__device__ float frexpf(float x, int* nptr); __device__ float hypotf(float x, float y); __device__ float ilogbf(float x); __device__ int isfinite(float a); @@ -146,7 +146,7 @@ __device__ double fma(double x, double y, double z); __device__ double fmax(double x, double y); __device__ double fmin(double x, double y); __device__ double fmod(double x, double y); -//__device__ double frexp(double x, int *nptr); +__device__ double frexp(double x, int *nptr); __device__ double hypot(double x, double y); __device__ double ilogb(double x); __device__ int isfinite(double x); diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp index f4f7ab0479..0a81d111c2 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp @@ -69,7 +69,7 @@ __device__ void double_precision_math_functions() fmax(0.0, 0.0); fmin(0.0, 0.0); fmod(0.0, 1.0); -// frexp(0.0, &iX); + frexp(0.0, &iX); hypot(1.0, 0.0); ilogb(1.0); isfinite(0.0); diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp index de3dec35ef..a52b1a22fe 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp @@ -70,7 +70,7 @@ __device__ void single_precision_math_functions() fmaxf(0.0f, 0.0f); fminf(0.0f, 0.0f); fmodf(0.0f, 1.0f); - //frexpf(0.0f, &iX); + frexpf(0.0f, &iX); hypotf(1.0f, 0.0f); ilogbf(1.0f); isfinite(0.0f); From 952bd58407c6c387018e615dc0b4814e677ebede Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 21 Jul 2017 15:50:12 -0500 Subject: [PATCH 02/10] fixed device selection during compilation to use rocm_agent_enumerator 1. Changed hipcc to use rocm_agent_enumerator 2. Changed square sample test to use device variable [ROCm/clr commit: b3f8c3fb8e6c9724bc6d66a93d852ea7b77f48af] --- projects/clr/hipamd/bin/hipcc | 22 +++++++++++++++++++ .../samples/0_Intro/square/square.hipref.cpp | 5 +++-- 2 files changed, 25 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index bcd3e3a591..2dceaa295f 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -103,6 +103,28 @@ if ($HIP_PLATFORM eq "hcc") { $HIPLDFLAGS = `${HCC_HOME}/bin/hcc-config --ldflags`; + $ROCM_AGENT_ENUM = "${ROCM_PATH}/bin/rocm_agent_enumerator"; + + my $myAgents = `${ROCM_AGENT_ENUM} -t GPU`; + my @agentsLine = split('\n', $myAgents); + + foreach my $val (@agentsLine) { + if($val eq "gfx701") { + $target_gfx701 = 1; + } + if($val eq "gfx801") { + $target_gfx801 = 1; + } + if($val eq "gfx802") { + $target_gfx802 = 1; + } + if($val eq "gfx803") { + $target_gfx803 = 1; + } + if($val eq "gfx900") { + $target_gfx900 = 1; + } + } #### GCC system includes workaround #### $HCC_WA_FLAGS = " "; diff --git a/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp b/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp index e694bfb8a4..167cb135f4 100644 --- a/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp +++ b/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp @@ -54,9 +54,10 @@ int main(int argc, char *argv[]) float *A_h, *C_h; size_t N = 1000000; size_t Nbytes = N * sizeof(float); - + static int device = 0; + CHECK(hipSetDevice(device)); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + CHECK(hipGetDeviceProperties(&props, device/*deviceID*/)); printf ("info: running on device %s\n", props.name); #ifdef __HIP_PLATFORM_HCC__ printf ("info: architecture on AMD GPU device is: %d\n",props.gcnArch); From 339cf92137128fcdfa8cb0d28e7b21e04b0601d8 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 20 Oct 2017 14:49:29 +0000 Subject: [PATCH 03/10] Modify device properties to use pool API. - Also better error code checking [ROCm/clr commit: e738fa66b5af4c9c8c1d1578825dfca5d8b55b69] --- projects/clr/hipamd/src/hip_hcc.cpp | 29 +++++++++++++++++------------ 1 file changed, 17 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 2c98ac804f..aa8dbd0072 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -695,26 +695,26 @@ int checkAccess(hsa_agent_t agent, hsa_amd_memory_pool_t pool) return access; } -hsa_status_t get_region_info(hsa_region_t region, void* data) +hsa_status_t get_pool_info(hsa_amd_memory_pool_t pool, void* data) { hsa_status_t err; hipDeviceProp_t* p_prop = reinterpret_cast(data); uint32_t region_segment; - // Get region segment - err = hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, ®ion_segment); + // Get pool segment + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, ®ion_segment); ErrorCheck(err); switch(region_segment) { case HSA_REGION_SEGMENT_READONLY: - err = hsa_region_get_info(region, HSA_REGION_INFO_SIZE, &(p_prop->totalConstMem)); break; - /* case HSA_REGION_SEGMENT_PRIVATE: - cout<<"PRIVATE"<totalConstMem)); break; case HSA_REGION_SEGMENT_GROUP: - err = hsa_region_get_info(region, HSA_REGION_INFO_SIZE, &(p_prop->sharedMemPerBlock)); break; + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &(p_prop->sharedMemPerBlock)); + printf ("shared_mem err=%d mem=%zu\n", err, p_prop->sharedMemPerBlock); + break; default: break; } - return HSA_STATUS_SUCCESS; + return err; } @@ -748,8 +748,10 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) // Set some defaults in case we don't find the appropriate regions: prop->totalGlobalMem = 0; prop->totalConstMem = 0; - prop-> maxThreadsPerMultiProcessor = 0; + prop->maxThreadsPerMultiProcessor = 0; prop->regsPerBlock = 0; + prop->totalConstMem = 0; + prop->sharedMemPerBlock = 0; if (_hsaAgent.handle == -1) { return hipErrorInvalidDevice; @@ -849,15 +851,18 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) prop-> maxThreadsPerMultiProcessor = prop->warpSize*max_waves_per_cu; // Get memory properties - err = hsa_agent_iterate_regions(_hsaAgent, get_region_info, prop); + err = hsa_amd_agent_iterate_memory_pools(_hsaAgent, get_pool_info, prop); + if (err == HSA_STATUS_INFO_BREAK) { + err = HSA_STATUS_SUCCESS; + } DeviceErrorCheck(err); - // Get the size of the region we are using for Accelerator Memory allocations: + // Get the size of the pool we are using for Accelerator Memory allocations: hsa_region_t *am_region = static_cast(_acc.get_hsa_am_region()); err = hsa_region_get_info(*am_region, HSA_REGION_INFO_SIZE, &prop->totalGlobalMem); DeviceErrorCheck(err); // maxSharedMemoryPerMultiProcessor should be as the same as group memory size. - // Group memory will not be paged out, so, the physical memory size is the total shared memory size, and also equal to the group region size. + // Group memory will not be paged out, so, the physical memory size is the total shared memory size, and also equal to the group pool size. prop->maxSharedMemoryPerMultiProcessor = prop->totalGlobalMem; // Get Max memory clock frequency From 277454248728483febc0fd43f55313f4b45917da Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 20 Oct 2017 15:50:18 +0000 Subject: [PATCH 04/10] Update device properties. - clear properties to defined initial state. - enable some property flags which are now supported. [ROCm/clr commit: ed85b15c3e0afb93406ab74cdfb870931af5f3f8] --- .../clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 3 +-- projects/clr/hipamd/src/hip_hcc.cpp | 12 +++--------- 2 files changed, 4 insertions(+), 11 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 379fc05f5b..9fcc02b2d7 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -108,13 +108,12 @@ extern int HIP_TRACE_API; #if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0) // Device compile and not host compile: -//TODO-HCC enable __HIP_ARCH_HAS_ATOMICS__ when HCC supports these. // 32-bit Atomics: #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (1) #define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (1) #define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (1) #define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (1) -#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0) +#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (1) // 64-bit Atomics: #define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (1) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index aa8dbd0072..d2d28c2a93 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -745,13 +745,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) hipError_t e = hipSuccess; hsa_status_t err; - // Set some defaults in case we don't find the appropriate regions: - prop->totalGlobalMem = 0; - prop->totalConstMem = 0; - prop->maxThreadsPerMultiProcessor = 0; - prop->regsPerBlock = 0; - prop->totalConstMem = 0; - prop->sharedMemPerBlock = 0; + memset(prop, 0, sizeof(hipDeviceProp_t)); if (_hsaAgent.handle == -1) { return hipErrorInvalidDevice; @@ -882,7 +876,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) prop->arch.hasGlobalFloatAtomicExch = 1; prop->arch.hasSharedInt32Atomics = 1; prop->arch.hasSharedFloatAtomicExch = 1; - prop->arch.hasFloatAtomicAdd = 0; + prop->arch.hasFloatAtomicAdd = 1; // supported with CAS loop, but is supported prop->arch.hasGlobalInt64Atomics = 1; prop->arch.hasSharedInt64Atomics = 1; prop->arch.hasDoubles = 1; @@ -890,7 +884,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) prop->arch.hasWarpBallot = 1; prop->arch.hasWarpShuffle = 1; prop->arch.hasFunnelShift = 0; // TODO-hcc - prop->arch.hasThreadFenceSystem = 0; // TODO-hcc + prop->arch.hasThreadFenceSystem = 1; prop->arch.hasSyncThreadsExt = 0; // TODO-hcc prop->arch.hasSurfaceFuncs = 0; // TODO-hcc prop->arch.has3dGrid = 1; From b5e3de043cb23b6370896affc5fac05f5cbdbda0 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 20 Oct 2017 13:24:04 -0700 Subject: [PATCH 05/10] Remove printf [ROCm/clr commit: b4c7876244dd35b6ef2e36aa90ef37abb2d55fac] --- projects/clr/hipamd/src/hip_hcc.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index d2d28c2a93..6a614b1826 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -710,7 +710,6 @@ hsa_status_t get_pool_info(hsa_amd_memory_pool_t pool, void* data) err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &(p_prop->totalConstMem)); break; case HSA_REGION_SEGMENT_GROUP: err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &(p_prop->sharedMemPerBlock)); - printf ("shared_mem err=%d mem=%zu\n", err, p_prop->sharedMemPerBlock); break; default: break; } From 5f24981c6128259a53c2b0f89440f45144b81dc8 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 23 Oct 2017 16:47:40 -0400 Subject: [PATCH 06/10] replace __hcc_workweek__ with HC_FEATURE_PRINTF flag [ROCm/clr commit: da7c37947c81e214fd8aac4072a80fed76a5d4ff] --- projects/clr/hipamd/src/hip_hcc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 2d67c31fe7..ae53d7ae45 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -616,7 +616,7 @@ void ihipDevice_t::locked_reset() //FIXME - Calling am_memtracker_reset is really bad since it destroyed all buffers allocated by the HCC runtime as well //such as the printf buffer. Re-initialze the printf buffer as a workaround for now. -#if (__hcc_workweek__ >= 17423) +#ifdef HC_FEATURE_PRINTF Kalmar::getContext()->initPrintfBuffer(); #endif }; From 1e4a2a7dd0d8ef42c2c9b62d7dbd465f97e9c2e1 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 23 Oct 2017 16:59:14 -0400 Subject: [PATCH 07/10] add a printf device function [ROCm/clr commit: 36c2795dd4045b74fe38f6429bcecf098bfd80fd] --- .../include/hip/hcc_detail/hip_runtime.h | 20 +++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 379fc05f5b..07c7a0d4b3 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -50,10 +50,16 @@ THE SOFTWARE. #include +// define HIP_ENABLE_PRINTF to enable printf +#ifdef HIP_ENABLE_PRINTF + #define HCC_ENABLE_ACCELERATOR_PRINTF 1 +#endif + //--- // Remainder of this file only compiles with HCC #if defined __HCC__ #include +#include "hc_printf.hpp" //TODO-HCC-GL - change this to typedef. //typedef grid_launch_parm hipLaunchParm ; @@ -420,6 +426,20 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) } +#ifdef __HCC_ACCELERATOR__ + +#ifdef HC_FEATURE_PRINTF +template +static inline __device__ void printf(const char* format, All... all) { + hc::printf(format, all...); +} +#else +template +static inline __device__ void printf(const char* format, All... all) { } +#endif + +#endif + #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) From 9997997fe5c509a7206fef721db585bda5bf6a98 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 23 Oct 2017 18:31:13 -0400 Subject: [PATCH 08/10] add an unit test for printf [ROCm/clr commit: f1992c5b113b6c271b9b135cef9c2a29e8317384] --- .../tests/src/kernel/hipPrintfKernel.cpp | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp diff --git a/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp b/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp new file mode 100644 index 0000000000..482098fd54 --- /dev/null +++ b/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp @@ -0,0 +1,38 @@ +/* +Copyright (c) 2015-2017 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 + * RUN: %t + * HIT_END + */ + +#define HIP_ENABLE_PRINTF + +#include"test_common.h" + +__global__ void run_printf(hipLaunchParm lp){ + printf("Hello World\n"); +} + +int main(){ +hipLaunchKernel(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0); +hipDeviceSynchronize(); +passed(); +} From f3b483341349a7bb21018cb0d495f5d0519dbf2d Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Tue, 24 Oct 2017 00:55:41 -0400 Subject: [PATCH 09/10] printf support for module API [ROCm/clr commit: 4ed3e9e15f44670487becde77f9d6846fd1eb55c] --- projects/clr/hipamd/src/hip_module.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 74b0092b84..85c88a1826 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -259,6 +259,11 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); CHECKLOG_HSA(status, hipErrorNotInitialized); + status = hsa_executable_agent_global_variable_define((*module)->executable, agent, + "_ZN2hc13printf_bufferE", + Kalmar::getContext()->getPrintfBufferPointerVA()); + CHECKLOG_HSA(status, hipErrorNotInitialized); + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); CHECKLOG_HSA(status, hipErrorNotInitialized); From d97b4b55cc46bb4765a56f371291c8682d088a49 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 25 Oct 2017 12:00:02 -0400 Subject: [PATCH 10/10] add HC_FEATURE_PRINTF around the printf buffer definition [ROCm/clr commit: 085d63099c2767323a410d29d7a5a2916a92178c] --- projects/clr/hipamd/src/hip_module.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 85c88a1826..962b6084f1 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -259,10 +259,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); CHECKLOG_HSA(status, hipErrorNotInitialized); +#ifdef HC_FEATURE_PRINTF status = hsa_executable_agent_global_variable_define((*module)->executable, agent, "_ZN2hc13printf_bufferE", Kalmar::getContext()->getPrintfBufferPointerVA()); CHECKLOG_HSA(status, hipErrorNotInitialized); +#endif status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); CHECKLOG_HSA(status, hipErrorNotInitialized);