【cuda学习日记】2.cuda编程模型

2.1 在CPU上执行sum array

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx< N; idx ++)
    {
        C[idx] = A[idx] + B[idx];
    }
}

void initialData(float *ip, int size)
{
    time_t t;
    srand((unsigned int) time(&t));

    for (int i = 0; i < size; i++) {
        ip[i] = (float) (rand() & 0xff) / 10.0f;
    }
}

int main(int argc , char **argv)
{
    int nElem = 1024;
    size_t nBytes = nElem * sizeof(float);
    float *h_A, *h_B, *h_C;

    h_A = (float *) malloc (nBytes);
    h_B = (float *) malloc (nBytes);
    h_C = (float *) malloc (nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);

    sumArraysOnHost(h_A, h_B, h_C, nElem);

    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

2.2 查看grid, block的索引维度

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <cuda_runtime.h>

__global__ void checkIndex(void)
{
    printf("threadidx: (%d ,%d ,%d) blockidx:(%d ,%d ,%d) blockdim: (%d ,%d ,%d) gridDim: (%d ,%d ,%d)\n", 
    threadIdx.x, threadIdx.y, threadIdx.z,
    blockIdx.x, blockIdx.y, blockIdx.z,
    blockDim.x,blockDim.y,blockDim.z,
    gridDim.x, gridDim.y, gridDim.z
    );
}


int main(int argc , char **argv)
{
    int nElem = 6;
    
    dim3 block(3);
    dim3 grid ((nElem + block.x -1)/block.x);

    printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);
    printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);

    checkIndex<<<grid, block>>>();

    cudaDeviceReset();

    return 0;
}

输出
grid.x 2 grid.y 1 grid.z 1
block.x 3 block.y 1 block.z 1
threadidx: (0 ,0 ,0) blockidx:(1 ,0 ,0) blockdim: (3 ,1 ,1) gridDim: (2 ,1 ,1)
threadidx: (1 ,0 ,0) blockidx:(1 ,0 ,0) blockdim: (3 ,1 ,1) gridDim: (2 ,1 ,1)
threadidx: (2 ,0 ,0) blockidx:(1 ,0 ,0) blockdim: (3 ,1 ,1) gridDim: (2 ,1 ,1)
threadidx: (0 ,0 ,0) blockidx:(0 ,0 ,0) blockdim: (3 ,1 ,1) gridDim: (2 ,1 ,1)
threadidx: (1 ,0 ,0) blockidx:(0 ,0 ,0) blockdim: (3 ,1 ,1) gridDim: (2 ,1 ,1)
threadidx: (2 ,0 ,0) blockidx:(0 ,0 ,0) blockdim: (3 ,1 ,1) gridDim: (2 ,1 ,1)

2.4 cuda核函数
kernel_name<<<grid, block>>> (argument list);

2.5 核函数限定符
global 在设备端执行，主机端或计算能力>3的设备端调用，必须有void 返回类型
device 在设备端执行，设备端调用
host 在主机端执行，主机端调用

2.6 在GPU中sum，和HOST sum array做对比

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <sys/time.h>

#define CHECK(call) \
 {\
    const cudaError_t error = call; \
    if (error != cudaSuccess)\
    {\
        printf("Error: %s: %d\n", __FILE__, __LINE__);\
        printf("code :%d reason :%s\n", error , cudaGetErrorString(error));\
        exit(1);\
    }\
}

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;
    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i])> epsilon)
        {
            match = 0;
            printf("Array do not match\n");
            printf("host %5.2f gpu % 5.2f at current %d\n", hostRef[i], gpuRef[i], i);
            break;

        }
    }
    if (match) printf("array matches\n");
}

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx< N; idx ++)
    {
        C[idx] = A[idx] + B[idx];
    }
}

void initialData(float *ip, int size)
{
    time_t t;
    srand((unsigned int) time(&t));

    for (int i = 0; i < size; i++) {
        ip[i] = (float) (rand() & 0xff) / 10.0f;
    }
}

