logo标签列表

优化CUDA字节操作

c++optimizationcudabyteabsolute-value
3

我对Cuda还比较新，正在尝试编写一个核函数，计算查询向量和大型向量数据库之间的绝对差值之和。两者的元素都必须是8位无符号int类型。我以Nvidia的示例并行约简核心为基础来设计我的核函数，也阅读了这个线程。

我只能获得约5GB/s的速度，这与快速CPU相比并没有太大优势，也远远达不到DDR5 GT640理论带宽80GB/s的水平。

我的数据集由1024字节的查询向量和100,000 x 1024字节的数据库组成。

我有100,000个块，每个块有128个线程。如果每个块访问相同的1024字节查询向量，会导致性能更差吗？因为每个块都在访问同一个内存位置。

blockSize和shared memory都设置为128和128*sizeof(int)，其中128被定义为THREADS_PER_BLOCK。

template<UINT blockSize> __global__ void reduction_sum_abs( BYTE* query_vector, BYTE* db_vector, uint32_t* result )
{
    extern __shared__ UINT sum[]; 
    UINT db_linear_index = (blockIdx.y*gridDim.x) + blockIdx.x ; 
    UINT i = threadIdx.x; 

    sum[threadIdx.x] = 0; 

    int* p_q_int = reinterpret_cast<int*>(query_vector); 
    int* p_db_int = reinterpret_cast<int*>(db_vector); 

    while( i < VECTOR_SIZE/4 ) {

        /* memory transaction */
        int q_int = p_q_int[i]; 
        int db_int = p_db_int[db_linear_index*VECTOR_SIZE/4 + i]; 

        uchar4 a0 = *reinterpret_cast<uchar4*>(&q_int); 
        uchar4 b0 = *reinterpret_cast<uchar4*>(&db_int); 

        /* sum of absolute difference */ 
        sum[threadIdx.x] += abs( (int)a0.x - b0.x ); 
        sum[threadIdx.x] += abs( (int)a0.y - b0.y ); 
        sum[threadIdx.x] += abs( (int)a0.z - b0.z ); 
        sum[threadIdx.x] += abs( (int)a0.w - b0.w ); 

        i += THREADS_PER_BLOCK; 

    }

    __syncthreads(); 

    if ( blockSize >= 128 ) {
        if ( threadIdx.x < 64 ) { 
            sum[threadIdx.x] += sum[threadIdx.x + 64]; 
        }
    }

    /* reduce the final warp */
    if ( threadIdx.x < 32 ) {        
        if ( blockSize >= 64 ) { sum[threadIdx.x] += sum[threadIdx.x + 32]; } __syncthreads(); 

        if ( blockSize >= 32 ) { sum[threadIdx.x] += sum[threadIdx.x + 16]; } __syncthreads(); 

        if ( blockSize >= 16 ) { sum[threadIdx.x] += sum[threadIdx.x + 8 ]; } __syncthreads(); 

        if ( blockSize >= 8  ) { sum[threadIdx.x] += sum[threadIdx.x + 4 ]; } __syncthreads(); 

        if ( blockSize >= 4  ) { sum[threadIdx.x] += sum[threadIdx.x + 2 ]; } __syncthreads(); 

        if ( blockSize >= 2  ) { sum[threadIdx.x] += sum[threadIdx.x + 1 ]; } __syncthreads(); 

    }


    /* copy the sum back to global */
    if ( threadIdx.x == 0 ) {
        result[db_linear_index] = sum[0]; 
    }
}

