From 62138712cf8866f108b21789dba690baa077a6eb Mon Sep 17 00:00:00 2001
From: Sean Keely <Sean.Keely@amd.com>
Date: Sat, 21 Nov 2020 21:37:58 -0600
Subject: [PATCH] Add asserts and minimum values for kernarg alignment and
 utility functions.

Kernel argument size and alignment queries are not supported on
code object v3.

Change-Id: I1bdd34e2e62132f912ac39d80355efd3456df87c


[ROCm/ROCR-Runtime commit: 6182abf5e992185cca8eb63830d04e8117128f9b]
---
 projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc       | 3 +++
 projects/rocr-runtime/rocrtst/common/helper_funcs.cc          | 2 ++
 .../rocrtst/samples/binary_search/binary_search.cc            | 4 ++++
 3 files changed, 9 insertions(+)

diff --git a/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc b/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc
index 848fca253c..7631211d6d 100755
--- a/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc
+++ b/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc
@@ -335,6 +335,7 @@ hsa_status_t LoadKernelFromObjFile(BaseRocR* test, hsa_agent_t* agent) {
   RET_IF_HSA_UTILS_ERR(err);
   test->set_group_segment_size(val);
 
+  // Remaining queries only supported on code object v3.
   err = hsa_executable_symbol_get_info(kern_sym,
                 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &val);
   RET_IF_HSA_UTILS_ERR(err);
@@ -343,6 +344,8 @@ hsa_status_t LoadKernelFromObjFile(BaseRocR* test, hsa_agent_t* agent) {
   err = hsa_executable_symbol_get_info(kern_sym,
            HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT, &val);
   RET_IF_HSA_UTILS_ERR(err);
+  assert(val >= 16 && "Reported kernarg size is too small.");
+  val = (val == 0) ? 16 : val;
   test->set_kernarg_align(val);
 
   return HSA_STATUS_SUCCESS;
diff --git a/projects/rocr-runtime/rocrtst/common/helper_funcs.cc b/projects/rocr-runtime/rocrtst/common/helper_funcs.cc
index 6e86f6bbc5..f145ffe49e 100755
--- a/projects/rocr-runtime/rocrtst/common/helper_funcs.cc
+++ b/projects/rocr-runtime/rocrtst/common/helper_funcs.cc
@@ -48,6 +48,7 @@
 #ifndef _WIN32
 #include <unistd.h>
 #endif
+#include <assert.h>
 #include <cmath>
 #include <iostream>
 #include <string>
@@ -175,6 +176,7 @@ Compare(const double* refData, const double* data,
 
 intptr_t
 AlignDown(intptr_t value, size_t alignment) {
+    assert(alignment != 0 && "Zero alignment");
     return (intptr_t) (value & ~(alignment - 1));
 }
 
diff --git a/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc b/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc
index 07db231034..627790a32f 100755
--- a/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc
+++ b/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc
@@ -419,6 +419,7 @@ hsa_status_t LoadKernelFromObjFile(BinarySearch* bs) {
                                                      &bs->group_segment_size);
   RET_IF_HSA_ERR(err);
 
+  // Remaining queries not supported on code object v3.
   err = hsa_executable_symbol_get_info(kern_sym,
                       HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
                                                            &bs->kernarg_size);
@@ -428,6 +429,8 @@ hsa_status_t LoadKernelFromObjFile(BinarySearch* bs) {
                  HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT,
                                                           &bs->kernarg_align);
   RET_IF_HSA_ERR(err);
+  assert(bs->kernarg_align >= 16 && "Reported kernarg size is too small.");
+  bs->kernarg_align = (bs->kernarg_align == 0) ? 16 : bs->kernarg_align;
 
   return err;
 }
@@ -465,6 +468,7 @@ hsa_status_t AgentMemcpy(void* dst, const void* src,
 // return a value that has the specified alignment.
 static intptr_t
 AlignDown(intptr_t value, size_t alignment) {
+  assert(alignment != 0 && "Zero alignment");
   return (intptr_t) (value & ~(alignment - 1));
 }
 static void*