快速上手

// 文件名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;
}