我可以在运行内核时将实际的绝对差值计算代码注释掉，这样可以获得约4倍的带宽增加，显然这会导致错误的答案，但我认为至少有一部分时间是在那里花费的。
我的访问字节的方式是否会创建银行冲突？如果是，我能避免吗？
我使用reinterpret_cast的方法正确吗？
有没有更好的方法进行8位无符号计算？
还有哪些（我认为很多，因为我是个完全的新手）优化可以做？
谢谢。
编辑：
我的机器规格如下：
Windows XP 2002 SP3
Intel 6600 2.40GHz
2GB RAM
GT640 GDDR5 1GB
Visual C++ 2010 Express
你的 BYTE 是如何定义的？ - Michal Hosala
这只是一个uint8_t。 - user3678912
2
如果没有可以快速编译和运行的代码，很难对性能做出太多评价。在这段代码中，银行冲突不应该是一个问题。你是否通过NVIDIA分析器运行了代码？那通常会让你对正在发生的事情有很好的了解。此外，我认为你在第一步缩减和整个warp缩减之间缺少了一个__syncthreads()。 - Jez
2个回答
8
这类问题的最佳做法是提供完整的代码，使得其他人可以直接编译和运行，而无需添加或更改任何内容。一般来说，SO期望您提供完整的this。由于您的问题还涉及性能问题，因此您还应在完整的代码中包含实际的计时测量方法。
修复错误：
您的代码中至少有两个错误，其中@Jez已经指出了一个。在进行“部分约简”之后，还有另一个错误：
if ( blockSize >= 128 ) {
    if ( threadIdx.x < 64 ) { 
        sum[threadIdx.x] += sum[threadIdx.x + 64]; 
    }
}

