DCU异构程序——GEMM

最新推荐文章于 2025-05-01 19:20:50 发布
最新推荐文章于 2025-05-01 19:20:50 发布
阅读量442 收藏 8
点赞数 4
分类专栏： 异构 文章标签： linux 架构 分布式 c++
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.youkuaiyun.com/L_peanut/article/details/144984146
版权

目录

一、概述

二、程序实现

三、编译运行

一、概述

        HIP属于显式编程模型，需要在程序中明确写出并行控制语句，包括数据传输、核函数启动等。核函数是运行在DCU上的函数，在CPU端运行的部分称为主机端（主要是执行管理和启动），DCU端运行的部分称为设备端（用于执行计算）。大概的流程如下图：

HIP程序流程

        ①主机端将需要并行计算的数据通过hipMemcpy()传递给DCU（将CPU存储的内容传递给DCU的显存）；

        ②调用核函数启动函数hipLaunchKernelGGL()启动DCU，开始执行计算；

        ③设备端将计算好的结果数据通过hipMemcpy()从DCU复制回CPU。

        hipMemcpy()是阻塞式的，数据复制完成后才可以执行后续的程序；hipLanuchKernelGGL()是非阻塞式的，执行完后程序继续向后执行，但是在Kernel没有计算完成之前，最后一个hipMemcpy()是不会开始的，这是由于HIP的Stream机制。

二、程序实现

        下面是对GEMM的具体实现，GEMM.cpp：

#include <stdio.h>
#include <string.h>
#include "hip/hip_runtime.h"
#include <sys/time.h>

#define NUM 256
#define SCALAR_ZERO (double)(0)
#define LDS_NUM_ELEMENTS 2048
#define LDS_OFFSET_B 512
#define LDS_OFFSET_BLK 1024
#define DEPTH 8
#define MT0I 64
#define MT1J 64

#define MAC(A,B,DST) DST += A*B
#define TYPE_MAC(MULA,MULB,DST) DST = MAC(MULA,MULB,DST);
#define TT 4
#define SIZE_HALF_A 32
#define SIZE_HALF_B 32
#define SIZE_HALF_C 32

#define MAC_4x4\
    TYPE_MAC(rA[0],rB[0],rC[0]); \
    TYPE_MAC(rA[0],rB[1],rC[1]); \
    TYPE_MAC(rA[0],rB[2],rC[2]); \
    TYPE_MAC(rA[0],rB[3],rC[3]); \
    TYPE_MAC(rA[1],rB[0],rC[4]); \
    TYPE_MAC(rA[1],rB[1],rC[5]); \
    TYPE_MAC(rA[1],rB[2],rC[6]); \
    TYPE_MAC(rA[1],rB[3],rC[7]); \
    TYPE_MAC(rA[2],rB[0],rC[8]); \
    TYPE_MAC(rA[2],rB[1],rC[9]); \
    TYPE_MAC(rA[2],rB[2],rC[10]); \
    TYPE_MAC(rA[2],rB[3],rC[11]); \
    TYPE_MAC(rA[3],rB[0],rC[12]); \
    TYPE_MAC(rA[3],rB[1],rC[13]); \
    TYPE_MAC(rA[3],rB[2],rC[14]); \
    TYPE_MAC(rA[3],rB[3],rC[15]); \

#define MAC_4x4_BLK\
    TYPE_MAC(rA[0+TT],rB[0+TT],rC[0]); \
    TYPE_MAC(rA[0+TT],rB[1+TT],rC[1]); \
    TYPE_MAC(rA[0+TT],rB[2+TT],rC[2]); \
    TYPE_MAC(rA[0+TT],rB[3+TT],rC[3]); \
    TYPE_MAC(rA[1+TT],rB[0+TT],rC[4]); \
    TYPE_MAC(rA[1+TT],rB[1+TT],rC[5]); \
    TYPE_MAC(rA[1+TT],rB[2+TT],rC[6]); \
    TYPE_MAC(rA[1+TT],rB[3+TT],rC[7]); \
    TYPE_MAC(rA[2+TT],rB[0+TT],rC[8]); \
    TYPE_MAC(rA[2+TT],rB[1+TT],rC[9]); \
    TYPE_MAC(rA[2+TT],rB[2+TT],rC[10]); \
    TYPE_MAC(rA[2+TT],rB[3+TT],rC[11]); \
    TYPE_MAC(rA[3+TT],rB[0+TT],rC[12]); \
    TYPE_MAC(rA[3+TT],rB[1+TT],rC[13]); \
    TYPE_MAC(rA[3+TT],rB[2+TT],rC[14]); \
    TYPE_MAC(rA[3+TT],rB[3+TT],rC[15]); \

