一个完整的手工构建的cuda动态链接库工程 03记

news2025/1/13 10:33:45

1， 源代码

仅仅是加入了模板函数和对应的 .cuh文件，当前的目录结构如下：



icmm/gpu/add.cu

#include <stdio.h>
#include <cuda_runtime.h>

#include "inc/add.cuh"

// different name in this level for different typename, as extern "C" can not decorate template function that is in C++;

extern "C"  void  vector_add_gpu_s(float *A, float *B, float *C, int n)
{
  dim3 grid, block;

  block.x = 256;
  grid.x = (n + block.x - 1) / block.x;
  printf("CUDA kernel launch with %d blocks of %d threads\n", grid.x, block.x);

  vector_add_kernel<><<<grid, block>>>(A, B, C, n);
}

extern "C"  void  vector_add_gpu_d(double* A, double* B, double* C, int n)
{
  dim3 grid, block;

  block.x = 256;
  grid.x = (n + block.x - 1) / block.x;
  printf("CUDA kernel launch with %d blocks of %d threads\n", grid.x, block.x);

  vector_add_kernel<><<<grid, block>>>(A, B, C, n);
}

icmm/gpu/add.h

#pragma once

extern "C"  void  vector_add_gpu_s(float *A, float *B, float *C, int n);
extern "C"  void  vector_add_gpu_d(double* A, double* B, double* C, int n);

icmm/gpu/inc/add.cuh

#pragma once

template<typename T>
__global__ void vector_add_kernel(T *A, T *B, T *C, int n)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x;

  if (i < n)
  {
    C[i] = A[i] + B[i] + 0.0f;
  }
}

icmm/gpu/inc/sub.cuh

#pragma once

template<typename T>
__global__ void vector_sub_kernel(T *A, T *B, T *C, int n)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x;

  if (i < n)
  {
    C[i] = A[i] - B[i] + 0.0f;
  }
}

icmm/gpu/sub.cu

#include <stdio.h>
#include <cuda_runtime.h>
#include "inc/sub.cuh"

extern "C"  void  vector_sub_gpu_s(float *A, float *B, float *C, int n)
{
  dim3 grid, block;

  block.x = 256;
  grid.x = (n + block.x - 1) / block.x;
  printf("CUDA kernel launch with %d blocks of %d threads\n", grid.x, block.x);

  vector_sub_kernel<><<<grid, block>>>(A, B, C, n);
}

extern "C"  void  vector_sub_gpu_d(double *A, double *B, double *C, int n)
{
  dim3 grid, block;

  block.x = 256;
  grid.x = (n + block.x - 1) / block.x;
  printf("CUDA kernel launch with %d blocks of %d threads\n", grid.x, block.x);

  vector_sub_kernel<><<<grid, block>>>(A, B, C, n);
}

icmm/gpu/sub.h

#pragma once

extern "C"  void  vector_sub_gpu_s(float *A, float *B, float *C, int n);
extern "C"  void  vector_sub_gpu_d(double *A, double *B, double *C, int n);

icmm/include/icmm.h


#pragma once
#include<cuda_runtime.h>

void hello_print();
void ic_S_add(float* A, float* B, float *C, int n);
void ic_D_add(double* A, double* B, double* C, int n);

void ic_S_sub(float* A, float* B, float *C, int n);
void ic_D_sub(float* A, float* B, float *C, int n);

icmm/Makefile

#libicmm.so

TARGETS = libicmm.so
GPU_ARCH= -arch=sm_70


all: $(TARGETS)



sub.o: gpu/sub.cu
	nvcc    -Xcompiler -fPIC $(GPU_ARCH) -c $<

add.o: gpu/add.cu
	nvcc    -Xcompiler -fPIC $(GPU_ARCH) -c $<
#-dc
#-rdc=true

add_link.o: add.o
	nvcc   -Xcompiler -fPIC  $(GPU_ARCH) -dlink   -o $@  $<  -L/usr/local/cuda/lib64 -lcudart -lcudadevrt

ic_add.o: src/ic_add.cpp
	g++ -fPIC -c $<  -L/usr/local/cuda/lib64 -I/usr/local/cuda/include -lcudart -lcudadevrt -I./

