Quick Start

This section briefly describes how to configure the heterogeneous programming environment and use the compiler, for you to quickly verify environment information and get familiar with the use of the BiSheng Compiler.

Installation and Environment Configuration

The BiSheng Compiler is released with the CANN package. After the CANN package is installed, the BiSheng Compiler is stored in ${INSTALL_DIR}/compiler/ccec_compiler.

Replace ${INSTALL_DIR} with the actual CANN component directory. If the Ascend-CANN-Toolkit package is installed as the root user, the CANN component directory is /usr/local/Ascend/ascend-toolkit/latest.

Before programming, configure the environment variables related to the binary program of the BiSheng Compiler in either of the following methods:

Method 1: Configure the CANN environment variable.
The CANN package provides a process-level environment variable setting script to automatically set the environment variable. In the following example command, the default installation paths are under the root or non-root user. Replace them with actual installation paths.
```
# Install Toolkit as the root user.
. /usr/local/Ascend/ascend-toolkit/set_env.sh 
# Install Toolkit as a non-root user.
. ${HOME}/Ascend/ascend-toolkit/set_env.sh 
```

Method 2: Configure the PATH environment variable.

# Obtain the BiSheng Compiler installation directory in the CANN package. For example:
$ export PATH=${INSTALL_DIR}/compiler/ccec_compiler/bin/:$PATH

Example of Compiling a Heterogeneous Program

This example demonstrates a heterogeneous program. It starts the kernel function of four blocks with each block writing its own data. The host uses the ACL Runtime APIs for runtime management.

// File name: QuickStartDemo.cce
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>

#ifdef ASCENDC_CPU_DEBUG
#define __aicore__
#else
#define __aicore__ [aicore]
#endif

#define BLOCKS 4
#define CACHELINE_SZ 64

// Define a kernel
__global__ __aicore__ void foo(__gm__ uint8_t *Out, int Stride) {
    Out[block_idx * Stride] = block_idx;
}

int main(int argc, char *argv[]) {
    aclInit(nullptr);
    aclrtSetDevice(0);
    aclrtStream stream;
    aclrtCreateStream(&stream);

    uint8_t ExpectedValue[] = {0, 1, 2, 3};
    uint8_t *OutputValue = nullptr;
    aclrtMalloc((void **)&OutputValue, BLOCKS, ACL_MEM_MALLOC_HUGE_FIRST);

    uint8_t InitValue[BLOCKS] = {0};
    aclrtMemcpyAsync((void *)OutputValue, sizeof(InitValue), InitValue,
                     sizeof(InitValue), ACL_MEMCPY_HOST_TO_DEVICE, stream);
    aclrtSynchronizeStream(stream);

    // Invoke a kernel
    foo<<<BLOCKS, nullptr, stream>>>(OutputValue, CACHELINE_SZ);

    uint8_t *OutHost = nullptr;
    aclrtMallocHost((void **)&OutHost, BLOCKS * CACHELINE_SZ);
    aclrtMemcpyAsync(OutHost, BLOCKS * CACHELINE_SZ, OutputValue,
                     BLOCKS * CACHELINE_SZ, ACL_MEMCPY_DEVICE_TO_HOST, stream);
    aclrtSynchronizeStream(stream);

    for (int I = 0; I < sizeof(ExpectedValue) / sizeof(uint8_t); I++) {
        printf("i%d\t Expect: 0x%04x\t\t\t\tResult: 0x%04x\n", I, ExpectedValue[I],
                OutHost[I * CACHELINE_SZ]);
    }

    aclrtFreeHost(OutHost);
    aclrtFree(OutputValue);

    aclrtDestroyStream(stream);
    aclrtResetDevice(0);
    aclFinalize();
    return 0;
}

The compilation commands are as follows. For details, see Heterogeneous Compilation.

# Runtime path in the CANN package
export RT_INC=${INSTALL_DIR}/runtime/include
export RT_LIB=${INSTALL_DIR}/runtime/lib64

# Function: Compile the host and device code together to generate executable files, which must be linked to libascendcl.so and libruntime.so.
# Compilation options --cce-soc-version and --cce-soc-core-type are used to compile the vector core program on AscendXXXYY.
$bisheng -O2 --cce-soc-version=AscendXXXYY --cce-soc-core-type=VecCore  -I$RT_INC -L$RT_LIB -lascendcl -lruntime QuickStartDemo.cce  -o QuickStartDemo

The execution result is as follows:

$ ./QuickStartDemo
i0       Expect: 0x0000                         Result: 0x0000
i1       Expect: 0x0001                         Result: 0x0001
i2       Expect: 0x0002                         Result: 0x0002
i3       Expect: 0x0003                         Result: 0x0003

Example of Compiling an Ascend C Operator

The following code implements a vector operator Add by using Ascend C.

// File name: QuickStartDemoVecAdd.cce
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>

#ifdef ASCENDC_CPU_DEBUG
#define __aicore__
#else
#define __aicore__ [aicore]
#endif

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // seperate to 2 parts, due to double buffer