#define TYPE_MAC_WRITE(DST,SRC,ALPHA,REG,BETA) DST = 0 != (BETA) ? (ALPHA)*(REG) + (BETA)*(SRC) : (ALPHA)*(REG)

__device__ uint64_t inline readtime()
{
    uint64_t clock;
    asm volatile("s_waitcnt lgkmcnt(0)\n\t"
                "s_memtime %0\n\t"
                 "s_waitcnt lgkmcnt(0)\n\t"
                    : "=s" (clock));
    return clock;
}

__global__ void global_depth8_lds_2_bank(double *src_a, double *src_b, double *dst_c, double alpha, double beta, int size_m, int size_n, int size_k, uint64_t *clock_cycle)
{
    uint64_t cycle1 = readtime();
    __shared__ double localMemory[LDS_NUM_ELEMENTS];

    unsigned int serial = 0;
    unsigned int grj = (serial >> 5);
    unsigned int gri = (serial & 31);

    unsigned int goa = grj * size_m + gri * 2;
    unsigned int gob = grj * size_n + gri * 2;
    
    unsigned int lwa = serial * 2;
    unsigned int lwb = serial * 2 + LDS_OFFSET_B;

    double *local_write_A = localMemory + lwa;
    double *local_write_B = localMemory + lwb;

    unsigned int lrj =  (serial >> 4);
    unsigned int lri = (serial &15);

    unsigned int lra = lri * 2;
    unsigned int lrb = lrj * 2 + LDS_OFFSET_B;

    double *local_read_A = localMemory + lra;
    double *local_read_B = localMemory + lrb;

    unsigned int goc = (lri * size_n + lrj) * 2;
    double *global_address_C = dst_c + goc;
    double *global_address_A = src_a + goa;
    double *global_address_B = src_b + gob;

    int i,j,l;

    double rA[8], rB[8], rC[16];
    double global_a0, global_a1, global_b0, global_b1;

    rC[0] = SCALAR_ZERO;
    rC[1] = SCALAR_ZERO;
    rC[2] = SCALAR_ZERO;
    rC[3] = SCALAR_ZERO;
    rC[4] = SCALAR_ZERO;
    rC[5] = SCALAR_ZERO;
    rC[6] = SCALAR_ZERO;
    rC[7] = SCALAR_ZERO;
    rC[8] = SCALAR_ZERO;
    rC[9] = SCALAR_ZERO;
    rC[10] = SCALAR_ZERO;
    rC[11] = SCALAR_ZERO;
    rC[12] = SCALAR_ZERO;
    rC[13] = SCALAR_ZERO;
    rC[14] = SCALAR_ZERO;
    rC[15] = SCALAR_ZERO;

    global_a0 = *(global_address_A + 0);
    global_a1 = *(global_address_A + 1);
    global_b0 = *(global_address_B + 0);
    global_b1 = *(global_address_B + 1);

    global_address_A += DEPTH * size_m;
    global_address_B += DEPTH * size_n;
    
    *(local_write_A + 0) = global_a0;
    *(local_write_A + 1) = global_a1;
    *(local_write_B + 0) = global_b0;
    *(local_write_B + 1) = global_b1;

    lwa = (lwa + LDS_OFFSET_BLK) % LDS_NUM_ELEMENTS;
    lwb = (lwb + LDS_OFFSET_BLK) % LDS_NUM_ELEMENTS;
    
    local_write_A = localMemory + lwa;
    local_write_B = localMemory + lwb;

    __syncthreads();

    rA[0] = *(local_read_A + 0);
    rA[1] = *(local_read_A + 1);
    rA[2] = *(local_read_A + 0 + SIZE_HALF_A);
    rA[3] = *(local_read_A + 1 + SIZE_HALF_A);

    rB[0] = *(local_read_B + 0);
    rB[1] = *(local_read_B + 1);
    rB[2] = *(local_read_B + 0 + SIZE_HALF_B);
    rB[3] = *(local_read_B + 1 + SIZE_HALF_B);

    local_read_A += MT0I;
    local_read_B += MT1J;

    for(i = 0; i < size_k; i += 8)
    {
        global_a0 = *(global_address_A + 0);
        global_a1 = *(global_address_A + 1);
        global_b0 = *(global_address_B + 0);
        global_b1 = *(global_address_B + 1);
    
        global_address_A += DEPTH * size_m;
        global_address_B += DEPTH * size_n;

        rA[0+TT] = *(local_read_A + 0);
        rA[1+TT] = *(local_read_A + 1);
        rA[2+TT] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3+TT] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0+TT] = *(local_read_B + 0);
        rB[1+TT] = *(local_read_B + 1);
        rB[2+TT] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3+TT] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;
        MAC_4x4

        rA[0] = *(local_read_A + 0);
        rA[1] = *(local_read_A + 1);
        rA[2] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0] = *(local_read_B + 0);
        rB[1] = *(local_read_B + 1);
        rB[2] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;
        MAC_4x4_BLK

        rA[0+TT] = *(local_read_A + 0);
        rA[1+TT] = *(local_read_A + 1);
        rA[2+TT] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3+TT] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0+TT] = *(local_read_B + 0);
        rB[1+TT] = *(local_read_B + 1);
        rB[2+TT] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3+TT] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;
        MAC_4x4

        rA[0] = *(local_read_A + 0);
        rA[1] = *(local_read_A + 1);
        rA[2] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3] = *(local_read_A + 1 + SIZE_HALF_A);
    
        rB[0] = *(local_read_B + 0);
        rB[1] = *(local_read_B + 1);
        rB[2] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;
        MAC_4x4_BLK

        rA[0+TT] = *(local_read_A + 0);
        rA[1+TT] = *(local_read_A + 1);
        rA[2+TT] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3+TT] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0+TT] = *(local_read_B + 0);
        rB[1+TT] = *(local_read_B + 1);
        rB[2+TT] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3+TT] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;
        MAC_4x4

        rA[0] = *(local_read_A + 0);
        rA[1] = *(local_read_A + 1);
        rA[2] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0] = *(local_read_B + 0);
        rB[1] = *(local_read_B + 1);
        rB[2] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;
        MAC_4x4_BLK

        rA[0+TT] = *(local_read_A + 0);
        rA[1+TT] = *(local_read_A + 1);
        rA[2+TT] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3+TT] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0+TT] = *(local_read_B + 0);
        rB[1+TT] = *(local_read_B + 1);
        rB[2+TT] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3+TT] = *(local_read_B + 1 + SIZE_HALF_B);

        *(local_write_A + 0) = global_a0;
        *(local_write_A + 1) = global_a1;
        *(local_write_B + 0) = global_b0;
        *(local_write_B + 1) = global_b1;

        lwa = (lwa + LDS_OFFSET_BLK) % LDS_NUM_ELEMENTS;
        lwb = (lwa + LDS_OFFSET_BLK) % LDS_NUM_ELEMENTS;

        local_write_A = localMemory + lwa;
        local_write_B = localMemory + lwb;

        lra = (lra + LDS_OFFSET_BLK) % LDS_NUM_ELEMENTS;
        lrb = (lrb + LDS_OFFSET_BLK) % LDS_NUM_ELEMENTS;

        local_read_A = localMemory + lra;
        local_read_B = localMemory + lrb;
        MAC_4x4

        __syncthreads();

        rA[0] = *(local_read_A + 0);
        rA[1] = *(local_read_A + 1);
        rA[2] = *(local_read_A + 0 + SIZE_HALF_A);
        rA[3] = *(local_read_A + 1 + SIZE_HALF_A);

        rB[0] = *(local_read_B + 0);
        rB[1] = *(local_read_B + 1);
        rB[2] = *(local_read_B + 0 + SIZE_HALF_B);
        rB[3] = *(local_read_B + 1 + SIZE_HALF_B);

        local_read_A += MT0I;
        local_read_B += MT1J;

        MAC_4x4_BLK;
    }

    TYPE_MAC_WRITE(*(global_address_C+0), *(global_address_C+0), alpha, rC[0], beta);
	TYPE_MAC_WRITE(*(global_address_C+1), *(global_address_C+1), alpha, rC[1], beta);
	TYPE_MAC_WRITE(*(global_address_C+0+SIZE_HALF_C), *(global_address_C+0+SIZE_HALF_C), alpha, rC[2], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+SIZE_HALF_C), *(global_address_C+1+SIZE_HALF_C), alpha, rC[3], beta);
	
	TYPE_MAC_WRITE(*(global_address_C+0+size_n), *(global_address_C+0+size_n), alpha, rC[4], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+size_n), *(global_address_C+1+size_n), alpha, rC[5], beta);
	TYPE_MAC_WRITE(*(global_address_C+0+SIZE_HALF_C+size_n), *(global_address_C+0+SIZE_HALF_C+size_n), alpha, rC[6], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+SIZE_HALF_C+size_n), *(global_address_C+1+SIZE_HALF_C+size_n), alpha, rC[7], beta);
	
	TYPE_MAC_WRITE(*(global_address_C+0+size_n*SIZE_HALF_C), *(global_address_C+0+size_n*SIZE_HALF_C), alpha, rC[8], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+size_n*SIZE_HALF_C), *(global_address_C+1+size_n*SIZE_HALF_C), alpha, rC[9], beta);
	TYPE_MAC_WRITE(*(global_address_C+0+SIZE_HALF_C+size_n*SIZE_HALF_C), *(global_address_C+0+SIZE_HALF_C+size_n*SIZE_HALF_C), alpha, rC[10], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+SIZE_HALF_C+size_n*SIZE_HALF_C), *(global_address_C+1+SIZE_HALF_C+size_n*SIZE_HALF_C), alpha, rC[11], beta);
	
	TYPE_MAC_WRITE(*(global_address_C+0+size_n*(SIZE_HALF_C+1)), *(global_address_C+0+size_n*(SIZE_HALF_C+1)), alpha, rC[12], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+size_n*(SIZE_HALF_C+1)), *(global_address_C+1+size_n*(SIZE_HALF_C+1)), alpha, rC[13], beta);
	TYPE_MAC_WRITE(*(global_address_C+0+SIZE_HALF_C+size_n*(SIZE_HALF_C+1)), *(global_address_C+0+SIZE_HALF_C+size_n*(SIZE_HALF_C+1)), alpha, rC[14], beta);
	TYPE_MAC_WRITE(*(global_address_C+1+SIZE_HALF_C+size_n*(SIZE_HALF_C+1)), *(global_address_C+1+SIZE_HALF_C+size_n*(SIZE_HALF_C+1)), alpha, rC[15], beta);
	
    __syncthreads();
    *clock_cycle = readtime() - cycle1;
}