__global__ void sumArraysOnGPU(float *A, float *B, float *C)
{
    int i  = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];
}


double cpusec()
{
    struct timeval tp;
    gettimeofday(&tp, NULL);
    return ((double) tp.tv_sec + (double)tp.tv_usec* 1.e-6);
}

int main(int argc , char **argv)
{
    printf("%s starting\n", argv[0]);

    int dev = 0;
    cudaSetDevice(dev);


    //set up data
    int nElem = 32;
    size_t nBytes = nElem * sizeof(float);
    float *h_A, *h_B, *hostRef, *gpuRef;

    h_A = (float *) malloc (nBytes);
    h_B = (float *) malloc (nBytes);
    hostRef = (float *) malloc (nBytes);
    gpuRef = (float *) malloc (nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);

    memset(hostRef,0, nBytes);
    memset(gpuRef,0, nBytes);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    cudaMalloc((float**)&d_A, nBytes);
    cudaMalloc((float**)&d_B, nBytes);
    cudaMalloc((float**)&d_C, nBytes);

    //transfer data from host to device
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

    dim3 block(nElem);
    dim3 grid(nElem/block.x);

    sumArraysOnGPU<<<grid,block>>>(d_A, d_B, d_C);
    printf("execution config <<<%d, %d>>>\n", grid.x, block.x);

    //copy kernel result back to host
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    checkResult(hostRef, gpuRef, nElem);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return 0;
}

输出
execution config <<<1, 32>>>
array matches

2.8 尝试用nvprof 检查执行时间，报错找不到dll文件，可以参考
https://blog.csdn.net/qq_41607336/article/details/126741908
解决，
但最后还是报warning， 跟显卡相关
======= Warning: nvprof is not supported on devices with compute capability 8.0 and higher.

2.9 加上device信息打印，以及cudaevent 计时， 参考https://blog.csdn.net/qq_42681630/article/details/144895351

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <windows.h>

typedef unsigned long DWORD;

#define CHECK(call) \
 {\
    const cudaError_t error = call; \
    if (error != cudaSuccess)\
    {\
        printf("Error: %s: %d\n", __FILE__, __LINE__);\
        printf("code :%d reason :%s\n", error , cudaGetErrorString(error));\
        exit(1);\
    }\
}

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;
    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i])> epsilon)
        {
            match = 0;
            printf("Array do not match\n");
            printf("host %5.2f gpu % 5.2f at current %d\n", hostRef[i], gpuRef[i], i);
            break;

        }
    }
    if (match) printf("array matches\n");
}

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx< N; idx ++)
    {
        C[idx] = A[idx] + B[idx];
    }
}

void initialData(float *ip, int size)
{
    time_t t;
    srand((unsigned int) time(&t));

    for (int i = 0; i < size; i++) {
        ip[i] = (float) (rand() & 0xff) / 10.0f;
    }
}

__global__ void sumArraysOnGPU(float *A, float *B, float *C)
{
    int i  = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];
}