ic_sub.o: src/ic_sub.cpp
	g++ -fPIC -c $<  -L/usr/local/cuda/lib64 -I/usr/local/cuda/include -lcudart -lcudadevrt -I./

$(TARGETS): sub.o ic_sub.o add.o ic_add.o add_link.o
	mkdir -p lib
	g++ -shared -fPIC  $^  -o lib/libicmm.so -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcudart -lcudadevrt 
	-rm -f *.o


.PHONY:clean
clean:
	-rm -f *.o lib/*.so test ./bin/test
	-rm -rf lib bin

icmm/makefile_bin

# executable
TARGET = test
GPU_ARCH = -arch=sm_70

all: $(TARGET)

add.o: gpu/add.cu
	nvcc -dc -rdc=true $(GPU_ARCH) -c $<

sub.o: gpu/sub.cu
	nvcc -dc -rdc=true $(GPU_ARCH) -c $<

add_link.o: add.o
	nvcc $(GPU_ARCH) -dlink   -o $@  $<  -L/usr/local/cuda/lib64 -lcudart -lcudadevrt

sub_link.o: sub.o
	nvcc $(GPU_ARCH) -dlink   -o $@  $<  -L/usr/local/cuda/lib64 -lcudart -lcudadevrt

ic_add.o: src/ic_add.cpp
	g++ -c $<  -L/usr/local/cuda/lib64 -I/usr/local/cuda/include -lcudart -lcudadevrt -I./

ic_sub.o: src/ic_sub.cpp
	g++ -c $<  -L/usr/local/cuda/lib64 -I/usr/local/cuda/include -lcudart -lcudadevrt -I./

test.o: testing/test.cpp
	g++ -c $< -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcudart -lcudadevrt -I./include

test: sub.o ic_sub.o sub_link.o add.o ic_add.o test.o add_link.o
	g++ $^ -L/usr/local/cuda/lib64 -lcudart -lcudadevrt   -o test
	mkdir ./bin
	cp ./test ./bin/
	-rm -f *.o

.PHONY:clean
clean:
	-rm -f *.o bin/* $(TARGET)

icmm/src/ic_add.cpp

#include <stdio.h>
#include <cuda_runtime.h>
#include "gpu/add.h"
//extern void vector_add_gpu(float *A, float *B, float *C, int n);

void hello_print()
{
  printf("hello world!\n");
}

//void ic_add(float* A, float* B, float *C, int n){  vector_add_gpu(A, B, C, n);}
void ic_S_add(float* A, float* B, float *C, int n)
{
  vector_add_gpu_s(A, B, C, n);
}

void ic_D_add(double* A, double* B, double* C, int n)
{
  vector_add_gpu_d(A, B, C, n);
}

icmm/src/ic_sub.cpp

#include <stdio.h>
#include <cuda_runtime.h>

#include "gpu/sub.h"
//extern void vector_add_gpu(float *A, float *B, float *C, int n);
void ic_S_sub(float* A, float* B, float *C, int n)
{
  vector_sub_gpu_s(A, B, C, n);
}

void ic_D_sub(double* A, double* B, double *C, int n)
{
  vector_sub_gpu_d(A, B, C, n);
}

icmm/testing/Makefile

#test

TARGET = test

all: $(TARGET)

CXX_FLAGS = -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcudart -lcudadevrt  -I../include -L../

test.o: test.cpp
	g++  -c $< $(CXX_FLAGS)

$(TARGET):test.o
	g++ $< -o $@ -L/usr/local/cuda/lib64 -lcudart -lcudadevrt  -L../lib  -licmm
	@echo "to execute: export LD_LIBRARY_PATH=${PWD}/../lib"

.PHONY:clean
clean:
	-rm -f *.o $(TARGET)

icmm/testing/test.cpp


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

#include "icmm.h"

void add_test_s(float* A, float* B, float* C, int n)
{
  ic_S_add(A, B, C, n);

  printf("Copy output data from the CUDA device to the host memory\n");

  float* h_C = (float*)malloc(n*sizeof(float));
  cudaMemcpy(h_C, C, n*sizeof(float), cudaMemcpyDeviceToHost);

  for (int i = 0; i < n; ++i)
  {
    printf("%3.2f ", h_C[i]);
    // if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {      fprintf(stderr, "Result verification failed at element %d!\n", i);      exit(EXIT_FAILURE);    }
  }

  printf("\nTest PASSED\n");

  free(h_C);
}

/**/
void add_test_d(double* A, double* B, double* C, int n)
{
  ic_D_add(A, B, C, n);

  printf("Copy output data from the CUDA device to the host memory\n");
  
  float *h_C = (float *)malloc(n*sizeof(double));
  cudaMemcpy(h_C, C, sizeof(double), cudaMemcpyDeviceToHost);

  for (int i = 0; i < n; ++i)
  {
    printf("%3.2f ", h_C[i]);
    // if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {      fprintf(stderr, "Result verification failed at element %d!\n", i);      exit(EXIT_FAILURE);    }
  }

  printf("\nTest PASSED\n");

  free(h_C);
}