我们需要在继续进行之前加上__syncthreads();。通过以上更改，我能够让您的内核生成可重复的结果，与我的 naive host 实现相匹配。另外，由于您有像这样不在线程块内评估的条件代码：
if ( threadIdx.x < 32 ) {

在条件代码块中使用 __syncthreads() 语句是非法的

  if ( blockSize >= 64 ) { sum[threadIdx.x] += sum[threadIdx.x + 32]; } __syncthreads();

（同样适用于执行相同操作的后续行）。因此，建议修复这个问题。我们可以采用几种方法来解决这个问题，其中之一是切换到使用volatile类型的指针来引用共享数据。由于我们现在正在一个warp内运行，volatile限定符强制编译器执行我们想要的操作：

volatile UINT *vsum = sum;
if ( threadIdx.x < 32 ) {        
    if ( blockSize >= 64 ) vsum[threadIdx.x] += vsum[threadIdx.x + 32];
    if ( blockSize >= 32 ) vsum[threadIdx.x] += vsum[threadIdx.x + 16]; 
    if ( blockSize >= 16 ) vsum[threadIdx.x] += vsum[threadIdx.x + 8 ];
    if ( blockSize >= 8  ) vsum[threadIdx.x] += vsum[threadIdx.x + 4 ];
    if ( blockSize >= 4  ) vsum[threadIdx.x] += vsum[threadIdx.x + 2 ]; 
    if ( blockSize >= 2  ) vsum[threadIdx.x] += vsum[threadIdx.x + 1 ];
}

CUDA 并行规约样例代码 和 相关pdf 可能是您的好复习资料。

时间/性能分析：

我恰好有一台GT 640，cc3.5设备。 当我在它上面运行bandwidthTest时，我观察到设备间传输大约为32GB/s。 这个数字代表了当设备内核访问设备内存时可实现带宽的合理近似上限。 此外，当我添加基于cudaEvent的计时并围绕您所展示的内容构建一个样例代码，并使用模拟数据，我观察到吞吐量约为16GB/s，而不是5GB/s。 因此，您的实际测量技术在这里将是有用的信息（事实上，完整的代码可能是分析您的内核计时和您的计时之间差异所需的）。

问题仍然存在，那么它可以改进吗？ （假设~32GB/s是近似的上限）。

你的问题：

我访问字节的方式会不会造成银行冲突？如果是，我可以避免冲突吗？

由于您的内核实际上将字节有效地加载为32位数量(uchar4)，并且每个线程正在加载相邻、连续的32位数量，因此我认为您的内核没有任何银行冲突访问问题。

我的reinterpret_cast用法正确吗？

是的,它看起来是正确的(下面是我的示例代码，经过上述修复，验证了您的内核生成的结果与朴素的主机函数实现相匹配)。

是否有更好的方法进行8位无符号计算？

有的，正如@njuffa所指出的那样，在这种情况下，SIMD intrinsic可以处理这个问题，事实证明，只需要单个指令(__vsadu4(), 请参见下面的示例代码)。

还有哪些优化方法可以使用？（我想会有很多，因为我是一个完全的新手）

  1. 使用@MichalHosala提出的cc3.0 warp-shuffle减少方法。

  2. 使用SIMD内置函数__vsadu4()来简化和改善对字节数量的处理，由@njuffa提出。

  3. 重新组织数据库向量数据以列主要存储。这使我们可以放弃普通的并行约简方法（即1中提到的方法），转而使用直接的for-loop读取内核，一个线程计算整个向量比较。在这种情况下（cc3.5 GT640），这使得我们的内核达到了设备的大约内存带宽。

这里是代码和样例运行，展示了3种实现方式：你的原始实现（加上上述“修复”，以获得正确的结果），opt1内核修改了你的实现，包括上述列表中的1和2项，而opt2内核则使用了上述列表中的2和3项。根据我的测量，你的内核达到了16GB/s，约为GT640带宽的一半，opt1内核运行速度约为24GB/s（增加大约相等于上述1和2项的部分），而经过数据重组的opt2内核运行速度接近全带宽（36GB/s）。
$ cat t574.cu
#include <stdio.h>
#include <stdlib.h>
#define THREADS_PER_BLOCK 128
#define VECTOR_SIZE 1024
#define NUM_DB_VEC 100000

typedef unsigned char BYTE;
typedef unsigned int UINT;
typedef unsigned int uint32_t;


template<UINT blockSize> __global__ void reduction_sum_abs( BYTE* query_vector, BYTE* db_vector, uint32_t* result )
{
    extern __shared__ UINT sum[];
    UINT db_linear_index = (blockIdx.y*gridDim.x) + blockIdx.x ;
    UINT i = threadIdx.x;

    sum[threadIdx.x] = 0;

    int* p_q_int = reinterpret_cast<int*>(query_vector);
    int* p_db_int = reinterpret_cast<int*>(db_vector);

    while( i < VECTOR_SIZE/4 ) {

        /* memory transaction */
        int q_int = p_q_int[i];
        int db_int = p_db_int[db_linear_index*VECTOR_SIZE/4 + i];

        uchar4 a0 = *reinterpret_cast<uchar4*>(&q_int);
        uchar4 b0 = *reinterpret_cast<uchar4*>(&db_int);

        /* sum of absolute difference */
        sum[threadIdx.x] += abs( (int)a0.x - b0.x );
        sum[threadIdx.x] += abs( (int)a0.y - b0.y );
        sum[threadIdx.x] += abs( (int)a0.z - b0.z );
        sum[threadIdx.x] += abs( (int)a0.w - b0.w );

        i += THREADS_PER_BLOCK;

    }

    __syncthreads();

    if ( blockSize >= 128 ) {
        if ( threadIdx.x < 64 ) {
            sum[threadIdx.x] += sum[threadIdx.x + 64];
        }
    }
    __syncthreads(); // **
    /* reduce the final warp */
    if ( threadIdx.x < 32 ) {
        if ( blockSize >= 64 ) { sum[threadIdx.x] += sum[threadIdx.x + 32]; } __syncthreads();

        if ( blockSize >= 32 ) { sum[threadIdx.x] += sum[threadIdx.x + 16]; } __syncthreads();

        if ( blockSize >= 16 ) { sum[threadIdx.x] += sum[threadIdx.x + 8 ]; } __syncthreads();

        if ( blockSize >= 8  ) { sum[threadIdx.x] += sum[threadIdx.x + 4 ]; } __syncthreads();

        if ( blockSize >= 4  ) { sum[threadIdx.x] += sum[threadIdx.x + 2 ]; } __syncthreads();

        if ( blockSize >= 2  ) { sum[threadIdx.x] += sum[threadIdx.x + 1 ]; } __syncthreads();

    }


    /* copy the sum back to global */
    if ( threadIdx.x == 0 ) {
        result[db_linear_index] = sum[0];
    }
}

__global__ void reduction_sum_abs_opt1( BYTE* query_vector, BYTE* db_vector, uint32_t* result )
{
  __shared__ UINT sum[THREADS_PER_BLOCK];
  UINT db_linear_index = (blockIdx.y*gridDim.x) + blockIdx.x ;
  UINT i = threadIdx.x;

  sum[threadIdx.x] = 0;

  UINT* p_q_int = reinterpret_cast<UINT*>(query_vector);
  UINT* p_db_int = reinterpret_cast<UINT*>(db_vector);

  while( i < VECTOR_SIZE/4 ) {

    /* memory transaction */
    UINT q_int = p_q_int[i];
    UINT db_int = p_db_int[db_linear_index*VECTOR_SIZE/4 + i];
    sum[threadIdx.x] += __vsadu4(q_int, db_int);

    i += THREADS_PER_BLOCK;

    }
  __syncthreads();
  // this reduction assumes THREADS_PER_BLOCK = 128
  if (threadIdx.x < 64) sum[threadIdx.x] += sum[threadIdx.x+64];
  __syncthreads();

  if ( threadIdx.x < 32 ) {
    unsigned localSum = sum[threadIdx.x] + sum[threadIdx.x + 32];
    for (int i = 16; i >= 1; i /= 2)
      localSum = localSum + __shfl_xor(localSum, i);
    if (threadIdx.x == 0) result[db_linear_index] = localSum;
    }
}

__global__ void reduction_sum_abs_opt2( BYTE* query_vector, UINT* db_vector_cm, uint32_t* result)
{
  __shared__ UINT qv[VECTOR_SIZE/4];
  if (threadIdx.x < VECTOR_SIZE/4) qv[threadIdx.x] = *(reinterpret_cast<UINT *>(query_vector) + threadIdx.x);
  __syncthreads();
  int idx = threadIdx.x + blockDim.x*blockIdx.x;
  while (idx < NUM_DB_VEC){
    UINT sum = 0;
    for (int i = 0; i < VECTOR_SIZE/4; i++)
      sum += __vsadu4(qv[i], db_vector_cm[(i*NUM_DB_VEC)+idx]);
    result[idx] = sum;
    idx += gridDim.x*blockDim.x;}
}

unsigned long compute_host_result(BYTE *qvec, BYTE *db_vec){

  unsigned long temp = 0;
  for (int i =0; i < NUM_DB_VEC; i++)
    for (int j = 0; j < VECTOR_SIZE; j++)
      temp += (unsigned long) abs((int)qvec[j] - (int)db_vec[(i*VECTOR_SIZE)+j]);
  return temp;
}

int main(){

  float et;
  cudaEvent_t start, stop;
  BYTE *h_qvec, *d_qvec, *h_db_vec, *d_db_vec;
  uint32_t *h_res, *d_res;
  h_qvec =   (BYTE *)malloc(VECTOR_SIZE*sizeof(BYTE));
  h_db_vec = (BYTE *)malloc(VECTOR_SIZE*NUM_DB_VEC*sizeof(BYTE));
  h_res = (uint32_t *)malloc(NUM_DB_VEC*sizeof(uint32_t));
  for (int i = 0; i < VECTOR_SIZE; i++){
    h_qvec[i] = rand()%256;
    for (int j = 0; j < NUM_DB_VEC; j++) h_db_vec[(j*VECTOR_SIZE)+i] = rand()%256;}
  cudaMalloc(&d_qvec, VECTOR_SIZE*sizeof(BYTE));
  cudaMalloc(&d_db_vec, VECTOR_SIZE*NUM_DB_VEC*sizeof(BYTE));
  cudaMalloc(&d_res, NUM_DB_VEC*sizeof(uint32_t));
  cudaMemcpy(d_qvec, h_qvec, VECTOR_SIZE*sizeof(BYTE), cudaMemcpyHostToDevice);
  cudaMemcpy(d_db_vec, h_db_vec, VECTOR_SIZE*NUM_DB_VEC*sizeof(BYTE), cudaMemcpyHostToDevice);
  cudaEventCreate(&start); cudaEventCreate(&stop);

// initial run

  cudaMemset(d_res, 0, NUM_DB_VEC*sizeof(uint32_t));
  cudaEventRecord(start);
  reduction_sum_abs<THREADS_PER_BLOCK><<<NUM_DB_VEC, THREADS_PER_BLOCK, THREADS_PER_BLOCK*sizeof(int)>>>(d_qvec, d_db_vec, d_res);
  cudaEventRecord(stop);
  cudaDeviceSynchronize();
  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&et, start, stop);
  cudaMemcpy(h_res, d_res, NUM_DB_VEC*sizeof(uint32_t), cudaMemcpyDeviceToHost);
  unsigned long h_result = 0;
  for (int i = 0; i < NUM_DB_VEC; i++) h_result += h_res[i];
  printf("1: et: %.2fms, bw: %.2fGB/s\n", et, (NUM_DB_VEC*VECTOR_SIZE)/(et*1000000));
  if (h_result == compute_host_result(h_qvec, h_db_vec)) printf("Success!\n");
  else printf("1: mismatch!\n");

// optimized kernel 1
  cudaMemset(d_res, 0, NUM_DB_VEC*sizeof(uint32_t));
  cudaEventRecord(start);
  reduction_sum_abs_opt1<<<NUM_DB_VEC, THREADS_PER_BLOCK>>>(d_qvec, d_db_vec, d_res);
  cudaEventRecord(stop);
  cudaDeviceSynchronize();
  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&et, start, stop);
  cudaMemcpy(h_res, d_res, NUM_DB_VEC*sizeof(uint32_t), cudaMemcpyDeviceToHost);
  h_result = 0;
  for (int i = 0; i < NUM_DB_VEC; i++) h_result += h_res[i];
  printf("2: et: %.2fms, bw: %.2fGB/s\n", et, (NUM_DB_VEC*VECTOR_SIZE)/(et*1000000));
  if(h_result == compute_host_result(h_qvec, h_db_vec)) printf("Success!\n");
  else printf("2: mismatch!\n");

