Передача массивов / указателей в качестве параметров шаблона

Question

Я пытаюсь создать шаблонную функцию следующего вида:

template <bool isHorizontal, float* kernel>
__global__ void smoothFilterColumns(const TwImageCUDA_Device* source, TwImageCUDA_Device* 
destination)
{
// code...
}

(Не беспокойтесь о __global__; это реализовано в CUDA.)

К сожалению, это не позволит мне создавать экземпляры этой функции следующим образом:

float ptrKernel[] = {1, 2, 1};
smoothFilterColumns<true, ptrKernel>(dxBuffer->cuda_image, dxOutput->cuda_image);

Я пробовал все виды float* с и float[] с, с модификатором const и без него. Можно ли вообще создать шаблон такого рода?

Заранее спасибо.

NB. Ядро передается как параметр шаблона, а не как параметр обычной функции, потому что это позволяет мне создавать более эффективный код в CUDA путем развертывания циклов.

Обновление Указатели на поплавки работают как параметры шаблона со стандартным C ++, но, очевидно, нет способа заставить их работать с функциями устройства CUDA, так как они ожидают указателей на адреса устройств, и их нельзя определить внешне. Если кто-то заставил это работать, пожалуйста, дайте мне знать.

Answer 1

Я сомневаюсь, что ты заставишь это работать. Как отмечают другие, стандарт C ++ говорит, что любой объект или функция, передаваемые в качестве параметра шаблона, должны иметь внешнюю связь (поэтому она не определена в области действия текущей единицы перевода). Проблема заключается в том, что CUDA в настоящее время вообще не поддерживает внешнюю связь - каждый символ, используемый в коде устройства, должен иметь внутреннюю связь (т. Е. Определенную в одной и той же единице перевода). Основная причина этого ограничения заключается в том, что в CUDA в настоящее время нет компоновщика для кода устройства.

Answer 2

Я предполагаю, что переменная ptrKernel, которую вы передаете в качестве аргумента шаблона, является локальной переменной. В любом случае, существует ограничение на то, что вы можете передать в качестве аргумента нетипичного шаблона. Согласно стандарту C ++ (14.3.2) допускается следующее:

• целочисленное константное выражение целого типа или типа перечисления
• имя параметра шаблона нетипичного типа
• имя объекта или функции с внешней связью
• адрес объекта или функции с внешней связью
• указатель на элемент

Убедитесь, что переменная ptrKernel соответствует этим требованиям (опять же, я предполагаю, что она не является переменной с внешней связью, то есть глобальным или статическим членом класса).

Answer 3

Пожалуйста, убедитесь, что ptrKernel имеет внешнюю связь.

// static float ptrKernel[] = { ... };
// ^ won't work.

// const float ptrKernel[] = { ... };
// ^ won't work.

float ptrKernel[] = { ... };
// ^ ok.

void func() {
   // float ptrKernel[] = { ... };
   // ^ won't work (not global variable).
   ...
}

---

Это ограничение нетипичного шаблона, как описано в § [temp.arg.nontype] / 1:

A шаблон-аргумент для нетипового, не шаблонного шаблон-параметра должен быть одним из:

• целое число константа-выражение целочисленного или перечислительного типа; или
• имя нетипичного шаблона-параметра ; или
• имя объекта или функции с внешней связью , включая шаблоны функций и функции template-id , но исключая нестатические члены класса, выражается как id -expression ; или
• адрес объекта или функции с внешней связью , включая шаблоны функций и функции template-id , но исключая нестатические члены класса, выраженные как & id-выражение , где & является необязательным, если имя относится к функции или массиву; или
• указатель на член, выраженный как описано в 5.3.1.

Answer 4

не сработает.Вы пытаетесь передать указатель CPU-RAM на ядро ​​GPU-RAM.

Вы можете сделать это по-разному: 1) встроить все постоянные значения с использованием нескольких шаблонов в зависимости от длины ядра;или иначе вы создаете класс функтора, который обрабатывает детали преобразования, которое вы хотите применить:

Вот рабочий пример, чтобы вы поняли.Не забудьте спецификатор device .

// with 3 int
template<int amount, int k0,int k1, int k2>
__global__ void apply_kernel(const float *input, float *output){


}

// with four int
template<int amount, int k0,int k1, int k2, int k3>
__global__ void apply_kernel(const float *input, float *output){


}

// with five int 
template<int amount, int k0,int k1, int k2, int k3, int k4>
__global__ void apply_kernel(const float *input, float *output){


}

class KernelOperator {
public:
      __host__ __device__ KernelOperator() {
      }
      __host__ __device__ int operator*(int value){
            return value * 2;
      }
};


// with KernelOperator
template<class T>
__global__ void apply_kernel(const float *input, float *output){
           T value;

}

int main(){
    apply_kernel<0, 1,2,1><<<10, 20>>>(NULL,NULL);

    apply_kernel< KernelOperator ><<<10, 20>>>(NULL,NULL);
}

Answer 5

Будет ли это работать в CUDA?

template <bool isHorizontal, class Kernel>
__global__ void smoothFilterColumns(
    const TwImageCUDA_Device* source, TwImageCUDA_Device* destination)
{
    const float *kernel = Kernel::ptr();
    // code...
}

struct Kernel_1_2_1
{
    static const float *ptr()
    {
        static const float kernel[] = {1, 2, 1};
        return kernel;
    }
}

smoothFilterColumns<true, Kernel_1_2_1>(
    dxBuffer->cuda_image, dxOutput->cuda_image);

Возможно, вы сможете сделать ядро ​​членом данных struct. И вы можете захотеть добавить механизм для передачи размера ядра.