HIP dot product fails to build

[psteinb]

Hi, I just wanted to start adding a bare HC stream and started off with the HIPified version in cc90cef and got this:

$ $ make -f HIP.make 
hipcc  -std=c++11 -DHIP main.cpp HIPStream.cu  -o hip-stream
HIPStream.cu:185:10: error: cannot combine with previous 'extern' declaration specifier
  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
         ^
/opt/rocm/include/hip/hcc_detail/host_defines.h:49:24: note: expanded from macro '__shared__'
#define __shared__     tile_static
                       ^
/opt/rocm/hcc/include/hc_defines.h:52:21: note: expanded from macro 'tile_static'
#define tile_static static __attribute__((section("clamp_opencl_local")))
                    ^
HIPStream.cu:185:31: error: expected parameter declarator
  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
                              ^
HIPStream.cu:185:31: error: expected ')'
HIPStream.cu:185:30: note: to match this '('
  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
                             ^
HIPStream.cu:185:21: error: C++ requires a type specifier for all declarations
  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
  ~~~~~~~~~~~~~~~~~ ^
HIPStream.cu:185:21: error:  tile_static can only be applied to a variable declaration
HIPStream.cu:185:41: error: expected ';' at end of declaration
  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
                                        ^
                                        ;
HIPStream.cu:185:56: error: definition of variable with array type needs an explicit size or an initializer
  extern __shared__ __align__(sizeof(T)) unsigned char smem[];
                                                       ^
HIPStream.cu:188:11: error: use of undeclared identifier 'blockDim'
  int i = blockDim.x * blockIdx.x + threadIdx.x;
          ^
HIPStream.cu:188:24: error: use of undeclared identifier 'blockIdx'
  int i = blockDim.x * blockIdx.x + threadIdx.x;
                       ^
HIPStream.cu:188:37: error: use of undeclared identifier 'threadIdx'
  int i = blockDim.x * blockIdx.x + threadIdx.x;
                                    ^
HIPStream.cu:189:26: error: use of undeclared identifier 'threadIdx'
  const size_t local_i = threadIdx.x;
                         ^
HIPStream.cu:192:31: error: use of undeclared identifier 'blockDim'
  for (; i < array_size; i += blockDim.x*gridDim.x)
                              ^
HIPStream.cu:192:42: error: use of undeclared identifier 'gridDim'
  for (; i < array_size; i += blockDim.x*gridDim.x)
                                         ^
HIPStream.cu:195:21: error: use of undeclared identifier 'blockDim'
  for (int offset = blockDim.x / 2; offset > 0; offset /= 2)
                    ^
HIPStream.cu:205:9: error: use of undeclared identifier 'blockIdx'
    sum[blockIdx.x] = tb_sum[local_i];
        ^
15 errors generated.
Died at /opt/rocm//bin/hipcc line 378.
make: *** [hip-stream] Error 1

which comes from this line in HIPStream.cu (btw, I think this file needs a suffix rename as it technically doesn't contain any CUDA code):

extern __shared__ __align__(sizeof(T)) unsigned char smem[];
T *tb_sum = reinterpret_cast<T*>(smem);

[jrprice]

The HIP.make probably needs a little more work to get it to build with HIP via HCC. At present, it works fine for me when targeting NVIDIA devices via nvcc, but I don't currently have a ROCm box set-up in order to make the changes necessary for HCC.

Contributions welcome!

[psteinb]

ok, I understand, maybe @whchung, @bensander et al could help?

[psteinb]

fixed with #26

[tomdeakin]

Closed after merge of #26 into master. Thanks for the contribution!