Существует ли обходной путь для негибкого использования ссылок на текстуры в CUDA

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

texture<unsigned char, 2, cudaReadModeElementType> g_tex_a;
texture<unsigned char, 2, cudaReadModeElementType> g_tex_b;

__global__ void gpu_kernel_1()
{
    // long complicated kernel
    foo = tex2D(g_tex_a, x,y);
}

__global__ void gpu_kernel_2()
{
    // long complicated kernel
    bar = tex2D(g_tex_a, x,y);
}

main()
{
    gpu_kernel_1<<<grid, block>>>();
    gpu_kernel_2<<<grid, block>>>();
}

Нельзя ли каким-то образом передать ядру текстуру, которую оно должно считывать с помощью tex2D, например g_tex_a или g_tex_b? Похоже, что компилятор должен знать ссылку на текстуру во время компиляции. Я хотите иметь возможность повторно использовать код gpu_kernel_1 и gpu_kernel_2 для работы с различными текстурами.

Я действительно не понимаю, как текстурный ссылочный объект используется кодом хоста и устройства. На данный момент у меня есть ужасное решение репликации всего кода для каждого ядра с единственным изменением, заключающимся в том, что функция tex2D использует другую ссылку на текстуру, например gpu_kernel_1_with_tex_a (), gpu_kernel_1_with_tex_b().

Есть ли какое-либо другое решение? Спасибо.