int main(int argc , char **argv)
{
    printf("%s starting\n", argv[0]);

    int dev = 0;
    cudaDeviceProp deviceprop;
    CHECK(cudaGetDeviceProperties(&deviceprop,dev));
    printf("Using Device %d : %s\n", dev, deviceprop.name);
    CHECK(cudaSetDevice(dev));
    

    printf( "Compute capability:  %d.%d\n", deviceprop.major, deviceprop.minor );
    printf( "Clock rate:  %d\n", deviceprop.clockRate );
    printf( "Memory Clock rate:  %d\n", deviceprop.memoryClockRate );
    printf( "Memory busWidth:  %d\n", deviceprop.memoryBusWidth );
    printf( "   --- Memory Information for device  ---\n");
    // printf( "Total global mem:  %ld\n", prop.totalGlobalMem );
    printf( "Total global mem:  %zu\n", deviceprop.totalGlobalMem );
    printf( "Total constant Mem:  %ld\n", deviceprop.totalConstMem );
    printf( "Max mem pitch:  %ld\n", deviceprop.memPitch );
    printf( "Texture Alignment:  %ld\n", deviceprop.textureAlignment );
    
    printf( "   --- MP Information for device  ---\n" );
    printf( "Multiprocessor count:  %d\n",
                deviceprop.multiProcessorCount );
    printf( "Shared mem per mp:  %ld\n", deviceprop.sharedMemPerBlock );
    printf( "Registers per mp:  %d\n", deviceprop.regsPerBlock );
    printf( "Threads in warp:  %d\n", deviceprop.warpSize );
    printf( "Max threads per block:  %d\n",
                deviceprop.maxThreadsPerBlock );
    printf( "Max thread dimensions:  (%d, %d, %d)\n",
                deviceprop.maxThreadsDim[0], deviceprop.maxThreadsDim[1],
                deviceprop.maxThreadsDim[2] );
    printf( "Max grid dimensions:  (%d, %d, %d)\n",
                deviceprop.maxGridSize[0], deviceprop.maxGridSize[1],
                deviceprop.maxGridSize[2] );
    printf( "\n" );

    //set up data
    int nElem = 1<<24;
    size_t nBytes = nElem * sizeof(float);
    float *h_A, *h_B, *hostRef, *gpuRef;

    h_A = (float *) malloc (nBytes);
    h_B = (float *) malloc (nBytes);
    hostRef = (float *) malloc (nBytes);
    gpuRef = (float *) malloc (nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);

    memset(hostRef,0, nBytes);
    memset(gpuRef,0, nBytes);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    cudaMalloc((float**)&d_A, nBytes);
    cudaMalloc((float**)&d_B, nBytes);
    cudaMalloc((float**)&d_C, nBytes);

    //transfer data from host to device
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);


    int Ilen = 1024;
    dim3 block(Ilen);
    dim3 grid((nElem + block.x - 1)/block.x);
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    cudaEventRecord(start);

    sumArraysOnGPU<<<grid,block>>>(d_A, d_B, d_C);
    printf("execution config <<<%d, %d>>>\n", grid.x, block.x);
    
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("Kernel execution time: %f ms\n", milliseconds);



    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    
    //copy kernel result back to host
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    checkResult(hostRef, gpuRef, nElem);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return 0;
}

输出：
Using Device 0 : NVIDIA GeForce RTX 4090
Compute capability: 8.9
Clock rate: 2520000
Memory Clock rate: 10501000
Memory busWidth: 384
— Memory Information for device —
Total global mem: 25756696576
Total constant Mem: 65536
Max mem pitch: 2147483647
Texture Alignment: 512
— MP Information for device —
Multiprocessor count: 128
Shared mem per mp: 49152
Registers per mp: 65536
Threads in warp: 32
Max threads per block: 1024
Max thread dimensions: (1024, 1024, 64)
Max grid dimensions: (2147483647, 65535, 65535)

execution config <<<16384, 1024>>>
execution config <<<16384, 1024>>>
Kernel execution time: 0.558752 ms
array matches

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：/a/949863.html

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

141.环形链表 142.环形链表II

141.环形链表 142.环形链表II

141.环形链表 & 142.环形链表II 141.环形链表 思路&#xff1a;快慢指针 or 哈希表 快慢指针代码&#xff1a; class Solution { public:bool hasCycle(ListNode *head) {if(headnullptr||head->nextnullptr)return false;ListNode *fasthead->next; //不能设置成…
阅读更多...
信用租赁系统助力企业实现免押金租赁新模式

信用租赁系统助力企业实现免押金租赁新模式

内容概要 在现代商业环境中&#xff0c;信用租赁正在迅速崛起。通过结合大数据与区块链技术&#xff0c;信用租赁系统彻底改变了传统的租赁流程。什么是信用租赁呢&#xff1f;简单说&#xff0c;就是不需要押金&#xff0c;你也能够租到你想要的物品&#xff0c;这对企业和消…
阅读更多...
el-select下拉框在弹框里面错位

el-select下拉框在弹框里面错位

