In CUDA kernel template function, how to test types?

user_2525479 0
Viewed 215

I have a CUDA kernel template function like this:

    template <typename scalar_t, typename accscalar_t, typename index_type, int indexing_kind>
    __global__ void lstm_cell_forward(
                TensorInfo<scalar_t, index_type> input,
                TensorInfo<scalar_t, index_type> hidden,
                TensorInfo<scalar_t, index_type> bias1,
                TensorInfo<scalar_t, index_type> bias2,
                TensorInfo<scalar_t, index_type> _cx,
                TensorInfo<scalar_t, index_type> _hy,
                TensorInfo<scalar_t, index_type> _cy,
                TensorInfo<scalar_t, index_type> workspace,
                index_type hsz,
                index_type totalElements) {
...
 scalar_t iig = DEVICE_LINEAR_GET(input, offset+0*hsz);
      scalar_t ifg = DEVICE_LINEAR_GET(input, offset+1*hsz);
      scalar_t icg = DEVICE_LINEAR_GET(input, offset+2*hsz);
      scalar_t iog = DEVICE_LINEAR_GET(input, offset+3*hsz);
...
}

I want to add printf() inside this kernel function to print the value of iig, ifg, icg, iog only when scalar_t is float. I tried to use typeid(float) == typeid(iig) to accomplish this, but "typeinfo.h" is not supported in CUDA code appearently.

If I just ommit the if statement

   printf("iig = %f, ifg = %f, icg = %f, iog = %f\n",
    iig, ifg, icg, iog);

And then go ahead and compile, the compiler will try to compile different combinations of template parameter types even for c10:Half type, a pytorch defined datatype. This will throw an error.

So my question is how to write the comparison of scalar_t to check if it equals to float?

I adopted the suggestions answered by @Robert Crovella, and got an error:

liwei.dai@854380cd7bb1:~/tests/cuda-debug-case/cuda/tmp$ !b
bash bi_make_and_run.sh 
test.cu:7:20: error: no template named 'is_same_v' in namespace 'cuda::std'; did you mean 'is_same'?
  if (::cuda::std::is_same_v<T, float>) printf("val is a float: %f\n", val);
      ~~~~~~~~~~~~~^~~~~~~~~
                   is_same
/opt/sw_home/local/cuda/include/cuda/std/std/detail/libcxx/include/type_traits:877:65: note: 'is_same' declared here
template <class _Tp, class _Up> struct _LIBCUDACXX_TEMPLATE_VIS is_same           : public false_type {};
                                                                ^
test.cu:7:39: error: expected unqualified-id
  if (::cuda::std::is_same_v<T, float>) printf("val is a float: %f\n", val);

But if I tried to use cuda::std::is_same<T, float>::value, and it worked. I checked the source code in the type_traits file, they just use is_same_v to call is_same::value. But I don't know why.

cuda c++
1 Answers

So my question is how to write the comparison of scalar_t to check if it equals to float?

I believe you can use features of libcu++ (type traits) for this:

$ cat t1868.cu
#include <cuda/std/type_traits>
#include <cstdio>
template <typename T>
__global__ void k(T val){
  if (::cuda::std::is_same_v<T, float>) printf("val is a float: %f\n", val);
  else printf("val is not a float\n");
}

int main(){

  float f = 1.2345;
  k<<<1,1>>>(f);
  double d = 1.2345;
  k<<<1,1>>>(d);
  cudaDeviceSynchronize();
}

$ nvcc -std=c++14 t1868.cu -o t1868
$ cuda-memcheck ./t1868
========= CUDA-MEMCHECK
val is a float: 1.234500
val is not a float
========= ERROR SUMMARY: 0 errors
$
user_16959600
Related