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

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 源代码的最顶层；

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

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

相关文章

java企业财务管理系统springboot+jsp

java企业财务管理系统springboot+jsp

1、基本内容 &#xff08;1&#xff09;搭建基础环境&#xff0c;下载JDK、开发工具eclipse/idea。 &#xff08;2&#xff09;通过HTML/CSS/JS搭建前端框架。 &#xff08;3&#xff09;下载MySql数据库&#xff0c;设计数据库表&#xff0c;用于存储系统数据。 &#xff08;4…
阅读更多...
LongAddr

LongAddr

目录 1. 引言 2. AtomicInteger的局限性 3. AtomicInteger与LongAdder 的性能差异 4.LongAdder 的结构 LongAddr架构 Striped64中重要的属性 Striped64中一些变量或者方法的定义 Cell类 5. 分散热点的原理 具体流程图 6. 在实际项目中的应用 7. 总结 1. 引言 在这一…
阅读更多...
【risc-v】易灵思efinix FPGA riscv 时钟配置的一些总结

【risc-v】易灵思efinix FPGA riscv 时钟配置的一些总结

系列文章目录 分享一些fpga内使用riscv软核的经验&#xff0c;共大家参考。后续内容比较多&#xff0c;会做成一个系列。 本系列会覆盖以下FPGA厂商 易灵思 efinix 赛灵思 xilinx 阿尔特拉 Altera 本文内容隶属于【易灵思efinix】系列。 文章目录 系列文章目录前言一、pan…
阅读更多...
【算法】单调栈题单（矩阵系列、字典序最小、贡献法）⭐

【算法】单调栈题单（矩阵系列、字典序最小、贡献法）⭐

文章目录 题单来源经典题单496. 下一个更大元素 I&#xff08;单调栈模板题&#xff09;503. 下一个更大元素 II&#xff08;单调栈循环数组&#xff09;2454. 下一个更大元素 IV&#xff08;第二个更大的元素&#xff1a;两个单调栈&#xff09;456. 132 模式&#xff08;单调…
阅读更多...
java学习part19接口

java学习part19接口

113-面向对象(高级)-接口的使用_哔哩哔哩_bilibili 1.接口概念 个人认为是一种能力&#xff0c;某个类是否具有某种能力。一个类实现了一个接口就相当于学会了某些功能。 2.使用 接口里的属性都是全局常量public static final&#xff0c;即便不写也会自动加上。 3.多实现 4.接…
阅读更多...
Python---函数递归---练习：斐波那契数列（本文以递归算法为主）

Python---函数递归---练习：斐波那契数列（本文以递归算法为主）

编程思想&#xff1a; 如何利用数学模型&#xff0c;来解决对应的需求问题&#xff1b;然后利用代码实现对应的数据模型。 算法&#xff1a;使用代码实现对应的数学模型&#xff0c;从而解决对应的业务问题 程序 算法 数据结构 在经常使用的算法中&#xff0c;有两种非常…
阅读更多...
OGG实现Oracle19C到postgreSQL14的实时同步

OGG实现Oracle19C到postgreSQL14的实时同步

&#x1f4e2;&#x1f4e2;&#x1f4e2;&#x1f4e3;&#x1f4e3;&#x1f4e3; 哈喽&#xff01;大家好&#xff0c;我是【IT邦德】&#xff0c;江湖人称jeames007&#xff0c;10余年DBA及大数据工作经验 一位上进心十足的【大数据领域博主】&#xff01;&#x1f61c;&am…
阅读更多...
如何从 Jira 成功迁移到极狐GitLab，看这个就够了！

如何从 Jira 成功迁移到极狐GitLab，看这个就够了！

内容来源&#xff1a;https://about.gitlab.com/blog 作者&#xff1a;Melissa Ushakov Atlassian 之前表示&#xff0c;到 2024 年 2 月会全面终止对于其服务器端产品的支持。 随着 Jira Server 的生命周期即将结束&#xff0c;众多组织都在考虑将其敏捷项目管理工具从Jira 迁…
阅读更多...
51单片机应用从零开始（十）·指针

51单片机应用从零开始（十）·指针

指针 C语言指针是一种保存变量地址的数据类型。它可以让程序直接访问内存中的数据&#xff0c;而不需要通过变量名来访问。指针变量存储的是一个地址&#xff0c;这个地址指向内存中的某个位置&#xff0c;该位置存储了一个值。 在C语言中&#xff0c;可以使用&运算符取得一…
阅读更多...
网络安全现状

网络安全现状

