Dynamic shared memory failes to compile

[jakub-homola]

Trying to compile the HIP program

#include <hip/hip_runtime.h>

__global__ void my_kernel()
{
    extern __shared__ int dyn_shmem[];
}

int main()
{
    int dyn_shmem_size = 64;
    hipLaunchKernelGGL(my_kernel, 4, 32, dyn_shmem_size, 0);

    hipDeviceSynchronize();

    return 0;
}

using the command

g++ -g -O2 -std=c++17 -I/home/jakub/apps/HIP-CPU/include source.hip.cpp -o program.x -ltbb -lpthread

produces the following compiler error

source.hip.cpp: In function ‘void my_kernel()’:
source.hip.cpp:5:37: error: conflicting specifiers in declaration of ‘dyn_shmem’
    5 |     extern __shared__ int dyn_shmem[];
      |                                     ^

This is probably cause by the extern and static specifiers being combined, as __shared__ is defined as #define __shared__ thread_local static. I understand, that using the HIP_DYNAMIC_SHARED macro would solve the issue, but as written in the description of the HIP-CPU library, it "allows CPUs to execute unmodified HIP code", and the extern __shared__ int dyn_shmem[]; is now correct HIP code.

[AlexVlx]

Thank you for reporting this. Sadly, it is rather hard to eschew use of the macro without compiler support, as for your example some special handling is required (allocating storage of the correct type etc., see the macro's expansion). I believe it is doable strictly from user-space, but it's not something that is imminent, and thus for the time being usage of the macro will be required; I will correct the documentation to reflect it.