/**/
void sub_test_s(float* A, float* B, float* C, int n)
{
  ic_S_sub(A, B, C, n);

  printf("Copy output data from the CUDA device to the host memory\n");

  float* h_C = (float*)malloc(n*sizeof(float));
  cudaMemcpy(h_C, C, n*sizeof(float), cudaMemcpyDeviceToHost);

  for (int i = 0; i < n; ++i)
  {
    printf("%3.2f ", h_C[i]);
    // if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {      fprintf(stderr, "Result verification failed at element %d!\n", i);      exit(EXIT_FAILURE);    }
  }

  printf("\nTest PASSED\n");

  free(h_C);
}


int main(void)
{
  int n = 50;
  size_t size = n * sizeof(float);

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

  for (int i = 0; i < n; ++i)
  {
    h_A[i] =  rand() / (float)RAND_MAX;
    h_B[i] =  rand() / (float)RAND_MAX;
  }

  float *d_A = NULL;
  float *d_B = NULL;
  float *d_C = NULL;

  cudaMalloc((void **)&d_A, size);
  cudaMalloc((void **)&d_B, size);
  cudaMalloc((void **)&d_C, size);
  cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
/*
  int threadsPerBlock = 256;
  int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
  printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);

  vector_add_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
*/

  //ic_add(d_A, d_B, d_C, n);
  add_test_s(d_A, d_B, d_C, n);
  sub_test_s(d_A, d_B, d_C, n);

  cudaFree(d_A);
  cudaFree(d_B);
  cudaFree(d_C);
  free(h_A);
  free(h_B);
  free(h_C);

  printf("Done\n");
  return 0;
}

2. 总结

.cu 代码给 g++ 的 .cpp 的代码需要使用 extern "C" 来修饰，所以一template 函数的实例化不能一直贯彻到 .cu 源代码的最顶层；

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

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈，一经查实，立即删除！

相关文章

nodejs微信小程序＋python＋PHP健身房信息管理系统的设计与实现-计算机毕业设计推荐

nodejs微信小程序＋python＋PHP健身房信息管理系统的设计与实现-计算机毕业设计推荐

目 录 摘 要 I ABSTRACT II 目 录 II 第1章 绪论 1 1.1背景及意义 1 1.2 国内外研究概况 1 1.3 研究的内容 1 第2章 相关技术 3 2.1 nodejs简介 4 2.2 express框架介绍 6 2.4 MySQL数据库 4 第3章 系统分析 5 3.1 需求分析 5 3.2 系统可行性分析 5 3.2.1技术可行性&#xff1a;…
阅读更多...
Redis常见类型

Redis常见类型

常用类型String字符串类型Hash字典类型List列表类型Set集合类型ZSet有序集合类型 Java程序操作Redis代码操作Redis 常用类型 String字符串类型 使用方式&#xff1a; 使用场景&#xff1a; Hash字典类型 字典类型(Hash) 又被成为散列类型或者是哈希表类型&#xff0c;它…
阅读更多...
【预测工具】不须编码的预测和数据可视化工具

【预测工具】不须编码的预测和数据可视化工具