威胁不断演变&#xff1a; 攻击者不断变化和改进攻击方法&#xff0c;采用更复杂、更隐秘的技术&#xff0c;以逃避检测和追踪。这包括新型的勒索软件、零日漏洞利用和社交工程攻击等。 供应链攻击&#xff1a; 攻击者越来越关注供应链的弱点&#xff0c;通过在供应链中植入恶…
阅读更多...
5_企业架构LNMP高可用负载均衡服务器

5_企业架构LNMP高可用负载均衡服务器

企业架构LNMP高可用负载均衡服务器之Nginx 学习目标和内容 1、能够描述负载均衡的作用 2、能够了解负载均衡常见实现方式 3、能够使用Nginx实现负载均衡 4、能够描述Nginx的常见负载均衡算法 一、背景描述及其方案设计 1、业务背景描述 时间&#xff1a;2011.6.-2013.9 发布产…
阅读更多...
移动平均滤波的原理和C代码

移动平均滤波的原理和C代码

移动平均滤波是一种简单有效的平滑信号的方法&#xff0c;它通过计算一系列数据点的平均值来减小信号中的波动。基本的移动平均滤波方法有两种&#xff1a;简单移动平均&#xff08;SMA&#xff09;和指数加权移动平均&#xff08;EWMA&#xff09;。 简单移动平均滤波&#xf…
阅读更多...
Stream API 方法使用总结

Stream API 方法使用总结

文章目录 1.1、Stream介绍1.2、Stream创建对象&#xff08;1&#xff09;empty()方法&#xff08;2&#xff09;of()方法&#xff08;3&#xff09;Arrays.stream()方法&#xff08;4&#xff09;list.stream()方法 1.3、Stream中间方法&#xff08;1&#xff09;filter()方法&…
阅读更多...
SpringBoot之自定义Starter

SpringBoot之自定义Starter

目录 一、自己的理解 1. 理解一 2. 理解二 二、自定义starter&#xff08;重点&#xff09; 三、以mybatis-spring-boot-starter为例进行分析 1. 写好自己的自动配置类逻辑 2. 创建自己的starter项目并引入自动配置类项目的依赖 3. 在其它项目中使用自定义的starter 一…
阅读更多...
如何开启Windows Server 2016 远端桌面

如何开启Windows Server 2016 远端桌面

使用GUI 设定 服务器管理器–> 本地服务器–> 远端桌面 启用远端桌面 远端–> 允许远端连线至此电脑 会提示防火墙设定跟电源设定 防火墙之前已经关闭了 完成
阅读更多...
设计基于STM32的温度传感器实时数据采集和显示系统

设计基于STM32的温度传感器实时数据采集和显示系统

温度传感器作为常见的传感器之一&#xff0c;被广泛应用于各种领域&#xff0c;如工业自动化、家电控制等。为了实时监测和控制温度&#xff0c;设计一个基于STM32的温度传感器实时数据采集和显示系统是很有必要的。本文将详细介绍如何设计这样一个系统&#xff0c;并提供相应的…
阅读更多...
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;…
阅读更多...
Gitee拉取代码报错You hasn‘t joined this enterprise! fatal unable to access

Gitee拉取代码报错You hasn‘t joined this enterprise! fatal unable to access

文章目录 一、问题二、解决2.1、进入**控制面板**2.2、进入**用户账户**2.3、进入**管理Windows凭据**2.4、**普通凭据**2.4.1、添加2.4.2、编辑 2.5、重新拉取|推送代码 三、最后 一、问题 Gitee拉取仓库代码的时候报错You hasnt joined this enterprise! fatal unable to ac…
阅读更多...
二十五、DSL查询文档(全文检索查询、精确查询、地理查询、复合查询)

二十五、DSL查询文档(全文检索查询、精确查询、地理查询、复合查询)

目录 一、全文检索查询 1、match查询 语法: 2、multi_match查询 语法: 3、match和mult_match的区别 二、精确查询 1、term查询&#xff1a; 语法&#xff1a; 2、range查询&#xff1a;&#xff08;范围查询&#xff09; 语法&#xff1a; 三、地理查询 1、geo_bou…
阅读更多...
SSM新闻发布管理系统

SSM新闻发布管理系统

SSM毕设分享 序号1&#xff1a;SSM新闻发布管理系统 1 项目简介 Hi&#xff0c;各位同学好&#xff0c;这里是郑师兄&#xff01; 今天向大家分享一个毕业设计项目作品【SSM新闻发布管理系统】 师兄根据实现的难度和等级对项目进行评分(最低0分&#xff0c;满分5分) 难度系数…
阅读更多...
最新文章

Copyright @ 2022~2023