diff --git a/projects/rocr-runtime/samples/common/hsa_base_util.cpp b/projects/rocr-runtime/samples/common/hsa_base_util.cpp index 2e31e20520..001546b53d 100644 --- a/projects/rocr-runtime/samples/common/hsa_base_util.cpp +++ b/projects/rocr-runtime/samples/common/hsa_base_util.cpp @@ -114,7 +114,6 @@ bool HSA_UTIL::HsaInit() check("Error in freezing executable object", err); // Get symbol handle. - hsa_executable_symbol_t kernelSymbol; err = hsa_executable_get_symbol(hsaExecutable, NULL, hsa_kernel_name, device, 0, &kernelSymbol); check("get symbol handle", err); @@ -173,6 +172,14 @@ double HSA_UTIL::Run(int dim, int group_x, int group_y, int group_z, int s_size, local_dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; local_dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; local_dispatch_packet.kernel_object = codeHandle; + + // Specify amount of private segment size (in bytes) that is needed per work-item + // Retrieve the amount of private memory needed + uint32_t private_mem_size = 0; + hsa_executable_symbol_get_info(kernelSymbol, + HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &private_mem_size); + local_dispatch_packet.private_segment_size = private_mem_size; + /////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// /* diff --git a/projects/rocr-runtime/samples/common/hsa_base_util.h b/projects/rocr-runtime/samples/common/hsa_base_util.h index e9076cf967..807899cbc4 100644 --- a/projects/rocr-runtime/samples/common/hsa_base_util.h +++ b/projects/rocr-runtime/samples/common/hsa_base_util.h @@ -48,6 +48,7 @@ class HSA_UTIL{ hsa_ext_module_t module; hsa_ext_program_t hsa_program; hsa_executable_t hsaExecutable; + hsa_executable_symbol_t kernelSymbol; hsa_code_object_t code_object; uint64_t codeHandle; hsa_signal_t hsa_signal; diff --git a/projects/rocr-runtime/samples/common/hsa_test.cpp b/projects/rocr-runtime/samples/common/hsa_test.cpp index a659c5eb71..52f99d32b5 100644 --- a/projects/rocr-runtime/samples/common/hsa_test.cpp +++ b/projects/rocr-runtime/samples/common/hsa_test.cpp @@ -151,16 +151,16 @@ HsaTest::Kernel::Kernel(hsa_agent_t agent, std::string hsail_text) HsaTest::Kernel::~Kernel() { Cleanup(); } uint64_t HsaTest::Kernel::GetCodeHandle(const char* kernel_name) { - hsa_executable_symbol_t kernel_symbol = {0}; + kernel_symbol_ = {0}; if (HSA_STATUS_SUCCESS != hsa_executable_get_symbol(executable_, NULL, kernel_name, agent_, 0, - &kernel_symbol)) { + &kernel_symbol_)) { return 0; } uint64_t code_handle = 0; if (HSA_STATUS_SUCCESS != - hsa_executable_symbol_get_info(kernel_symbol, + hsa_executable_symbol_get_info(kernel_symbol_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &code_handle)) { return 0; @@ -169,6 +169,14 @@ uint64_t HsaTest::Kernel::GetCodeHandle(const char* kernel_name) { return code_handle; } +hsa_status_t HsaTest::Kernel::GetScratchSize(uint32_t* size) { + + hsa_status_t status; + status = hsa_executable_symbol_get_info(kernel_symbol_, + HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, size); + return status; +} + void HsaTest::Kernel::Initialize() { CreateProgramFromHsailFile(); CreateCodeObjectAndExecutable(); diff --git a/projects/rocr-runtime/samples/common/hsa_test.h b/projects/rocr-runtime/samples/common/hsa_test.h index 96c191574b..8fd34ab22a 100644 --- a/projects/rocr-runtime/samples/common/hsa_test.h +++ b/projects/rocr-runtime/samples/common/hsa_test.h @@ -74,6 +74,7 @@ class HsaTest { virtual ~Kernel(); uint64_t GetCodeHandle(const char* kernel_name); + hsa_status_t GetScratchSize(uint32_t* size); protected: virtual void Initialize(); @@ -92,6 +93,7 @@ class HsaTest { hsa_ext_program_t program_; hsa_code_object_t code_object_; hsa_executable_t executable_; + hsa_executable_symbol_t kernel_symbol_; std::string hsail_file_; };