некоторые математические операции в CUDA

Question

У меня есть 2D-матрица, содержащая 0,1 и 2. Я пишу ядро ​​cuda, где количество потоков равно размеру матрицы, и каждый поток будет работать на каждом элементе матрицы. Теперь мне нужны были математические операции, которые могли бы содержать 0 и 1 как есть, но преобразовали бы 2 в 1. Это математическая операция без каких-либо if-else, которая сделала бы следующее преобразование: 0 -> 0; 1 -> 1; 2 -> 1. Существует ли какой-либо возможный способ использования математических операторов, которые могли бы выполнить вышеупомянутое преобразование. Любая помощь будет чрезвычайно оценена. Спасибо.

Answer 1

Это не вопрос cuda.

int A; 
// set A to 0, 1, or 2 
int a = (A + (A>>1)) & 1; 
// a is now 0 if A is 0, or 1 if A is 1 or 2 

или как макрос:

#define fix01(x) ((x+(x>>1))&1) 

int a = fix01(A); 

Это также похоже на работу:

#define fix01(x) ((x&&1)&1) 

Я не знаю, если использование логического значения и оператора (&&) подходит для вашего определение «математических операций».

Answer 2

Поскольку речь шла о «математических» функциях, которые я предлагаю следующий 2-й полином порядка:

int f(int x) { return ((3-x)*x)/2; } 

Но если вы хотите избежать ветвлений, чтобы максимизировать скорость: Существует мин инструкция, так как PTX ISA 1.0. (См табл. 36 в 3.1 руководстве PTX ISA.) Таким образом, следующий CUDA код

__global__ void test(int *x, int *y) 
{ 
    *y = *x <= 1 ? *x : 1; 
} 

компилирует к следующему PTX ассемблере в моем тесте (только называется NVCC от CUDA 5 без каких-либо опций арочных)

code for sm_10 
      Function : _Z4testPiS_ 
    /*0000*/  /*0x1000c8010423c780*/  MOV R0, g [0x4]; 
    /*0008*/  /*0xd00e000580c00780*/  GLD.U32 R1, global14 [R0]; 
    /*0010*/  /*0x1000cc010423c780*/  MOV R0, g [0x6]; 
    /*0018*/  /*0x30800205ac400780*/  IMIN.S32 R1, R1, c [0x1] [0x0]; 
    /*0020*/  /*0xd00e0005a0c00781*/  GST.U32 global14 [R0], R1; 

Таким образом, реализация min() с использованием условного?: Фактически компилируется в одну инструкцию IMIN.S32 PTX без какого-либо разветвления. Поэтому я бы рекомендовал это для каких-либо реальных приложений:

int f(int x) { return x <= 1 ? x : 1; } 

Но вернемся к вопросу об использовании только неветвящейся операции:

Другая форма получения этого результата в С помощью двух не операторы:

int f(int x) { return !!x; } 

Или просто сравнить с нулем:

int f(int x) { return x != 0; } 

(результаты! nd! = гарантируется 0 или 1, ср. 6.5.3.3 Пар. 5 и п. 6.5.9 Пар. 3 стандарта C99, ISO/IEC 9899: 1999. Afair эта гарантия также содержится в CUDA.)