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: 6182abf5e9]

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;

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

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*