void mul_cpu(double *src_a, double *src_b, double *dst_c, double alpha, double beta, int size_m, int size_n, int size_k)
{
    int i, j, k;
    for(i = 0; i < size_m; i++)
    {
        for(j = 0; j < size_n; j++)
        {
            double sum = 0;
            for(k = 0; k < size_k; k++)
            {
                sum += src_a[k*size_m + i] * src_b[k*size_n + j];
            }
            dst_c[i*size_n + j] = alpha * sum + beta * dst_c[i*size_n + j];
        }
    }
}

int main(int argc, char *argv[])
{
    double *src_a, *src_b, *out_cpu;
    double *a_device, *b_device, *c_device, *out_gpu;
    double alpha = 2.0, beta = 3.0;

    int size_m = 64, size_n = 64, size_k = 128;
    int m = NUM;
    int i, error = 0;
    uint64_t *clock_cycle;
    uint64_t time[1];

    src_a = (double *)malloc(size_m * size_k * sizeof(double));
    src_b = (double *)malloc(size_k * size_n * sizeof(double));
    out_cpu = (double *)malloc(size_m * size_n * sizeof(double));
    out_gpu = (double *)malloc(size_m * size_n * sizeof(double));

    hipMalloc((void**)&a_device, size_m * size_k * sizeof(double));
    hipMalloc((void**)&b_device, size_k * size_n * sizeof(double));
    hipMalloc((void**)&c_device, size_m * size_n * sizeof(double));
    hipMalloc((void**)&clock_cycle, 1 * sizeof(uint64_t));

    for(i = 0; i < size_m * size_k; i++)
    {
        src_a[i] = rand()%128;
    }
    for(i = 0; i < size_k * size_n; i++)
    {
        src_b[i] = rand()%128;
    }
    for(i = 0; i < size_m * size_n; i++)
    {
        out_gpu[i] = out_cpu[i] = rand()%128;
    }

    hipInit(0);
    hipDevice_t device;
    hipCtx_t context;
    hipDeviceGet(&device, 0);
    hipCtxCreate(&context, 0, device);

    hipMemcpy(a_device, src_a, size_m * size_k * sizeof(double), hipMemcpyHostToDevice);
    hipMemcpy(b_device, src_b, size_n * size_k * sizeof(double), hipMemcpyHostToDevice);
    hipMemcpy(c_device, out_cpu, size_m * size_n * sizeof(double), hipMemcpyHostToDevice);
    hipLaunchKernelGGL(global_depth8_lds_2_bank, dim3(1,1,1), dim3(256,1,1), 0, 0, a_device, b_device, c_device, alpha, beta, size_m, size_n, size_k, clock_cycle);
    hipMemcpy(time, clock_cycle, 1 * sizeof(uint64_t), hipMemcpyDeviceToHost);
    hipMemcpy(out_gpu, c_device, size_n * size_m * sizeof(double), hipMemcpyDeviceToHost);
    mul_cpu(src_a, src_b, out_cpu, alpha, beta, size_m, size_n, size_k);

    for(i = 0; i < size_n * size_m; i++)
    {
        if(fabs(out_gpu[i] - out_cpu[i]) > 1e-6)
        {
            error++;
            printf("%d, %lf, %lf\n", i, out_gpu[i], out_cpu[i]);
        }
    }

    if(error == 0)
        printf("**********result is ok!**********\n");

    double run_time = (double)time[0]/(1443.0*1000);
    printf("clock %lld, time is %f (ms), GFlops is %f GFlops\n\n", time[0], run_time, (double)(2*size_m*size_n*size_k)/(run_time*1000000));

    printf("Finish\n");

    return 0;
}