// convert db_vec to column-major storage for optimized kernel 2

  UINT *h_db_vec_cm, *d_db_vec_cm;
  h_db_vec_cm = (UINT *)malloc(NUM_DB_VEC*(VECTOR_SIZE/4)*sizeof(UINT));
  cudaMalloc(&d_db_vec_cm, NUM_DB_VEC*(VECTOR_SIZE/4)*sizeof(UINT));
  for (int i = 0; i < NUM_DB_VEC; i++)
    for (int j = 0; j < VECTOR_SIZE/4; j++)
      h_db_vec_cm[(j*NUM_DB_VEC)+i] = *(reinterpret_cast<UINT *>(h_db_vec + (i*VECTOR_SIZE))+j);
  cudaMemcpy(d_db_vec_cm, h_db_vec_cm, NUM_DB_VEC*(VECTOR_SIZE/4)*sizeof(UINT), cudaMemcpyHostToDevice);
  cudaMemset(d_res, 0, NUM_DB_VEC*sizeof(uint32_t));
  cudaEventRecord(start);
  reduction_sum_abs_opt2<<<64, 512>>>(d_qvec, d_db_vec_cm, d_res);
  cudaEventRecord(stop);
  cudaDeviceSynchronize();
  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&et, start, stop);
  cudaMemcpy(h_res, d_res, NUM_DB_VEC*sizeof(uint32_t), cudaMemcpyDeviceToHost);
  h_result = 0;
  for (int i = 0; i < NUM_DB_VEC; i++) h_result += h_res[i];
  printf("3: et: %.2fms, bw: %.2fGB/s\n", et, (NUM_DB_VEC*VECTOR_SIZE)/(et*1000000));
  if(h_result == compute_host_result(h_qvec, h_db_vec)) printf("Success!\n");
  else printf("3: mismatch!\n");

  return 0;
}

