Перегрузка функции перемешивания CUDA делает оригинальные невидимыми

Я пытаюсь реализовать свою собственную 64-битную функцию перемешивания в CUDA. Однако, если я сделаю это так:

static __inline__ __device__ double __shfl_xor(double var, int laneMask, int width=warpSize)
{
    int hi, lo;
    asm volatile( "mov.b64 { %0, %1 }, %2;" : "=r"(lo), "=r"(hi) : "d"(var) );
    hi = __shfl_xor( hi, laneMask, width );
    lo = __shfl_xor( lo, laneMask, width );
    return __hiloint2double( hi, lo );
}

Все последующие вызовы __shfl_xor будут создаваться из этой 64-битной версии, независимо от типа аргумента. Например, если я делаю

int a;
a = __shfl_xor( a, 16 );

Он по-прежнему будет использовать двойную версию. Обходной путь может заключаться в использовании разных имен функций. Но так как я вызываю эту функцию перемешивания из функции шаблона, использование разных имен означает, что мне нужно сделать другую версию для 64-битных с плавающей запятой, что не совсем аккуратно.

Итак, как я могу перегрузить функцию __shfl_xor(double,...) и в то же время убедиться, что __shfl_xor(int,...) можно вызвать соответствующим образом?


cuda shuffle ptx
person Rainn    schedule 11.04.2013    source источник
comment
@RoBiK спасибо за комментарий...   -  person Rainn    schedule 12.04.2013

Ответы (1)


arrow_upward
2
arrow_downward

Все целочисленные типы и числа с плавающей запятой могут быть преобразованы в двойные. Если у вас есть выбор между встроенной функцией и вашей специализированной двойной функцией, компилятор здесь может выбрать вашу для всех типов.

Пробовали ли вы создать функцию с другим именем и использовать ее как для создания своего специализированного двойного варианта, так и в качестве макетов для других типов?

Например:

static __inline__ __device__ double foo_shfl_xor(double var, int laneMask, int width=warpSize)
{
    // Your double shuffle implementation
}

static __inline__ __device__ int foo_shfl_xor(int var, int laneMask, int width=warpSize)
{
    // For every non-double data type you use
    // Just call the original shuffle function
    return __shfl_xor(var, laneMask, width);
}

// Your code that uses shuffle
double d;
int a;
foo_shfl_xor(d, ...); // Calls your custom shuffle
foo_shfl_xor(a, ...); // Calls default shuffle
person Ashwin Nanjappa    schedule 12.04.2013
comment
Оно работает. Большое спасибо. Тем не менее, мне все еще интересно, есть ли более простое решение. - person Rainn; 12.04.2013