I modified vector_copy.c sample from CLOC 0.9 to query the alignment for the kernel symbol segment and it returns the value 3:
* Extract the symbol from the executable.
err = hsa_executable_get_symbol(executable, "", "&__OpenCL_vector_copy_kernel", agent, 0, &symbol);
check(Extract the symbol from the executable, err);
* Extract dispatch information from the symbol
err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT, &kernarg_segment_align);
printf("\nKERNARG_SEGMENT_ALIGNMENT = %d", kernarg_segment_align);
According to the HSA Runtime Specs, this value should be: "the maximum of 16 and the maximum alignment of any of the kernel arguments."
Am I doing something wrong or is the returned value incorrect?
Solved! Go to Solution.
The return value is incorrect. The value should be 16. I have opened a issue on this problem and it will be corrected in a future release.
This is fixed in release version release-v1.0f.2.