$ nvcc -O3 -arch=sm_35 -o t574 t574.cu
$ ./run35 t574
1: et: 6.34ms, bw: 16.14GB/s
Success!
2: et: 4.16ms, bw: 24.61GB/s
Success!
3: et: 2.83ms, bw: 36.19GB/s
Success!
$

一些注意事项：

  1. 上述代码，特别是您的内核，必须编译为cc3.0或更高版本，以我设置测试用例的方式。这是因为我在一个单一的1D网格中创建了100,000个块，因此我们不能直接在cc2.0设备上运行此代码。
  2. 对于opt2内核，在不同的设备上运行时，可以通过修改网格和块参数进行微调。我将它们设置为64和512，这些值不应该是关键因素（尽管块应该是VECTOR_SIZE/4线程或更多），因为算法使用网格遍历循环来覆盖整个向量集。GT640只有2个SM，因此在这种情况下，64个线程块足以让设备保持繁忙（甚至32个也可以）。您可能需要修改这些参数以在更大的设备上获得最佳性能。
谢谢，这很棒，除了我正在计算查询向量与数据库中每个向量之间的各个绝对差之和，而不是它们所有的总和，这就是为什么我使用了100,000个块（实际上是一个10 * 10000网格）。除非我对您的代码解释有误。我通过将处理的字节数（1024 * 100,000 * 2）除以内核运行时由nvidia分析器测量来衡量性能。 - user3678912
有人偏爱使用volatile线程块规约方法而不是__shfl_xor吗？无论如何，代码必须编译为cc3.0。最近我在写论文时，似乎shfl比volatile好一点，而且更安全两位。所以只是好奇... - Michal Hosala
是的，我没有看到不能切换到列主序的任何理由，只要结果是正确答案，我愿意接受任何建议。为什么您不同意我的算法？此外，我使用cudaEvent得到的时间与nvidia分析工具中的时间不同。 - user3678912
我肯定不会接近6.4毫秒。主机硬件规格是否会影响GPU性能？不幸的是，我只能使用一台相当老旧的Windows XP机器。使用带宽利用工具，我记得我也获得了大约32GB / sec的结果，而Nvidia SDK缩减示例运行速度约为25GB / sec。 - user3678912
我已经修改了我的答案，以反映您需要单独计算所有向量差异的事实。提供了两个优化内核的示例，第二个基于您允许重新组织数据库向量数据。关于您的算术，这有点争议，但我们在因子2上存在分歧。您似乎是基于数据库向量加上查询向量来计算带宽的。至少我的第二个优化内核只需要加载几次查询向量。不幸的是，这并不能解释我们带宽测量结果的差异。 - Robert Crovella
显示剩余4条评论
1
一件事情立刻引起了我的注意：
if ( blockSize >= 128 ) {
    if ( threadIdx.x < 64 ) { 
        sum[threadIdx.x] += sum[threadIdx.x + 64]; 
    }
}

