За конкретные советы буду весьма благодарен.

Есть выражение, записанное в обратной польской записи (далее RPN), например такое:
(X + X) * (X + 2), что в RPN будет: X X + X 2 + *

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

То есть если считать, что размер множества выражений равен 1 (выражение единственное), а размер множества значений переменной равен 3, то мне надо получить 3 числа, каждое из которых будет результатом обсчёта выражения с подставленным конкретным значением переменной X.

В реальной жизни - выражений не одно, а много (десятки тысяч), возможных значений переменной X не 3, а тысячи.

Сейчас мой кернел работает по такой схеме:

__global__ void Kernel(int* pExpressions, int nNumExpressions,
float* pX, float* pOut, int nNumTestCases)
{
const int nThreadIdx = __mul24(blockDim.x, blockIdx.x) + threadIdx.x;
const int nThreadIdxInBlock = threadIdx.x;
const int nNumThreads = __mul24(blockDim.x, gridDim.x);

for (int nItem = nThreadIdx; ; nItem += nNumThreads)
{
int nExpression = nItem / nNumTestCases;
if (nExpression >= nNumExpressions)
break;

int nTestCase = nItem % nNumTestCases;
pOut[nExpression * nNumTestCases + nTestCase] = EvaluateTestCase(
pExpressions[nExpression], pX[nTestCase], nThreadIdxInBlock);
}
}

Внутри EvaluateTestCase заводится свой стек, живущий в shared mem, для каждого вызова EvaluateTestCase стек свой - посколько EvaluateTestCase вызывается для разных значений pX. Кол-во выражений и значений переменной кратно 32, каждый варп работает по одному выражению (divergence по идее нет).

Это всё щастье работает в два раза медленнее, чем на Athlon X2 4800+, считающем обоими ядрами. После соотнесения частот и прочих характеристик я прихожу к выводу, что эта задача будет считаться на GTX280 в 4-5 раз быстрее, чем на Core2Quad 3ГГц, что меня совсем не радует. По идее - задача исключительно параллелится, но пока что эффективность удручает.