Update Private Segment Size parameter of the dispatch packet

[git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1254638]


[ROCm/ROCR-Runtime commit: b93946790d]

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;
+
 	///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
 
 	/*

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;

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();

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_;
   };