第一个条件在任何情况下都成立，而第二个条件仅在前两个环节中成立。因此，您可以通过交换它们的顺序来获益：

if ( threadIdx.x < 64 ) {
    if ( blockSize >= 128 ) { 
        sum[threadIdx.x] += sum[threadIdx.x + 64]; 
    }
}

这将允许除前两个之外的所有warp更快地完成执行。
接下来，您可以使用__shfl_xor指令显着加速warp级别的约简。
/* reduce the final warp */
if ( threadIdx.x < 32 ) {
  auto localSum = sum[threadIdx.x] + sum[threadIdx.x + 32]); 
  for (auto i = 16; i >= 1; i /= 2)
  {
      localSum = localSum + __shfl_xor(localSum, i);
  }

  if (threadIdx.x == 0) result[db_linear_index] = localSum;
}

我并不是说这就是所有问题，你的代码可能还有其他问题，但是这些问题我很容易就能发现。我甚至还没有使用我的解决方案测试性能，但我相信它应该会有所改进。
编辑：另外，你似乎无必要地将数据写入共享内存四次。
/* sum of absolute difference */ 
sum[threadIdx.x] += abs( (int)a0.x - b0.x ); 
sum[threadIdx.x] += abs( (int)a0.y - b0.y ); 
sum[threadIdx.x] += abs( (int)a0.z - b0.z ); 
sum[threadIdx.x] += abs( (int)a0.w - b0.w );

为什么不简单地这样做呢？
    /* sum of absolute difference */ 
sum[threadIdx.x] += abs( (int)a0.x - b0.x )
    + abs( (int)a0.y - b0.y )
    + abs( (int)a0.z - b0.z ); 
    + abs( (int)a0.w - b0.w );
我现在看到我的答案远远没有涵盖所有提供的代码问题，但我现在无法正确地写下来，可能明天会回来。 - Michal Hosala
谢谢，编译器如何解释在单个语句中放置所有abs（）的内容？此外，根据https://developer.nvidia.com/cuda-gpus，GT640 GDDR5是一个计算3.5卡。 - user3678912
+= 被定义为 T& T::operator +=(const T2& b);，因此我认为它首先将 += 右侧的所有参数相加以创建 b，然后将其作为参数传递给 +=，从而使其结果成为对 sum 的单个赋值操作，而不是四个。啊，我错过了 GT640 的 GDDR3 和 GDDR5 版本的事实... 我删除了与加速小于 3.0 GPU 的 warp reduction 相关的部分。 - Michal Hosala
1
你可能需要研究使用SIMD设备函数（内部函数），直接操作打包字节，例如__vabsdiffu4()和__vsadu4()（用于无符号字节）或__vabsdiffs4()和__vsads4()（用于有符号字节）。有关详细信息，请参阅CUDA数学API文档：http://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__SIMD.html#group__CUDA__MATH__INTRINSIC__SIMD - njuffa
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接