问题出现 Element Plus 是一个基于 Vue 3 的组件库&#xff0c;el-select 是其中一个用于选择器的组件。在 el-select 组件中&#xff0c;teleported 属性用于控制下拉菜单的渲染位置。 解决方法 teleported 属性「element-plus」 popper-append-to-body属性「element」 ‌…
阅读更多...
IO进程day1

IO进程day1

一、思维导图
阅读更多...
力扣-21-合并两个有序链表

力扣-21-合并两个有序链表

思路&#xff1a; 因为是升序的两个链表&#xff0c;我们可以进行数据域比大小&#xff0c;然后把p3(自己创建的)的指针域指向小的那个 注&#xff1a;一定要先判断两个指针为0的情况
阅读更多...
人工智能的发展领域之GPU加速计算的应用概述、架构介绍与教学过程

人工智能的发展领域之GPU加速计算的应用概述、架构介绍与教学过程

文章目录 一、架构介绍GPU算力平台概述优势与特点 二、注册与登录账号注册流程GPU服务器类型配置选择指南内存和存储容量网络带宽CPU配置 三、创建实例实例创建步骤镜像选择与设置 四、连接实例SSH连接方法远程桌面配置 一、架构介绍 GPU算力平台概述 一个专注于GPU加速计算的…
阅读更多...
QT实现 端口扫描暂停和继续功能 3

QT实现 端口扫描暂停和继续功能 3

上篇QT给端口扫描工程增加线程2-CSDN博客 为按钮pushButton_Stop添加clicked事件&#xff0c;功能为暂停扫描&#xff0c;并在暂停后显示继续按钮&#xff0c;点击继续按钮之后继续扫描 1.更新UI 添加继续按钮 点击转到槽则会自动声明 2. 更新 MainWindow.h 需要新增的部分…
阅读更多...
汽车微处理器安全机制以及测试介绍

汽车微处理器安全机制以及测试介绍

本文介绍了三类汽车微处理器安全机制&#xff1a;硬件类、软件类和混合类&#xff0c;旨在提高系统的可靠性和安全性。硬件类安全机制包括逻辑内建自测试&#xff08;Logic-BIST&#xff09;、三重模块冗余&#xff08;TMR&#xff09;、内存内建自测试&#xff08;Memory-BIST…
阅读更多...
【Azure Redis 缓存】Azure Redis 遇见的连接不上问题和数据丢失的情况解答

【Azure Redis 缓存】Azure Redis 遇见的连接不上问题和数据丢失的情况解答

问题描述 PHP应用再连接Azure Redis服务时&#xff0c;出现Connection Timed out。当通过升级提高Azure Redis的性能时候&#xff0c;发现之前的数据丢失了。 image.png 问题解答 当Redis服务出现Timeout的情况时&#xff0c;可以从Redis服务的指标(Metrics)开始查看&#xff0…
阅读更多...
python学习笔记—15—数据容器之列表

python学习笔记—15—数据容器之列表

