为线性插值设置 CUDA 2D“无符号字符”纹理

Question

我有一个代表二维数组的无符号字符线性数组。我想将它放入 CUDA 2D 纹理并对其执行（浮点）线性插值，即让纹理调用获取 4 个最近的无符号字符邻居，在内部将它们转换为浮点数，在它们之间进行插值，并返回结果浮点值。

我在设置纹理并将其绑定到纹理参考时遇到了一些困难。我已经阅读了 CUDA 参考手册和附录，但我没有任何运气。

以下是设置和绑定 1) 浮点纹理和 2) 无符号字符纹理的可运行代码。浮点代码运行得很好。但是，如果您取消注释底部的两个注释 unsigned char 行，则会引发“无效参数”错误。

#include <cstdio>
#include <cuda_runtime.h>

typedef unsigned char uchar;

// Define (global) texture references; must use "cudaReadModeNormalizedFloat"
// for ordinal textures
texture<float, cudaTextureType2D, cudaReadModeNormalizedFloat> texRefFloat;
texture<uchar, cudaTextureType2D, cudaReadModeNormalizedFloat> texRefUChar;

// Define size of (row major) textures
size_t const WIDTH  = 1000;
size_t const HEIGHT = 1000;
size_t const TOT_PIX = WIDTH*HEIGHT;

int main(void)
{
   // Set texel formats
   cudaChannelFormatDesc descFloat = cudaCreateChannelDesc<float>();
   cudaChannelFormatDesc descUChar = cudaCreateChannelDesc<uchar>();

   // Choose to perform texture 2D linear interpolation
   texRefFloat.filterMode = cudaFilterModeLinear;
   texRefUChar.filterMode = cudaFilterModeLinear;

   // Allocate texture device memory
   float * d_buffFloat; cudaMalloc(&d_buffFloat, sizeof(float)*TOT_PIX);
   uchar * d_buffUChar; cudaMalloc(&d_buffUChar, sizeof(uchar)*TOT_PIX);

   // Bind texture references to textures
   cudaError_t errFloat = cudaSuccess;
   cudaError_t errUChar = cudaSuccess;

   errFloat = cudaBindTexture2D(0, texRefFloat, d_buffFloat, descFloat,
                  WIDTH, HEIGHT, sizeof(float)*WIDTH);
   // Uncomment the following two lines for an error
   //errUChar = cudaBindTexture2D(0, texRefUChar, d_buffUChar, descUChar,
   //               WIDTH, HEIGHT, sizeof(uchar)*WIDTH);

   // Check for errors during binding
   if (errFloat != cudaSuccess)
   {
      printf("Error binding float texture reference: %s\n",
          cudaGetErrorString(errFloat));
      exit(-1);
   }

   if (errUChar != cudaSuccess)
   {
      printf("Error binding unsigned char texture reference: %s\n",
          cudaGetErrorString(errUChar));
      exit(-1);
   }

   return 0;
}

任何帮助/见解将不胜感激！

亚伦

Answer 1

纹理的每一行都必须正确对齐。如果将纹理绑定到普通数组（与 CUDA 数组相反），通常无法保证这一点。要将普通内存绑定到 2D 纹理，您需要使用 cudaMallocPitch() 分配内存。这将设置行间距，使其适合绑定到纹理。请注意，将 0 作为第一个参数传递给纹理绑定 API 调用并不是一个好习惯。此参数用于 CUDA 向应用程序返回偏移量。如果偏移量不为零，则需要在纹理访问期间将其添加到纹理坐标中。

这是一个简单的示例，展示了如何从元素为 unsigned char 的纹理中读取插值。

#include <stdlib.h>
#include <stdio.h>

// Macro to catch CUDA errors in CUDA runtime calls
#define CUDA_SAFE_CALL(call)                                          \
do {                                                                  \
    cudaError_t err = call;                                           \
    if (cudaSuccess != err) {                                         \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString(err) );       \
        exit(EXIT_FAILURE);                                           \
    }                                                                 \
} while (0)
// Macro to catch CUDA errors in kernel launches
#define CHECK_LAUNCH_ERROR()                                          \
do {                                                                  \
    /* Check synchronous errors, i.e. pre-launch */                   \
    cudaError_t err = cudaGetLastError();                             \
    if (cudaSuccess != err) {                                         \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString(err) );       \
        exit(EXIT_FAILURE);                                           \
    }                                                                 \
    /* Check asynchronous errors, i.e. kernel failed (ULF) */         \
    err = cudaThreadSynchronize();                                    \
    if (cudaSuccess != err) {                                         \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString( err) );      \
        exit(EXIT_FAILURE);                                           \
    }                                                                 \
} while (0)

texture<unsigned char, 2, cudaReadModeNormalizedFloat> tex;

__global__ void kernel (int m, int n, float shift_x, float shift_y) 
{
    float val;
    for (int row = 0; row < m; row++) {
        for (int col = 0; col < n; col++) {
            val = tex2D (tex, col+0.5f+shift_x, row+0.5f+shift_y);
            printf ("%.2f  ", val);
        }
        printf ("\n");
    }
}

int main (void)
{
    int m = 4; // height = #rows
    int n = 3; // width  = #columns
    size_t pitch, tex_ofs;
    unsigned char arr[4][3]= {{11,12,13},{21,22,23},{31,32,33},{251,252,253}};
    unsigned char *arr_d = 0;

    CUDA_SAFE_CALL(cudaMallocPitch((void**)&arr_d,&pitch,n*sizeof(*arr_d),m));
    CUDA_SAFE_CALL(cudaMemcpy2D(arr_d, pitch, arr, n*sizeof(arr[0][0]),
                                n*sizeof(arr[0][0]),m,cudaMemcpyHostToDevice));
    tex.normalized = false;
    tex.filterMode = cudaFilterModeLinear;
    CUDA_SAFE_CALL (cudaBindTexture2D (&tex_ofs, &tex, arr_d, &tex.channelDesc,
                                       n, m, pitch));
    if (tex_ofs !=0) {
        printf ("tex_ofs = %zu\n", tex_ofs);
        return EXIT_FAILURE;
    }
    printf ("reading array straight\n");
    kernel<<<1,1>>>(m, n, 0.0f, 0.0f);
    CHECK_LAUNCH_ERROR();
    CUDA_SAFE_CALL (cudaDeviceSynchronize());
    printf ("reading array shifted in x-direction\n");
    kernel<<<1,1>>>(m, n, 0.5f, 0.0f);
    CHECK_LAUNCH_ERROR();
    CUDA_SAFE_CALL (cudaDeviceSynchronize());
    printf ("reading array shifted in y-direction\n");
    kernel<<<1,1>>>(m, n, 0.0f, 0.5f);
    CUDA_SAFE_CALL (cudaDeviceSynchronize());
    CUDA_SAFE_CALL (cudaFree (arr_d));
    return EXIT_SUCCESS;
}

这个程序的输出如下：

reading array straight
0.04  0.05  0.05
0.08  0.09  0.09
0.12  0.13  0.13
0.98  0.99  0.99
reading array shifted in x-direction
0.05  0.05  0.05
0.08  0.09  0.09
0.12  0.13  0.13
0.99  0.99  0.99
reading array shifted in y-direction
0.06  0.07  0.07
0.10  0.11  0.11
0.55  0.56  0.56
0.98  0.99  0.99