有一天&#xff0c;我的同事问我&#xff0c;他应该如何做一个快速预测模型而不是Excel&#xff0c;并产生比线性回归或Excel图中的那些简单方程更好的结果。这是我的答案。 TableCurve 2D (Image by author) Sigmaplot很早以前就推出了这个软件。它已被广泛用于在数据中寻找最…
阅读更多...
C#基础学习--命名空间和程序集

C#基础学习--命名空间和程序集

引用其他程序集 编译器接受源代码文件并生成一个名为程序集的输出文件。 在许多项目中&#xff0c;会想使用来自其他程序集的类或类型。这些程序集可能来自BCL或第三方供应商&#xff0c;或者自己创建的。这些程序集称为类库&#xff0c;而且它们的程序集文件的名称通常以dll…
阅读更多...
Linux（13）：例行性工作排程

Linux（13）：例行性工作排程

例行性工程 听谓的排程是将工作安排执行的流程之意。 Linux 排程就是透过 crontab 与 at 这两个东西。 两种工作排程的方式&#xff1a; 一种是例行性的&#xff0c;就是每隔一定的周期要来办的事项&#xff1b; 一种是突发性的&#xff0c;就是这次做完以后就没有的那一种&a…
阅读更多...
领域驱动架构（DDD）建模

领域驱动架构（DDD）建模

一、背景 常见的软件开发方式是拿到产品需求后&#xff0c;直接考虑数据库中表应该如何设计&#xff0c;这种方式已经将设计与业务需求脱节&#xff0c;而更多的是直接考虑应该如何实现了&#xff0c;这有点本末倒置。而DDD是从领域(问题域)为出发点进行的设计方法。 领域驱动…
阅读更多...
C++面试宝典第1题：爬楼梯

C++面试宝典第1题：爬楼梯

题目 小乐爬楼梯&#xff0c;一次只能上1级或者2级台阶。楼梯一共有n级台阶&#xff0c;请问总共有多少种方法可以爬上楼&#xff1f; 解析 这道题虽然是一道编程题&#xff0c;但实际上更是一道数学题&#xff0c;着重考察应聘者的逻辑思维能力和分析解决问题的能力。 当楼梯只…
阅读更多...
严蔚敏数据结构p17（2.19）——p18(2.24) (c语言代码实现)

严蔚敏数据结构p17（2.19）——p18(2.24) (c语言代码实现)

