cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

r_potter
Journeyman III

Finalizing HSA kernels with global variable declarations

I have a kernel that declares but not defines a global variable that is intended to be shared between CPU and GPU, with the aim of using hsa_executable_global_variable_define() to associate global_var with an existing variable in CPU code.

decl prog global_u32 &global_var;

prog kernel &class_write_global_kernel(kernarg_u64 %f)

{

  st_global_align(4)_u32 1, [&global_var];

  ret;

};

This causes a segfault in hsa_ext_program_finalize(). If I include the definition of the variable in the kernel (i.e. prog global_u32 &global_var) then the kernel finalizes successfully, but that approach prevents me from using hsa_executable_global_variable_define().

0 Likes
2 Replies
jedwards
Staff

What version of the runtime and finalizer are you using? Also, please provide the host code indicating how you define the global variable for the kernel.

Thank you.

0 Likes

It was with the May/June release of the runtime and KFD 1.4. However, whilst checking versions I just noticed that a new version was just released.

Using 1.0.3 and KFD 1.6 appears to resolve the issue, so please disregard this.

0 Likes