__brev

产品支持情况

• Ascend 950PR/Ascend 950DT：支持

• Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持

• Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持

• Atlas 200I/500 A2 推理产品：不支持

• Atlas 推理系列产品AI Core：不支持
• Atlas 推理系列产品Vector Core：不支持

• Atlas 训练系列产品：不支持

功能说明

将输入数据的位序反转，返回反转后的值。

函数原型

C++

unsigned int __brev(unsigned int x)

C++

unsigned long long __brev(unsigned long long x)

C++

unsigned long __brev(unsigned long x)

参数说明

表1 参数说明

参数名	输入/输出	描述
x	输入	源操作数。

返回值说明

输入数据的位反转值。当输入的类型为uint32_t时，返回值的第n位对应输入数据的第31-n位；当输入的类型为uint64_t时，返回值的第n位对应输入数据的第63-n位。

• 当x为0时，类型为uint32_t时，返回值为0。
• 当x为0时，类型为uint64_t时，返回值为0。
• 当x为1时，类型为uint32_t时，返回值为2147483648。
• 当x为1时，类型为uint64_t时，返回值为9223372036854775808。

约束说明

无

需要包含的头文件

使用该接口需要包含"simt_api/device_functions.h"头文件。

C++

#include "simt_api/device_functions.h"

调用示例

• SIMT编程场景：

  C++

  __global__ __launch_bounds__(1024) void KernelBrev(unsigned int* dst, unsigned int* x)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = __brev(x[idx]);
  }
  

  C++

  __global__ __launch_bounds__(1024) void KernelBrev(unsigned long long* dst, unsigned long long* x)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = __brev(x[idx]);
  }
  

• SIMD与SIMT混合编程场景：

  C++

  __simt_vf__ __launch_bounds__(1024) inline void KernelBrev(__gm__ unsigned int* dst, __gm__ unsigned int* x)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = __brev(x[idx]);
  }
  

  C++

  __simt_vf__ __launch_bounds__(1024) inline void KernelBrev(__gm__ unsigned long long* dst, __gm__ unsigned long long* x)
  {
      int idx = threadIdx.x + blockIdx.x * blockDim.x;
      dst[idx] = __brev(x[idx]);
  }