Is it possible to use std::conditional in a CUDA kernel?

Question

I'm looking to template a CUDA kernel. Neither Google, not SO were particularly forthcoming. Within this kernel, I would like to have a typedef that is dependent on the typename. I tried the following (assuming/hoping) it would work similarly to C++ but it didn't.

#include <type_traits>  // for std::conditional. Not sure if it's necessary
#include <cuda_runtime.h>
#include <cuComplex.h>

template <typename fType>
__global__ void DummyKernel() {
    typedef std::conditional<sizeof(fType) == sizeof(double), cuDoubleComplex, cuFloatComplex>::type cfType;
}

This produces the error

nontype "std::conditional<_Test, _Ty1, _Ty2>::type [with _Test=<expression>, _Ty1=cuDoubleComplex, _Ty2=cuFloatComplex]" is not a type name

Is there a way to do what I want? I'm using CUDA 5.5 and VS2012.

@catscradle I was about to say that I'm using VS and it doesn't comply with the standard, when I realized that the Nvidia compiler is separate. I have to break my stupid lazy habits. I'll accept this as the answer if posted. — Avi Ginsburg, Jan 13 '14 at 14:43

score 2 · Answer 1 · answered Jan 13 '14 at 14:59

Note that for those who want to achieve the same result without relying on C++11 (using C++11 with CUDA may not be as straightforward for now), you can use Boost MPL instead:

#include <boost/type_traits/conditional.hpp>

template <typename fType>
__global__ void DummyKernel() {
    typedef typename boost::conditional<sizeof(fType) == sizeof(double), cuDoubleComplex, cuFloatComplex>::type cfType;
// Which is an equivalent of:
// typedef typename boost::mpl::if_<boost::mpl::bool_<sizeof(fType) == sizeof(double)>, cuDoubleComplex, cuFloatComplex>::type cfType;

    // Quick size check
    printf("size = %u\n", sizeof (cfType));
}

Compilation is simply (adapt to your CUDA architecture):

nvcc test.cu -o test -gencode arch=compute_30,code=sm_30

Then, you get:

DummyKernel<float>  ---> size = 8
DummyKernel<double> ---> size = 16

Tested with CUDA 5.5 and Boost 1.55 on Arch Linux.

score 1 · Accepted Answer · answered Jan 13 '14 at 14:57

1

std::conditional depends on a template parameter fType, so you must put the typename keyword after the typedef.

answered Jan 13 '14 at 14:57

catscradle

1,709
11
18

2

Can you use `std::conditional` at all in a CUDA kernel? – Vitality Jan 13 '14 at 14:58
@JackOLantern I have no idea. I assumed from OP's comment that it worked for him... – catscradle Jan 13 '14 at 15:02
@JackOLantern All of the `std::conditional`'s code is in the header, so what is there to compile? – catscradle Jan 13 '14 at 15:06

Is it possible to use std::conditional in a CUDA kernel?

2 Answers2