std::numeric_limits<_T>::infinity() compilation problem

[ekondis]

Using the std::numeric_limits<_T>::infinity() function within a kernel returns compilation error:

error:  'std::numeric_limits<int>::infinity': no overloaded function has restriction specifiers that are compatible
      with the ambient context 'trivial_kernel'
        *data = std::numeric_limits<int>::infinity();
                                                  ^

To reproduce use the following HIP code :

#include <iostream>
#include <hip/hip_runtime.h>

__global__ void
trivial_kernel(hipLaunchParm lp, int *data){
	*data = std::numeric_limits<int>::infinity();
}

int main(void){
	int *data;
	hipMalloc((void**)&data, sizeof(int));
	hipLaunchKernel(HIP_KERNEL_NAME(trivial_kernel), dim3(1), dim3(1), 0, 0, data);
	int res = -1;
	hipMemcpy(&res, data, sizeof(int), hipMemcpyDeviceToHost);
	hipFree(data);
	std::cout << res << std::endl;

	return 0;
}

[pfultz2]

This is a problem in hcc actually:

#include <limits>

void f(int *data) [[hc]] [[cpu]] {
    *data = std::numeric_limits<int>::infinity();
}

And there is only an error when using both [[hc]] [[cpu]] together.

[ekondis]

Do I have to open an issue for hcc?

AFAIK, under ROCm v1.6 it was compiling fine.

[pfultz2]

Do I have to open an issue for hcc?

I would, so that they can see it. I dont know if they are all subscribed to HIP issues.

[pramenku]

Is defect opened againt HCC? Can someone please post the issue#.

[AlexVlx]

@ekondis We do not officially support this type of use of anything that is not __device__ or __host__ __device__ in a __global__ / __device__ context. While this may sometime "work", it is not something that we claim should work. Therefore, please do not do that - I observe that for CUDA numeric_limits is not used, therefore this will not yield asymmetry. Thank you.

P.S.: using hipLaunchParm is deprecated and should be purged with extreme prejudice, if you have the cycles.

[ekondis]

Thank you @AlexVlx. I'll replace this call with a simple macro/constant.