源码安装 HIPIFY 和应用示例,将cuda生态源码转化成HIP生态源码

1，源码下载

GitHub - ROCm/HIPIFY: HIPIFY: Convert CUDA to Portable C++ CodeHIPIFY: Convert CUDA to Portable C++ Code. Contribute to ROCm/HIPIFY development by creating an account on GitHub.icon-default.png?t=N7T8https://github.com/ROCm/HIPIFY.git

git clone --recursive https://github.com/ROCm/HIPIFY.git
sudo apt install clang-dev

 2，编译并安装

2.1 通常方式

hipify-clang 文档：

https://github.com/ROCm/HIPIFY/blob/amd-staging/docs/hipify-clang.md

编译命令：

cmake -DCMAKE_INSTALL_PREFIX=../dist -DCMAKE_BUILD_TYPE=Release  ..
make -j install

此时 hipify-clang 会被安装到 HIPIFY/dist/bin 中，

测试：

cd  ../dist/bin
hipify --help

如果系统中存在多个llvm版本，在执行翻译命令时，比如hipify-clang ./vectorAdd.cu  --cuda-path=/usr/local/cuda-12.1可能会发生错误，如下提示：

CommandLine Error: Option 'static-func-full-module-prefix' registered more than once!
LLVM ERROR: inconsistency in registered CommandLine option

这时需要使用自制的LLVM，如下2.2节所示。

2.2 自制LLVM的方式

2.2.1  下载llvm源码

wget https://github.com/llvm/llvm-project/archive/refs/tags/llvmorg-17.0.6.tar.gz

解压，tar zxf llvmorg.....

2.2.2  配置编译LLVM

cd llvmorg.....mkdir -p build ../dist/localcd buildcmake -G "Unix Makefiles" ../llvm      \
-DLLVM_ENABLE_PROJECTS="clang;clang-tools-extra;compiler-rt"           \
-DLLVM_BUILD_EXAMPLES=ON           -DLLVM_TARGETS_TO_BUILD="host"      \
-DCMAKE_BUILD_TYPE=Release           -DLLVM_ENABLE_ASSERTIONS=ON       \
-DLLVM_ENABLE_RUNTIMES=all             -DLLVM_BUILD_LLVM_DYLIB=ON      \
-DCMAKE_INSTALL_PREFIX=../../dist/local

make -j

make -j install

测试时，llvm 被install在如下文件夹：

/home/hipper/ex_dock_hipify/dist/local

ls /home/hipper/ex_dock_hipify/dist/local 如图：

2.2.3 配置编译HIPIFY

指定 LLVM 安装目录的配置方法：

-DCMAKE_PREFIX_PATH=/home/hipper/ex_dock_hipify/dist/local


cmake \-DCMAKE_BUILD_TYPE=Release \-DCMAKE_INSTALL_PREFIX=../dist \-DCMAKE_PREFIX_PATH=/home/hipper/ex_dock_hipify/dist/local \-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-12.1  ..

make -j install

3. 示例

3.1翻译 .cu 文件到 .hip 文件

命令：

/home/hipper/ex_dock_hipify/HIPIFY/dist/bin/hipify-clang ./vectorAdd.cu  --cuda-path=/usr/local/cuda-12.1

会在 ./ 目录中生成 vectoreAdd.cu.hip 的文件。

其中，hipify-clang 并不检查输入文件的扩展名，比如这里的.cu，它只检查文件内部的内容，将cuda生态的关键字有机地翻译成 hip生态的关键字，输出文件会在原文件名的基础上加上 .hip 后缀；

源代码分别如下。

使用 cuda samples中的vectoradd.cu为例，源码如下：

vectorAdd.cu

