search

CHIP-SPV / chipStar

chipStar is a tool for compiling and running HIP/CUDA on SPIR-V via OpenCL or Level Zero APIs.
Other
185 stars 30 forks source link

HeCBench: lebesgue-hip fails to compile due to incompatible atomicMax function #627

Open kballeda opened 12 months ago

kballeda commented 12 months ago

As part of HeCBench analysis, lebesgue-hip benchmark failed to compile due to atomicMax incompatibility. Equivalent benchmark in CUDA passes. I did a comparison of atomicMax support in CUDA and ChipStar found couple of differences in support.

ChipStar:

kballeda@x1002c6s6b0n0:~/ANL_WORK/HeCBench/lebesgue-hip> make
hipcc  -std=c++14 -Wall -O3 -c kernels.cu -o kernels.o
kernels.cu:40:3: error: no matching function for call to 'atomicMax'
  atomicMax(lmax, t);
  ^~~~~~~~~
/home/kballeda/local/chipstar_0914/include/hip/devicelib/atomics.hh:217:36: note: candidate function not viable: no known conversion from 'double *__restrict' to 'int *' for 1st argument
extern "C++" inline __device__ int atomicMax(int *address, int val) {
                                   ^
/home/kballeda/local/chipstar_0914/include/hip/devicelib/atomics.hh:223:45: note: candidate function not viable: no known conversion from 'double *__restrict' to 'unsigned int *' for 1st argument
extern "C++" inline __device__ unsigned int atomicMax(unsigned int *address,
                                            ^
/home/kballeda/local/chipstar_0914/include/hip/devicelib/atomics.hh:231:1: note: candidate function not viable: no known conversion from 'double *__restrict' to 'unsigned long long *' for 1st argument
atomicMax(unsigned long long *address, unsigned long long val) {

CUDA:

kballeda@ortce-a100-80G2:~/ANL_WORK/0906_cuda_2/HeCBench/lebesgue-cuda$ make
nvcc  -std=c++14 -Xcompiler -Wall -arch=sm_60 -O3 -c main.cpp -o main.o
nvcc  -std=c++14 -Xcompiler -Wall -arch=sm_60 -O3 -c utils.cpp -o utils.o
nvcc  -std=c++14 -Xcompiler -Wall -arch=sm_60 -O3 -c kernels.cu -o kernels.o
nvcc  -std=c++14 -Xcompiler -Wall -arch=sm_60 -O3 main.o utils.o kernels.o -o main

Path to lebesgue-hip in HeCBench: https://github.com/zjin-lcf/HeCBench/tree/master/lebesgue-hip

pjaaskel commented 12 months ago

We just seem to miss the implementation. There's one for unsigned long long. The restrict qual should not make a difference here.