CUDAのグローバルメモリとテクスチャメモリの速度を比較してみました。
対象とした処理は「平均値フィルタ」です。

コードは以下の通り。
// 平均値フィルタ(グローバルメモリ)
__global__ void meanFilterKernel(unsigned char *out, const unsigned char *in,
const int width, const int filterRadius, const float fscale)
{
for (int i = threadIdx.x; i < width; i += blockDim.x) {
int sum = 0;
for (int fy = -filterRadius; fy <= filterRadius; ++fy) {
int y = blockIdx.x + fy;
for (int fx = -filterRadius; fx <= filterRadius; ++fx) {
sum += in[y * width + i + fx];
}
}
out[blockIdx.x * width + i] = sum * fscale;
}
}

// 平均値フィルタ(テクスチャメモリ)
texture<unsigned char, 2> tex;
__global__ void meanFilterKernelTexture(unsigned char *out, const unsigned char *in,
const int width, const int filterRadius, const float fscale)
{
for (int i = threadIdx.x; i < width; i += blockDim.x) {
int sum = 0;
for (int y = -filterRadius; y <= filterRadius; ++y) {
for (int x = -filterRadius; x <= filterRadius; ++x) {
sum += tex2D(tex, i + x, blockIdx.x + y);
}
}
out[blockIdx.x * width + i] = sum * fscale;
}
}

void main()
{
cv::Mat img = cv::imread("Room_16000.bmp", 0);
int pixelNum = img.cols * img.rows;
int dataSize = pixelNum * sizeof(unsigned char);

unsigned char *hOut, *dIn, *dOut;
hOut = new unsigned char[pixelNum];
checkCudaErrors(cudaMalloc((unsigned char**)&dIn, dataSize));
checkCudaErrors(cudaMalloc((unsigned char**)&dOut, dataSize));
checkCudaErrors(cudaMemcpy(dIn, img.data, dataSize, cudaMemcpyHostToDevice));

int filterRadius = 5;
int filterDiameter = filterRadius * 2 + 1;
float fscale = 1.0f / filterDiameter / filterDiameter;
dim3 grid(img.rows);
dim3 block(1024);

LARGE_INTEGER freq;
QueryPerformanceFrequency(&freq);

// CPU
LARGE_INTEGER start, end;
QueryPerformanceCounter(&start);
for (int h = filterRadius; h < img.rows - filterRadius; ++h) {
for (int w = filterRadius; w < img.cols - filterRadius; ++w) {
int sum = 0;
for (int fy = -filterRadius; fy <= filterRadius; ++fy) {
int y = h + fy;
for (int fx = -filterRadius; fx <= filterRadius; ++fx) {
sum += img.data[y * img.cols + w + x];
}
}
hOut[h * img.cols + w] = sum * fscale;
}
}
QueryPerformanceCounter(&end);
 printf("CPU:%lf[ms]\n", (double)(end.QuadPart - start.QuadPart)
/ freq.QuadPart * 1000.0);

// GPU(グローバルメモリ)
QueryPerformanceCounter(&start);
meanFilterKernel <<<grid, block>>>(dOut, dIn, img.cols,
filterRadius, fscale);
checkCudaErrors(cudaDeviceSynchronize());
QueryPerformanceCounter(&end);
printf("GPU(global memory):%lf[ms]\n", (double)(end.QuadPart - start.QuadPart)
/ freq.QuadPart * 1000.0);

// GPU(テクスチャメモリ)
cudaArray *cuArray;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>();
checkCudaErrors(cudaMallocArray(&cuArray, &desc, img.cols, img.rows));
checkCudaErrors(cudaMemcpyToArray(cuArray, 0, 0, img.data,
dataSize, cudaMemcpyHostToDevice));
checkCudaErrors(cudaBindTextureToArray(tex, cuArray));

QueryPerformanceCounter(&start);
meanFilterKernelTexture <<<grid, block>>>(dOut, dIn, img.cols,
filterRadius, fscale);
checkCudaErrors(cudaDeviceSynchronize());
QueryPerformanceCounter(&end);
printf("GPU(texture memory):%lf[ms]\n", (double)(end.QuadPart - start.QuadPart
/ freq.QuadPart * 1000.0);
}

実験は以下の環境で行いました。
CPU:i7-2600 3.4GHz
メモリ:12GB
GPU:GTX970

結果は以下のようになりました(単位は[ms]、表の左端はフィルタサイズ)。

・画像解像度2000x1000
02000

・画像解像度16000x8000
16000

想定に反してテクスチャメモリの方が遅くなってしまいました。
使い方が間違っているのかなんなのか。

CPUGPUの処理時間比は50倍程度でしょうか。
シェアードメモリを使うような処理はもっと速くなるのかもしれないですが。

☆追記
試行回数を増やして、その平均処理時間をとっても、同様の傾向でした。