2.19已知线性表中的元素以值递增有序排列,并以单链表作存储结构。试写一高效的算法, 删除表中所有值大于 mink 且小于 maxk 的元素(若表中存在这样的元素&#xff09;同时释放被删结点空间, 并分析你的算法的时间复杂度(注意:mink 和 maxk 是给定的个参变量,它们的值可以和表中…
阅读更多...
QNX时钟调研

QNX时钟调研

SYSPAGE_ENTRY()的使用&#xff0c;SYSPAGE_ENTRY 测试QNX下printf(“poo\n”);的耗时 #include <sys/neutrino.h> #include <inttypes.h> #include <stdio.h> #include <stdlib.h> #include <sys/syspage.h>int main( void ) {uint64_t cps, …
阅读更多...
P7 链表 链表头前方插入新节点

P7 链表 链表头前方插入新节点

目录 前言 01 链表头插入数据 示例代码 02 指定节点前方插入新节点 测试代码 前言 &#x1f3ac; 个人主页&#xff1a;ChenPi &#x1f43b;推荐专栏1: 《C》✨✨✨ &#x1f525; 推荐专栏2: 《 Linux C应用编程&#xff08;概念类&#xff09;_ChenPi的博客-CSDN博客》✨…
阅读更多...
Linux系统配置深度学习环境之cudnn安装

Linux系统配置深度学习环境之cudnn安装

前言 一个针对深度学习应用优化的 GPU 加速库。它提供了高性能、高可靠性的加速算法&#xff0c;旨在加速深度神经网络模型的训练和推理过程。 cuDNN 提供了一系列优化的基本算法和函数&#xff0c;包括卷积、池化、规范化、激活函数等&#xff0c;以及针对深度学习任务的高级功…
阅读更多...
【MySQL的DQL查询语句】

【MySQL的DQL查询语句】

MySQL的DQL查询语句-----在Navicat下 将学生表导入Navicat中查询语句查询一整张表查询年龄大于22年龄大于22的女生查找文科的学生查找六班的学生计算学生的总分 &#xff08;group by&#xff09;合并两表 &#xff08;join on xxxx&#xff09;合并两张表 并求总分先合并在聚合…
阅读更多...
一站式自动化：Ansible Playbook的全面学习之旅

一站式自动化：Ansible Playbook的全面学习之旅

1 Playbook介绍 1.1 Playbook介绍 playbook 是由一个或多个play组成的列表 Playbook 文件使用YAML来写的 1.2 YAML 1.2.1 介绍 是一种表达资料序列的格式&#xff0c;类似XML Yet Another Markup Language 2001年首次发表 www.yaml.org 1.2.2 特点 可读性好 和脚本语言交…
阅读更多...
机器学习笔记 - 什么是模型量化压缩技术？

机器学习笔记 - 什么是模型量化压缩技术？

一、简述 我们都知道现实世界是连续的状态，而计算机世界是离散的状态，这是什么意思呢？我们看一下下图，最右边的马力欧（高清）的状态，可以想象现实世界是连续的状态，而电脑世界在图像上呈现的是一格一格子的状态（左图）是离散的状态。 所以在计算机世界如果想要…
阅读更多...
医疗器械设备模组的具体应用

医疗器械设备模组的具体应用

直线模组是一种高精度、高速度的精密传动元件&#xff0c;目前被广泛应用在各种工业自动化领域&#xff1b;尤其是在激光加工、电子制造、医疗设备、物流设备和机器人等行业中&#xff0c;都发挥着重要作用&#xff0c;接下来我们看看医疗器械设备模组的具体应用吧&#xff01;…
阅读更多...
智慧灯杆系统平台架构设计需要考虑的几个要点

智慧灯杆系统平台架构设计需要考虑的几个要点

智慧灯杆是一种集成了各种先进技术的道路照明设施。它不仅提供照明服务&#xff0c;还可以具有物联网技术、视频监控、环境监测、广播通讯、无线网络覆盖等多种功能。这些智能功能可以通过互联网进行控制和管理&#xff0c;从而实现智慧城市的建设。智慧灯杆能够提升城市的智能…
阅读更多...
Webgis学习总结

Webgis学习总结

前言&#xff1a; 作者跟随视频学习了webgis内容进行如下学习复习总结 参考&#xff1a;新中地学习笔记 WebGIS第一课&#xff1a;测试高德API并通过&#xff1a; 注册申请高德API成为开发者&#xff0c;创建自己的项目和key进行项目初始化&#xff0c;可以使用JS API官方文…
阅读更多...
PyQt6 QComboBox下拉组合框控件

PyQt6 QComboBox下拉组合框控件

​锋哥原创的PyQt6视频教程&#xff1a; 2024版 PyQt6 Python桌面开发 视频教程(无废话版) 玩命更新中~_哔哩哔哩_bilibili2024版 PyQt6 Python桌面开发 视频教程(无废话版) 玩命更新中~共计34条视频&#xff0c;包括&#xff1a;2024版 PyQt6 Python桌面开发 视频教程(无废话…
阅读更多...
算法基础--双指针

算法基础--双指针

前面已经写了两篇关于算法方面的文章&#xff0c;这几天想了下&#xff0c;决定把这个算法整理成一个系列&#xff0c;除了是帮助自己巩固算法知识外&#xff0c;还能够把自己总结的每种算法的套路保存下来并分享给大家&#xff0c;这样以后即使是哪天想要重拾起来&#xff0c;…
阅读更多...
Apache Doris 详细教程（三）

Apache Doris 详细教程（三）

7、监控和报警 Doris 可以使用 Prometheus 和 Grafana 进行监控和采集&#xff0c;官网下载最新版即可。 Prometheus 官网下载&#xff1a;https://prometheus.io/download/ Grafana 官网下载&#xff1a;https://grafana.com/grafana/download Doris 的监控数据通过 FE 和…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023