unified-runtime编译与验证

一.创建容器

docker run --gpus all --shm-size=32g -ti \
            -e NVIDIA_VISIBLE_DEVICES=all --privileged --net=host \
            --rm -it \
            -v $PWD:/home \
            -w /home ghcr.io/intel/llvm/ubuntu2204_build /bin/bash

二.编译unified-runtime

git clone https://github.com/oneapi-src/unified-runtime
cd unified-runtime
mkdir build
cd build
cmake -DUR_BUILD_ADAPTER_CUDA=ON -DUR_BUILD_ADAPTER_NATIVE_CPU=ON -DUMF_DISABLE_HWLOC=ON ..
make

三.生成一个cuda ptx kernel

tee cuda_copy.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>

__global__ void kernel_copy(float *input,float *output)
{
    unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;
    output[tid]=input[tid];
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -dc -lineinfo -arch=sm_86 -ptx cuda_copy.cu -o cuda_copy.ptx

四.API测试

tee um_query_device.cpp<<-'EOF'
#include <iostream>
#include <memory>
#include <stdlib.h>
#include <vector>
#include <iostream>
#include <fstream>
#include <sstream>
#include "ur_api.h"

#define ur_check(call) \
    do { \
        ur_result_t error = call; \
        if (error != UR_RESULT_SUCCESS) { \
            fprintf(stderr, " error in file '%s' in line %i: %d.\n", __FILE__, __LINE__,error); \
            exit(EXIT_FAILURE); \
        } \
    } while (0)
    
constexpr unsigned PAGE_SIZE = 4096;
template <typename T, size_t N> struct alignas(PAGE_SIZE) AlignedArray {
    T data[N];
};
	
int main(int, char *[]) {
    ur_result_t status;
    ur_check(urLoaderInit(0, nullptr));
    std::cout << "Platform initialized.\n";

    uint32_t adapterCount = 0;
    std::vector<ur_adapter_handle_t> adapters;
    uint32_t platformCount = 0;
    std::vector<ur_platform_handle_t> platforms;
    ur_check(urAdapterGet(0, nullptr, &adapterCount));
    adapters.resize(adapterCount);
    ur_check(urAdapterGet(adapterCount, adapters.data(), nullptr));
    ur_check(urPlatformGet(adapters.data(), adapterCount, 1, nullptr,&platformCount));
    platforms.resize(platformCount);
    ur_check(urPlatformGet(adapters.data(), adapterCount, platformCount,
                           platforms.data(), nullptr));
    for (auto p : platforms) {
        ur_api_version_t api_version = {};
        ur_check(urPlatformGetApiVersion(p, &api_version));
        std::cout << "API version: " << UR_MAJOR_VERSION(api_version) << "."
                  << UR_MINOR_VERSION(api_version) << std::endl;

        uint32_t deviceCount = 0;
        ur_check(urDeviceGet(p, UR_DEVICE_TYPE_GPU, 0, nullptr, &deviceCount));
        std::vector<ur_device_handle_t> devices(deviceCount);
        ur_check(urDeviceGet(p, UR_DEVICE_TYPE_GPU, deviceCount, devices.data(),
                             nullptr));
        for (auto d : devices) {
            ur_device_type_t device_type = UR_DEVICE_TYPE_ALL;
            ur_check(urDeviceGetInfo(
                d, UR_DEVICE_INFO_TYPE, sizeof(ur_device_type_t),
                static_cast<void *>(&device_type), nullptr));
            static const size_t DEVICE_NAME_MAX_LEN = 1024;
            char device_name[DEVICE_NAME_MAX_LEN] = {0};
            ur_check(urDeviceGetInfo(d, UR_DEVICE_INFO_NAME, DEVICE_NAME_MAX_LEN - 1,
                                static_cast<void *>(&device_name), nullptr));
            if (device_type == UR_DEVICE_TYPE_GPU) {
                std::cout << "Found a " << device_name << " gpu.\n";
            }
            ur_context_handle_t hContext;
            ur_check(urContextCreate(1, &d, nullptr, &hContext));
            std::ifstream inputFile("cuda_copy.ptx");
            std::ostringstream buffer;
            buffer << inputFile.rdbuf(); 
            std::string fileContent = buffer.str();
            inputFile.close();
            ur_program_handle_t hProgram;
            ur_check(urProgramCreateWithBinary(hContext, d, fileContent.length(), (const uint8_t *)fileContent.c_str(),nullptr,&hProgram));
            constexpr int a_size = 32;
            AlignedArray<float, a_size> a, b;
            for (auto i = 0; i < a_size; ++i) {
                a.data[i] = a_size - i;
                b.data[i] = 0;
            }
            status=urProgramBuild(hContext, hProgram, nullptr);
            ur_mem_handle_t dA, dB;
            ur_check(urMemBufferCreate(hContext, UR_MEM_FLAG_READ_WRITE,
                                       a_size * sizeof(int), nullptr, &dA));
            ur_check(urMemBufferCreate(hContext, UR_MEM_FLAG_READ_WRITE,
                                       a_size * sizeof(int), nullptr, &dB));
            ur_kernel_handle_t hKernel;
            ur_check(urKernelCreate(hProgram, "_Z11kernel_copyPfS_", &hKernel));
            ur_check(urKernelSetArgMemObj(hKernel, 0, nullptr, dA));
            ur_check(urKernelSetArgMemObj(hKernel, 1, nullptr, dB));
            ur_queue_handle_t queue;
            ur_check(urQueueCreate(hContext, d, nullptr, &queue));
            ur_check(urEnqueueMemBufferWrite(queue, dA, true, 0, a_size * sizeof(float),
                                             a.data, 0, nullptr, nullptr));
            ur_check(urEnqueueMemBufferWrite(queue, dB, true, 0, a_size * sizeof(float),
                                             b.data, 0, nullptr, nullptr));
            const size_t gWorkOffset[] = {0, 0, 0};
            const size_t gWorkSize[] = {a_size, 1, 1};
            const size_t lWorkSize[] = {1, 1, 1};
            ur_event_handle_t event;
            ur_check(urEnqueueKernelLaunch(queue, hKernel, 3, gWorkOffset, gWorkSize,
                                           lWorkSize, 0, nullptr, &event));

            ur_check(urEnqueueMemBufferRead(queue, dB, true, 0, a_size * sizeof(int),
                                            b.data, 1, &event, nullptr));

            ur_check(urQueueFinish(queue));
            ur_check(urContextRelease(hContext));
            for (auto i = 0; i < a_size; ++i) {
               printf("%.2f\n",b.data[i]);
            }
        }
    }
out:
    for (auto adapter : adapters) {
        urAdapterRelease(adapter);
    }
    urLoaderTearDown();
    return status == UR_RESULT_SUCCESS ? 0 : 1;
}
EOF
g++ -o um_query_device um_query_device.cpp -I../include  -L lib -lur_loader -lpthread
LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$PWD/lib ./um_query_device