... because, AFAIK, Rusticl doesn't support program scope global variables (or, IOW, crossWorkgroup variables in global scope in SPIR-V). Following debug message appears on a debug build of Rusticl while running HIP programs:
SPIR-V parsing FAILED:
In file ../../mesa/src/compiler/spirv/vtn_variables.c:2394
Initializer for CrossWorkgroup variable 6 not yet supported in Mesa.
692 bytes into the SPIR-V binary
This occurs on HIP device programs without global variables like on this simple case: __global__ void k(int *Out) { *Out = 123; }. That's because chipStar unconditionally inserts code for device-side dynamic memory allocations which introduces a global variable __chip_var___chipspv_device_heap into SPIR-V modules.
... because, AFAIK, Rusticl doesn't support program scope global variables (or, IOW, crossWorkgroup variables in global scope in SPIR-V). Following debug message appears on a debug build of Rusticl while running HIP programs:
This occurs on HIP device programs without global variables like on this simple case:
__global__ void k(int *Out) { *Out = 123; }. That's because chipStar unconditionally inserts code for device-side dynamic memory allocations which introduces a global variable__chip_var___chipspv_device_heapinto SPIR-V modules.