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 CANN component directory. For example, if the installation is performed by the root user, the default file storage path is /usr/local/Ascend/cann.
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.
# Configure environment variables after installing the Toolkit package as the root user. source /usr/local/Ascend/cann/set_env.sh # Configure environment variables after installing the Toolkit package as a non-root user. source ${HOME}/Ascend/cann/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 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 * CACHELINE_SZ, ACL_MEM_MALLOC_HUGE_FIRST);
uint8_t InitValue[BLOCKS * CACHELINE_SZ] = {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 compile commands are as follows. For details about the compile options, see Compile Options.
# 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. # Compile option --npu-arch is used to specify the AI processor architecture version. $bisheng -O2 --npu-arch=dav-2201 -I$RT_INC -L$RT_LIB -lascendcl -lruntime QuickStartDemo.cce -o QuickStartDemo
The execution result is as follows:
1 2 3 4 5 | $ ./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;
// Perform initialization.
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);
// Perform deinitialization.
aclrtDestroyStream(stream);
aclrtDestroyContext(context);
aclrtResetDevice(deviceId);
aclFinalize();
return 0;
}
The compile commands are as follows. For details about the compile options, see Compile Options.
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. # Compile option --npu-arch is used to specify the AI processor architecture version. $bisheng -O2 --npu-arch=dav-2201 -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:
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 | $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 |