HIP-CPU icon indicating copy to clipboard operation
HIP-CPU copied to clipboard

Dynamic shared memory failes to compile

Open jakub-homola opened this issue 2 years ago • 0 comments

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.

jakub-homola avatar Sep 30 '21 07:09 jakub-homola