Блог

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

Цель
Моя цель состоит в том, чтобы заставить мои указатели на функции выбирать, какую функцию применять к двум объектам в соответствии с внутренним свойством последнего. Короче говоря, я хочу эмулировать шаблоны C с помощью ядра CUDA — фактически не используя аргументы или switch предложения шаблона, а вместо этого указатели на функции и class / struct members .

Подход

• Определите мои пользовательские объекты struct customObj с помощью одного свойства ( int type ), которое будет эмулировать аргументы шаблона.
• Определите набор фиктивных функций ( Sum() , Subtract() , и т.д.) на выбор.
• Сохраните список функций для применения ( functionsList ) и соответствующих type элементов для поиска ( first_types , second_types ) в __constant__ памяти, чтобы функция functionsList[i](obj1,obj2) применялась к объектам с obj1.type == first_types[i] помощью и obj2.type == second_types[i] .

Рабочий код
Следующий код был скомпилирован для Linux x86_64 с CUDA 5.0 на графическом процессоре с вычислительными возможностями 3.0 (GeForce GTX 670) и работает.

 #include <stdio.h>
#include <iostream>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %dn", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

struct customObj
{
  int type;
  double d;
  // Constructors
  __device__ __host__ customObj() {}
  __device__ __host__ customObj(const intamp; _type, const doubleamp; _d) : type(_type), d(_d) {}
};

typedef void (*function_t)(customObjamp;, customObjamp;);
// Define a bunch of functions
__host__ __device__ void Sum(customObjamp; obj1, customObjamp; obj2) {printf("Sum chosen! d1   d2 = %fn", obj1.d   obj2.d);}
__host__ __device__ void Subtract(customObjamp; obj1, customObjamp; obj2) {printf("Subtract chosen! d1 - d2 = %fn", obj1.d - obj2.d);}
__host__ __device__ void Multiply(customObjamp; obj1, customObjamp; obj2) {printf("Multiply chosen! d1 * d2 = %fn", obj1.d * obj2.d);}

#define ARRAYLENGTH 3
__constant__ int first_type[ARRAYLENGTH] = {1, 2, 3};
__constant__ int second_type[ARRAYLENGTH] = {1, 1, 2};
__constant__ function_t functionsList[ARRAYLENGTH] = {Sum, Sum, Subtract};

// Kernel to loop through functions list
__global__ void choosefunction(customObj obj1, customObj obj2) {
   int i = 0;
   function_t f = NULL;
   do {
     if ((obj1.type == first_type[i]) amp;amp; (obj2.type == second_type[i])) {
       f = functionsList[i];
       break;
    }
    i  ;
  } while (i < ARRAYLENGTH);
  if (f == NULL) printf("No possible interaction!n");
  else f(obj1,obj2);
}

int main() {
  customObj obj1(1, 5.2), obj2(1, 2.6);
  choosefunction<<<1,1>>>(obj1, obj2);
  gpuErrchk(cudaPeekAtLastError());
  gpuErrchk(cudaDeviceSynchronize()); 

  return 0;
}
 

Проблема
Проблема, которую я обнаружил, заключается в том, что, как только я заменяю тип данных члена int type и связанных переменных и функций ( __constant__ int first_types[...] и так далее)… код компилируется, но перестает работать!

• Если я изменяю тип данных с int на char или int8_t , средство проверки памяти выдает error 4 ответ на мой вызов to cudaDeviceSynchronize() .
• Если я изменю тип данных на unsigned short int , я получу переполнение аппаратного стека.

Итак, есть ли у кого-нибудь подобные проблемы при работе с __constant__ памятью? Я действительно понятия не имею, что происходит. Насколько я знаю, char и int8_t являются встроенными типами длиной 1 байт, в то время как размер int составляет 4 байта, так что, возможно, речь идет о выравнивании данных, но я просто предполагаю здесь. Кроме того, предполагается, что CUDA поддерживает указатели на функции на графическом процессоре начиная с compute capability 2.0. Существуют ли какие-либо особые ограничения для указателей на функции в __constant__ памяти, которые мне не хватает?

Я смог воспроизвести проблему (ошибка 4, неопределенный сбой запуска) на CUDA 5.0 на 64-битном RHEL 5.5, но не на CUDA 6.0.

Пожалуйста, обновите / обновите до CUDA 6.