三、编译运行

        HIP程序采用hipcc编译。

运行结果：

国产加速器海光DCU&GPGPU深算处理器异构编程实战（上）
qq_27815483的博客
08-25 3544
我们也简单概括并比较了CPU和类GPU架构，可以看出CPU和类GPU架构在围绕低延迟和高吞吐量上的架构设计的巨大差异，这种差异必然会导致编程方式上的不同。DCU作为类GPU架构的一种也具有与类GPU相似的特性。因此，我们在详细开展DCU编程介绍之前，首先对DCU系统的硬件架构进行详细的介绍，这样大家在后续章节学习编程和程序优化的时候就会对DCU编程有更深入的理解。本章将会从DCU硬件架构和DCU节点系统架构两个方面开展详细的架构介绍。
国产加速器海光DCU&GPGPU深算处理器异构编程实战（下）
qq_27815483的博客
08-26 2008
Fortran语言自20世纪50年代诞生至今，一直被广泛应用于计算物理学、计算化学、流体力学、气候预测、有限元分析等计算密集的学科领域。同C/C++语言一样，Fortran语言也是科学/工程计算领域的主流编程语言，使用Fortran语言编写的应用程序运行在世界各地大大小小的计算集群上，为人类的各种研究与创新提供数据支撑。
学习笔记-Deepseek开源第三天： DeepGEMM 技术
地球空间
02-27 604
这只是个人的简单学习笔记，详细信息请登录官网和相关社区了解。
fortran调用MKL函数库中的gemm的fortran95接口计算矩阵相乘
chd_lkl的博客
07-02 2379
关于MKL函数库的使用参见下面视频 http://v.fcode.cn/video-use_library.html 下面给出示例代码 program TestMKLgemm use blas95 implicit none integer, parameter :: m = 2, k = 2, n = 2 real*8 :: A...
‌博士生存指南：如何用3个月从PyTorch进阶CUDA核函数开发？‌
像风一样自由
03-30 1612
3个月的CUDA进阶之路需要平衡理论学习与项目实践：前两周掌握基础语法，随后以性能优化为主线，最终通过混合编程打通落地方案。记住，每一个性能百分点的提升，都是对计算本质理解的深化。，覆盖内存优化、并行模式设计、混合编程接口三大核心模块，助你在3个月内构建高性能计算的核心竞争力。‌目标‌：掌握CUDA基础语法，实现首个性能超过PyTorch原生算子的自定义核函数。**‌目标‌：**实现高效GEMM（矩阵乘）核函数，性能达到cuBLAS的80%以上。益‌：FP16计算吞吐量可达FP32的8倍（理论值）
GEMM优化、并行优化、算子优化，从BLISlab项目入手！
weixin_45571628的博客
08-22 3328
GEMM优化在深度学习、科学计算等领域都有重要意义，本文通过视频、文字、转载等各种方式，让你对gemm有一个系统的、全面的、深入的认识。
GEMM算法及优化流程详解
热门推荐
独行侠
02-16 2万+
前言 神经网络前向耗时主要由卷积的耗时决定，参考賈杨青毕业论文，那么如何对卷积加速便成了重要的一个点，主流的加速方法有 以下几种： im2col+GEMM：目前几乎所有的主流计算框架包括 Caffe, MXNet 等都实现了该方法. 该方法把整个卷积过程转化成了GEMM过程，而GEMM在各种 BLAS 库中都是被极致优化的，一般来说，速度较快。 Winograd： Winograd 是存在...
海光DCU异构编程开发
08-06
DCU指的是深度计算处理器（Deep-learning ...资源内容包括异构计算技术与DTK开发套件、异构并行程序设计、异构程序常用性能优化、Fortran程序移植与hipfor、异构加速卡与OpenMP、OpenACC以及异构调试工具使用等。
DCU异构编程课程：软件调试与性能优化技术
资源摘要信息: "DCU异构编程与优化-调试与优化课程合集" 在IT行业，异构编程是一个涉及利用不同计算架构（如CPU、GPU、DSP等）进行程序设计与优化的领域。该领域的专家需要深入理解不同硬件平台的性能特点，并能够...
DCU异构编程实训代码：SampleChapter6
DCU编程指的是使用Delphi Component Unit（组件单元）进行的应用程序开发。Delphi是Borland公司推出的一款先进的可视化编程工具，它允许开发者以一种高效和直观的方式创建各种应用程序。DCU文件是Delphi编译后的二...
MKL 函数cblas_?gemm参数分析
u013806541的博客
04-03 3701
参考链接：https://software.intel.com/en-us/node/520775 函数原型： void cblas_sgemm (const CBLAS_LAYOUT Layout, const CBLAS_TRANSPOSE transa, constCBLAS_TRANSPOSE transb, const MKL_INT m, const MKL_INT n, cons
rt-linux下的cgroup cpu的死锁bug
zhaoxin1989的博客
04-27 1233
rt-linux下的cgroup cpu的死锁bug，rt-linux，优先级继承，cgroup，死锁，cgroup period timer，ktimer，rt-linux下的hrtimer
Linux Shell 重定向与管道符号（＞, ＞＞, |）的实现机制
最新发布
imhikaru的专栏
05-01 598
重定向目标 | 标准输出 → 文件 | 标准输出 → 文件末尾 | 命令输出 → 管道写端，命令输入 ← 管道读端 || 用途 | 输出重定向到文件（覆盖原内容） | 输出追加到文件末尾 | 将前一命令的输出作为后一命令的输入 || 项目 || 系统调用 |
Linux 服务管理两种方式service和systemctl
lza20001103的博客
04-29 1874
Linux 服务管理两种方式service和systemctl
Ubuntu平台使用aarch64-Linux交叉编译opencv库并移植RK3588S边缘端
没有余地 EliasJie
04-30 577
Ubuntu平台使用aarch64-Linux交叉编译opencv库并移植ARM边缘端
Linux常用指令
ya3288426755的博客
04-23 1928
最近在进行操作系统的实验，涉及到Linux系统的使用，下面总结一些常用指令，涵盖文件管理、系统信息、进程管理、用户与权限、网络工具、文本处理、压缩与归档、系统维护等核心功能。：将前一个命令的输出作为后一个命令的输入。：将命令输出保存到文件。
Chromium 134 编译指南 - Android 篇：从Linux版切换到Android版（六）
守城小轩的技术窝棚
05-01 462
通过本篇文章的指导，您已经成功将Linux版的Chromium源代码转换为同时支持Android平台的配置。这种方法不仅节省了重新下载整个代码库的时间和带宽，还让您的开发环境更加灵活，能够在不同平台间快速切换。理解了.gclient文件的作用和重要性修改了配置文件，添加了Android作为目标平台同步了Android平台所需的特定依赖项验证了配置是否成功完成这些步骤看似简单，但对于高效开发Chromium跨平台版本至关重要。
快速上手Linux的Web服务器的部署及优化
ZZZKKKRTSAE的博客
04-27 1377
www是world wide web的缩写，及万维网，也就是全球信息广播的意思通常说的上网就是使用www来查询用户所需要的信息www可以结合文字、图形、影像以及声音等多媒体，超链接的方式将信息以Internet传递到世界各处去当你连接www网站，该网站会提供一些数据，客户端要使用可以解析这些数据的软件来处理，那就是浏览器HTTP协议：全称为Hyper Text Transfer Protocol（超文本传输协议）简单来说http就是从服务器传输超文本（html）到本地浏览器的传输协议。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

猿核试Bug愁

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值