// ---------- Device side code ------------------------------
#include "kernel_operator.h"
__global__ __aicore__ void VecAdd(__gm__ float *x, __gm__ float *y, __gm__ float *z) {
    using namespace AscendC;

    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
    GlobalTensor<float> xGm;
    GlobalTensor<float> yGm;
    GlobalTensor<float> zGm;

    xGm.SetGlobalBuffer(x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    yGm.SetGlobalBuffer(y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    zGm.SetGlobalBuffer(z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(float));
    pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(float));
    pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(float));

    LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
    LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
    LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
    uint32_t loopCount = TILE_NUM * BUFFER_NUM;
    for (uint32_t i = 0; i < loopCount; i++) {
      DataCopy(xLocal, xGm[i * TILE_LENGTH], TILE_LENGTH);
      DataCopy(yLocal, yGm[i * TILE_LENGTH], TILE_LENGTH);
      inQueueX.EnQue(xLocal);
      inQueueY.EnQue(yLocal);

      xLocal = inQueueX.DeQue<float>();
      yLocal = inQueueY.DeQue<float>();
      Add(zLocal, xLocal, yLocal, TILE_LENGTH);
      outQueueZ.EnQue<float>(zLocal);

      zLocal = outQueueZ.DeQue<float>();
      DataCopy(zGm[i * TILE_LENGTH], zLocal, TILE_LENGTH);
    }
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
    outQueueZ.FreeTensor(zLocal);
}

int main(int argc, char *argv[]) {
  size_t inputByteSize = TOTAL_LENGTH * sizeof(float);
  size_t outputByteSize = TOTAL_LENGTH * sizeof(float);
  uint32_t blockDim = 8;

  // Initialize AscendCL.
  aclInit(nullptr);
  // Allocate runtime resources.
  aclrtContext context;
  int32_t deviceId = 0;
  aclrtSetDevice(deviceId);
  aclrtCreateContext(&context, deviceId);
  aclrtStream stream = nullptr;
  aclrtCreateStream(&stream);

  // Allocate the host memory.
  float *xHost, *yHost, *zHost;
  float *xDevice, *yDevice, *zDevice;
  aclrtMallocHost((void**)(&xHost), inputByteSize);
  aclrtMallocHost((void**)(&yHost), inputByteSize);
  aclrtMallocHost((void**)(&zHost), outputByteSize);
  // Allocate the device memory.
  aclrtMalloc((void**)&(xDevice), inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
  aclrtMalloc((void**)&(yDevice), inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
  aclrtMalloc((void**)&(zDevice), outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
  // Initialize the host memory.
  for (int i = 0; i < TOTAL_LENGTH; ++i) {
    xHost[i] = 1.0f;
    yHost[i] = 2.0f;
  }
  aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
  aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

  // Use the kernel launch symbol <<<>>> to call the kernel function to complete specified operations.
  VecAdd<<<USE_CORE_NUM, nullptr, stream>>>(xDevice, yDevice, zDevice);

  aclrtSynchronizeStream(stream);
  // Copy the computation result from the device to the host.
  aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
#undef printf
  for (int i = 0; i < TOTAL_LENGTH; i++) {
    printf("i%d\t Expect: %f\t\t\t\tResult: %f\n", i, 3.0f,
           zHost[i]);
  }
  // Release allocated resources.
  aclrtFree(xDevice);
  aclrtFree(yDevice);
  aclrtFree(zDevice);
  aclrtFreeHost(xHost);
  aclrtFreeHost(yHost);
  aclrtFreeHost(zHost);
  // Deinitialize AscendCL.
  aclrtDestroyStream(stream);
  aclrtDestroyContext(context);
  aclrtResetDevice(deviceId);
  aclFinalize();

  return 0;
}

The compilation commands are as follows. For details, see Heterogeneous Compilation.

export RT_INC=${INSTALL_DIR}/runtime/include
export RT_LIB=${INSTALL_DIR}/runtime/lib64

# Function: Compile the host and device code together to generate an executable file, which must be linked to libruntime.so.
# Compilation options --cce-soc-version and --cce-soc-core-type are used to compile the vector core program on AscendXXXYY.
$bisheng -O2 --cce-soc-version=AscendXXXYY --cce-soc-core-type=VecCore  -I$RT_INC -L$RT_LIB -lascendcl -lruntime QuickStartDemoVecAdd.cce  -o QuickStartDemoVecAdd -I${INSTALL_DIR}/compiler/tikcpp/tikcfw/ -I${INSTALL_DIR}/compiler/tikcpp/tikcfw/impl -I${INSTALL_DIR}/compiler/tikcpp/tikcfw/interface --std=c++17

The execution result is as follows:

$export LD_LIBRARY_PATH=$RT_LIB:$LD_LIBRARY_PATH
$ ./QuickStartDemoVecAdd
i0       Expect: 3.000000                               Result: 3.000000
i1       Expect: 3.000000                               Result: 3.000000
i2       Expect: 3.000000                               Result: 3.000000
i3       Expect: 3.000000                               Result: 3.000000
i4       Expect: 3.000000                               Result: 3.000000
i5       Expect: 3.000000                               Result: 3.000000
i6       Expect: 3.000000                               Result: 3.000000
i7       Expect: 3.000000                               Result: 3.000000
i8       Expect: 3.000000                               Result: 3.000000
i9       Expect: 3.000000                               Result: 3.000000
i10      Expect: 3.000000                               Result: 3.000000
...
i16383      Expect: 3.000000                               Result: 3.000000