1. 数据容器 列表(list)、元组(tuple)、字符串(str)、集合(set)、字典(dict) 2. 列表 (1) 定义 tmp_list ["super", "carry", "doinb"] print(f"tmp_list {tmp_list}, tmp_list type is {type(tmp_list)}") tmp_list1 ["doi…
阅读更多...
记录一次面试中被问到的问题 (HR面)

记录一次面试中被问到的问题 (HR面)

文章目录 一、你对公司的了解多少二、为什么对这个岗位感兴趣三、不能说的离职原因四、离职原因高情商回复五、你的核心优势是什么六、你认为你比其他面试候选人的优势是什么七、不要提及情感 一、你对公司的了解多少 准备要点&#xff1a; 在面试前&#xff0c;对公司进行充分…
阅读更多...
VLMs之Agent之CogAgent：《CogAgent: A Visual Language Model for GUI Agents》翻译与解读

VLMs之Agent之CogAgent：《CogAgent: A Visual Language Model for GUI Agents》翻译与解读

VLMs之Agent之CogAgent&#xff1a;《CogAgent: A Visual Language Model for GUI Agents》翻译与解读 导读&#xff1a;这篇论文介绍了CogAgent&#xff0c;一个专注于图形用户界面 (GUI) 理解和导航的视觉语言模型 (VLM)。这篇论文提出了一种新的视觉语言模型 CogAgent&#…
阅读更多...
linux audio(1)-pulseaudio模块数据流

linux audio(1)-pulseaudio模块数据流

本文主要讨论pulseaudio模块的数据流。这里的模块（module）主要限制在sink和source这两种类型。其他类型的数据流后续有空 再撰文讨论。 pulseaudio的模块一般会启动一路线程进行数据的搬运和处理。 下面的是module-null-source模块的数据搬运线程启动代码。 进入thread_func…
阅读更多...
ros2-4.1 服务通信介绍

ros2-4.1 服务通信介绍

服务是ROS图中节点之间的另一种通信方法。服务分为客户端和服务端&#xff0c;客户端发送请求给服务端&#xff0c;服务端可以根据客户端的请求做一些处理&#xff0c;然后返回结果给客户端。也称为为请求-响应模型。 服务和话题的不同之处&#xff0c;话题是没有返回的&#…
阅读更多...
微信小程序之历史上的今天

微信小程序之历史上的今天

微信小程序之历史上的今天 需求描述 今天我们再来做一个小程序&#xff0c;主要是搜索历史上的今天发生了哪些大事&#xff0c;结果如下 当天的历史事件或者根据事件选择的历史事件的列表&#xff1a; 点击某个详细的历史事件以后看到详细信息&#xff1a; API申请和小程序…
阅读更多...
数据库模型全解析：从文档存储到搜索引擎

数据库模型全解析：从文档存储到搜索引擎

目录 前言1. 文档存储&#xff08;Document Store&#xff09;1.1 概念与特点1.2 典型应用1.3 代表性数据库 2. 图数据库&#xff08;Graph DBMS&#xff09;2.1 概念与特点2.2 典型应用2.3 代表性数据库 3. 原生 XML 数据库&#xff08;Native XML DBMS&#xff09;3.1 概念与…
阅读更多...
Vue3+TS+vite项目笔记1

Vue3+TS+vite项目笔记1

vue2与vue3的比较 源码的升级 使用Proxy代替defineProperty实现响应式。 重写虚拟DOM的实现和Tree-Shaking。 新的特性 Composition API&#xff08;组合API&#xff09;&#xff1a; setup ref与reactive computed与watch ...... 新的内置组件&#xff1a; Fragment T…
阅读更多...
Spring5框架之SpringMVC

Spring5框架之SpringMVC

目录 1.SpringMVC的入门案例 1.1 通过maven构建一个web项目 1.2 添加对应的依赖及Tomcat插件 1.3 创建SpringMVC的配置文件 1.4 在web.xml中注册DispatchServlet 1.5 创建自定义的Controller 1.6 在Springmvc配置文件中注册 原理分析&#xff1a; 2.SpringMVC基于注解的…
阅读更多...
Android Audio基础（53）——PCM逻辑设备Write数据

Android Audio基础（53）——PCM逻辑设备Write数据

1. 前言 本文，我们将以回放（Playback，播放音频）为例，讲解PCM Data是如何从用户空间到内核空间，最后传递到Codec。 在 ASoC音频框架简介中，我们给出了回放（Playback）PCM数据流示意图。： 对于Linux来说，由于分为 user space 和kernel space，而且两者之间数据不能随便…
阅读更多...
【漫话机器学习系列】039.点积（dot product）

【漫话机器学习系列】039.点积（dot product）

点积&#xff08;Dot Product&#xff09; 点积是线性代数中的一种基本运算&#xff0c;用于两个向量的操作。它是将两个向量按分量相乘并求和的结果&#xff0c;用于衡量两个向量在同一方向上的相似性。 点积的定义 给定两个相同维度的向量 和 &#xff0c;它们的点积定义为…
阅读更多...
最新文章

Copyright @ 2022~2023