#include <stdio.h>
#include <cuda_runtime.h>__global__ void vectorAdd(const float *A, const float *B, float *C,int numElements) {int i = blockDim.x * blockIdx.x + threadIdx.x;if (i < numElements) {C[i] = A[i] + B[i] + 0.0f;}
}int main(void) {cudaError_t err = cudaSuccess;int numElements = 50000;size_t size = numElements * sizeof(float);printf("[Vector addition of %d elements]\n", numElements);float *h_A = (float *)malloc(size);float *h_B = (float *)malloc(size);float *h_C = (float *)malloc(size);if (h_A == NULL || h_B == NULL || h_C == NULL) {fprintf(stderr, "Failed to allocate host vectors!\n");exit(EXIT_FAILURE);}for (int i = 0; i < numElements; ++i) {h_A[i] = rand() / (float)RAND_MAX;h_B[i] = rand() / (float)RAND_MAX;}float *d_A = NULL;err = cudaMalloc((void **)&d_A, size);if (err != cudaSuccess) {fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}float *d_B = NULL;err = cudaMalloc((void **)&d_B, size);if (err != cudaSuccess) {fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}float *d_C = NULL;err = cudaMalloc((void **)&d_C, size);if (err != cudaSuccess) {fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}printf("Copy input data from the host memory to the CUDA device\n");err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);if (err != cudaSuccess) {fprintf(stderr,"Failed to copy vector A from host to device (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);if (err != cudaSuccess) {fprintf(stderr,"Failed to copy vector B from host to device (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}int threadsPerBlock = 256;int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,threadsPerBlock);vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);err = cudaGetLastError();if (err != cudaSuccess) {fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}printf("Copy output data from the CUDA device to the host memory\n");err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);if (err != cudaSuccess) {fprintf(stderr,"Failed to copy vector C from device to host (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}for (int i = 0; i < numElements; ++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("Test PASSED\n");err = cudaFree(d_A);if (err != cudaSuccess) {fprintf(stderr, "Failed to free device vector A (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}err = cudaFree(d_B);if (err != cudaSuccess) {fprintf(stderr, "Failed to free device vector B (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}err = cudaFree(d_C);if (err != cudaSuccess) {fprintf(stderr, "Failed to free device vector C (error code %s)!\n",cudaGetErrorString(err));exit(EXIT_FAILURE);}free(h_A);free(h_B);free(h_C);printf("Done\n");return 0;
}

生成的 vectorAdd.cu.hip :

#include <stdio.h>
#include <hip/hip_runtime.h>__global__ void vectorAdd(const float *A, const float *B, float *C,int numElements) {int i = blockDim.x * blockIdx.x + threadIdx.x;if (i < numElements) {C[i] = A[i] + B[i] + 0.0f;}
}int main(void) {hipError_t err = hipSuccess;int numElements = 50000;size_t size = numElements * sizeof(float);printf("[Vector addition of %d elements]\n", numElements);float *h_A = (float *)malloc(size);float *h_B = (float *)malloc(size);float *h_C = (float *)malloc(size);if (h_A == NULL || h_B == NULL || h_C == NULL) {fprintf(stderr, "Failed to allocate host vectors!\n");exit(EXIT_FAILURE);}for (int i = 0; i < numElements; ++i) {h_A[i] = rand() / (float)RAND_MAX;h_B[i] = rand() / (float)RAND_MAX;}float *d_A = NULL;err = hipMalloc((void **)&d_A, size);if (err != hipSuccess) {fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}float *d_B = NULL;err = hipMalloc((void **)&d_B, size);if (err != hipSuccess) {fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}float *d_C = NULL;err = hipMalloc((void **)&d_C, size);if (err != hipSuccess) {fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}printf("Copy input data from the host memory to the CUDA device\n");err = hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice);if (err != hipSuccess) {fprintf(stderr,"Failed to copy vector A from host to device (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}err = hipMemcpy(d_B, h_B, size, hipMemcpyHostToDevice);if (err != hipSuccess) {fprintf(stderr,"Failed to copy vector B from host to device (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}int threadsPerBlock = 256;int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,threadsPerBlock);vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);err = hipGetLastError();if (err != hipSuccess) {fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}printf("Copy output data from the CUDA device to the host memory\n");err = hipMemcpy(h_C, d_C, size, hipMemcpyDeviceToHost);if (err != hipSuccess) {fprintf(stderr,"Failed to copy vector C from device to host (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}for (int i = 0; i < numElements; ++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("Test PASSED\n");err = hipFree(d_A);if (err != hipSuccess) {fprintf(stderr, "Failed to free device vector A (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}err = hipFree(d_B);if (err != hipSuccess) {fprintf(stderr, "Failed to free device vector B (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}err = hipFree(d_C);if (err != hipSuccess) {fprintf(stderr, "Failed to free device vector C (error code %s)!\n",hipGetErrorString(err));exit(EXIT_FAILURE);}free(h_A);free(h_B);free(h_C);printf("Done\n");return 0;
}

3.2 编译运行 vectorAdd.cu.hip

编译：

$ /opt/rocm/bin/hipcc ./vectorAdd.cu.hip -o vectorAdd

运行效果如下图：

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

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

相关文章

springboot230基于Spring Boot在线远程考试系统的设计与实现

springboot230基于Spring Boot在线远程考试系统的设计与实现

在线远程考试系统设计与实现 摘 要 信息数据从传统到当代&#xff0c;是一直在变革当中&#xff0c;突如其来的互联网让传统的信息管理看到了革命性的曙光&#xff0c;因为传统信息管理从时效性&#xff0c;还是安全性&#xff0c;还是可操作性等各个方面来讲&#xff0c;遇到…
阅读更多...
基于JAVA的毕业设计分配选题系统 开源项目

基于JAVA的毕业设计分配选题系统 开源项目

目录 一、摘要1.1 项目介绍1.2 项目录屏 二、功能模块2.1 专业档案模块2.2 学生选题模块2.3 教师放题模块2.4 选题审核模块 三、系统展示四、核心代码4.1 查询专业4.2 新增专业4.3 选择课题4.4 取消选择课题4.5 审核课题 五、免责说明 一、摘要 1.1 项目介绍 基于JAVAVueSpri…
阅读更多...
vmware虚拟机centos中/dev/cl_server8/root 空间不够

vmware虚拟机centos中/dev/cl_server8/root 空间不够

在使用vmware时发现自己的虚拟机的/dev/cl_server8/root空间不够了&#xff0c;没办法安装新的服务。所以查了一下改空间的办法。 1.在虚拟机关闭的状态下&#xff0c;选中需要扩容的虚拟机->设置->硬件-> 硬盘->扩展->填写扩大到的值。 2.打开虚拟机&#xff…
阅读更多...
jxls——自定义命令设置动态行高

jxls——自定义命令设置动态行高

文章目录 前言依赖引入绘制 jxls 批注的 excel 模板测试类编写自定义命令关于自动换行 前言 之前的博客中都简单说了数据的渲染和导出excel文件。包括固定的 表头结构&#xff0c;以及动态 表头和表数据等方式。 本篇博客主要说明自定义命令的方式&#xff0c;控制输出excel文…
阅读更多...
Unity AssetBundle详解，加载本地包、加载网络包代码全分享

Unity AssetBundle详解，加载本地包、加载网络包代码全分享

在Unity中，AssetBundle（简称AB包）是一种将多个文件或资源打包到一个文件中的方式，用于优化资源的加载和管理。使用AB包，可以按需加载资源，减少应用的初始加载时间，并可以实现热更新等功能。下面是一个基本的流程，展示如何在Unity中加载AB包并显示其中的资源。 步骤1：…
阅读更多...
springboot 实现本地文件存储

springboot 实现本地文件存储

springboot 实现本地文件存储 实现过程 上传文件保存文件&#xff08;本地磁盘&#xff09;返回文件HTTP访问服务器路径给前端&#xff0c;进行效果展示 存储 服务端接收上传的目的是提供文件的访问服务&#xff0c;对于SpringBoot而言&#xff0c;其对静态资源访问提供了很…
阅读更多...
H3C防火墙安全授权导入

H3C防火墙安全授权导入

一、防火墙授权概述 前面我们已经了解了一些防火墙的基本概念&#xff0c;有讲过防火墙除了一些基本功能&#xff0c;还有一些高级安全防护&#xff0c;但是这些功能需要另外独立授权&#xff0c;不影响基本使用。这里以H3C防火墙为例进行大概了解下。 正常情况下&#xff0c;防…
阅读更多...
深度学习_15_过拟合欠拟合

深度学习_15_过拟合欠拟合

过拟合和欠拟合 过拟合和欠拟合是训练模型中常会发生的事&#xff0c;如所要识别手势过于复杂&#xff0c;如五角星手势&#xff0c;那就需要更改高级更复杂的模型去训练&#xff0c;若用比较简单模型去训练&#xff0c;就会导致模型未能抓住手势的全部特征&#xff0c;那简单…
阅读更多...
[云原生] K8s之pod进阶

[云原生] K8s之pod进阶

一、pod的状态说明 &#xff08;1&#xff09;Pod 一直处于Pending状态 Pending状态意味着Pod的YAML文件已经提交给Kubernetes&#xff0c;API对象已经被创建并保存在Etcd当中。但是&#xff0c;这个Pod里有些容器因为某种原因而不能被顺利创建。比如&#xff0c;调度不成功(…
阅读更多...
原神抢码，米游社抢码-首发

原神抢码，米游社抢码-首发

本文章仅供学习使用-侵权请联系删除_2023年3月14日08:17:06 本来在深渊12层打不过的我偶然在刷到了一个dy的直播间&#xff0c;看到主播在抢码上号帮忙打深渊还号称痛苦号打不满不送原石的旗号我就决定扫码试试&#xff0c;在直播间内使用了两部手机互相扫码在扫了一下午的码后…
阅读更多...
自动驾驶技术详解

自动驾驶技术详解

&#x1f3ac;个人简介&#xff1a;一个全栈工程师的升级之路&#xff01; &#x1f4cb;个人专栏&#xff1a;自动驾驶技术 &#x1f380;CSDN主页 发狂的小花 &#x1f304;人生秘诀&#xff1a;学习的本质就是极致重复! 目录 一 自动驾驶视觉感知算法 1目标检测 1.1 两阶…
阅读更多...
css背景图片属性

css背景图片属性

基础代码&#xff1a; div {width: 200px;height: 200px;background: url(./css-logo.png); }<div></div> 1、background-repeat&#xff1a;默认是repeat 设置背景图片在容器内是否平铺。 background-repeat: repeat-y; background-repeat: repeat-x; background…
阅读更多...
消息中间件之RocketMQ源码分析(二十四)

消息中间件之RocketMQ源码分析(二十四)

事务消息 事务消息机制。 事务消息的发送和处理总结为四个过程: 1.生产者发送事务消息和执行本地事务 2.Broker存储事务消息 3.Broker回查事务消息 4.Broker提交或回滚事务消息 生产者发送事务消息和执行本地事务。 发送过程分为两个阶段: 第一阶段,发送事务消息 第二阶段,发…
阅读更多...
Java泛型简介

Java泛型简介

Java泛型简介 Java泛型是在Java 5中引入的一个特性&#xff0c;它允许程序员在编译时指定类、接口或方法能够接受的类型。泛型的主要目的是提供编译时类型安全检查&#xff0c;避免在运行时因为类型转换错误而导致的ClassCastException。 在没有泛型之前&#xff0c;Java中的集…
阅读更多...
Ubuntu系统使用Docker搭建Jupyter Notebook并实现无公网ip远程连接

Ubuntu系统使用Docker搭建Jupyter Notebook并实现无公网ip远程连接

文章目录 1. 选择与拉取镜像2. 创建容器3. 访问Jupyter工作台4. 远程访问Jupyter工作台4.1 内网穿透工具安装4.2 创建远程连接公网地址4.3 使用固定二级子域名地址远程访问 本文主要介绍如何在Ubuntu系统中使用Docker本地部署Jupyter Notebook&#xff0c;并结合cpolar内网穿透…
阅读更多...
C语言系列（所需基础：大学C语言及格）-4-转义字符/注释/选择语句

C语言系列（所需基础：大学C语言及格）-4-转义字符/注释/选择语句

文章目录 一、转义字符二、注释三、选择语句 一、转义字符 加上\会讲原来的字符改变意思&#xff0c;即进行转义 例如\t会使t变成\t用于表示转义字符&#xff0c;使得t转义成水平制表符 其他转义字符&#xff1a; 三字母词&#xff08;展示\&#xff1f;的用处&#xff09;…
阅读更多...
k8s-001-Centos7内核升级

k8s-001-Centos7内核升级

1. 查看内核 [rootlocalhost ~]# uname -a 2. 执行的命令(安装最新版内核): 下载: rpm --import https://www.elrepo.org/RPM-GPG-KEY-elrepo.org 安装: rpm -Uvh http://www.elrepo.org/elrepo-release-7.0-2.el7.elrepo.noarch.rpm &#xff08; 查看最新版内核&…
阅读更多...
力扣hot100题解（python版33-35题）

力扣hot100题解（python版33-35题）

33、排序链表 给你链表的头结点 head &#xff0c;请将其按 升序 排列并返回 排序后的链表 。 示例 1&#xff1a; 输入&#xff1a;head [4,2,1,3] 输出&#xff1a;[1,2,3,4]示例 2&#xff1a; 输入&#xff1a;head [-1,5,3,4,0] 输出&#xff1a;[-1,0,3,4,5]示例 3&a…
阅读更多...
kafka架构详解

kafka架构详解

文章目录 概述kafaka架构Kafka的设计时什么样的Zookeeper 在 Kafka 中的作用知道 概述 Apache Kafka 是分布式发布 - 订阅消息系统&#xff0c;在 kafka 官网上对 kafka 的定义&#xff1a;一个分布式发布 - 订阅消息传递系统。 Kafka 最初由 LinkedIn 公司开发&#xff0c;Li…
阅读更多...
蚂蚁集团推动编制的全球首个隐私计算一体机国际标准发布

蚂蚁集团推动编制的全球首个隐私计算一体机国际标准发布

近日&#xff0c;IEEE 标准协会&#xff08;IEEE-SA&#xff09;正式发布并推行了由我国企业主导的全球首个隐私计算一体机国际标准《隐私计算一体机技术要求》&#xff08;IEEE 3156-2023&#xff09;。IEEE-SA是权威国际标准制定机构&#xff0c;该标准的成功发布意味着中国的…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023