源碼安裝 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.pswp.cn/news/711205.shtml
繁體地址，請注明出處：http://hk.pswp.cn/news/711205.shtml
英文地址，請注明出處：http://en.pswp.cn/news/711205.shtml

如若內容造成侵權/違法違規/事實不符，請聯系多彩編程網進行投訴反饋email:809451989@qq.com，一經查實，立即刪除！

相關文章

springboot230基于Spring Boot在線遠程考試系統的設計與實現

springboot230基于Spring Boot在線遠程考試系統的設計與實現

在線遠程考試系統設計與實現 摘 要 信息數據從傳統到當代&#xff0c;是一直在變革當中&#xff0c;突如其來的互聯網讓傳統的信息管理看到了革命性的曙光&#xff0c;因為傳統信息管理從時效性&#xff0c;還是安全性&#xff0c;還是可操作性等各個方面來講&#xff0c;遇到…
閱讀更多...
后端知識（理解背誦）

后端知識（理解背誦）

文章目錄 &#x1f37a; 來源&#x1f37a; C&#x1f37b; new 和 malloc 的區別&#xff1f;2&#x1f37b; delete 和 delete[] 的區別&#xff1f;0&#x1f37b; 內存泄漏是什么&#xff1f;如何避免&#xff1f;1 &#x1f37a; 計算機網絡&#x1f37b; URL 輸入后發生了…
閱讀更多...
【Day62】代碼隨想錄之單調棧_739. 每日溫度_496.下一個更大元素 I

【Day62】代碼隨想錄之單調棧_739. 每日溫度_496.下一個更大元素 I

文章目錄 1. 739. 每日溫度2. 496.下一個更大元素 I 1. 739. 每日溫度 參考文檔&#xff1a;代碼隨想錄 分析&#xff1a; 找下一個更高溫度出現在幾天后&#xff0c;即當前位置右側出現的一個比它更大的值&#xff0c;如果是暴力搜索&#xff0c;兩層for&#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 兩階…
閱讀更多...
代碼隨想錄算法訓練營第四五天 | dp[j] = min(dp[j], dp[j - coins[i]] + 1)

代碼隨想錄算法訓練營第四五天 | dp[j] = min(dp[j], dp[j - coins[i]] + 1)

目錄 爬樓梯 &#xff08;進階&#xff09;零錢兌換完全平方數總結 LeetCode 70. 爬樓梯 &#xff08;進階&#xff09; LeetCode 322. 零錢兌換 LeetCode 279.完全平方數 爬樓梯 &#xff08;進階&#xff09; 好做 import java.util.*;public class Main{// dp[i] 爬到有…
閱讀更多...
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提交或回滾事務消息 生產者發送事務消息和執行本地事務。 發送過程分為兩個階段: 第一階段,發送事務消息 第二階段,發…
閱讀更多...
Spring Expression Language (SpEL)

Spring Expression Language (SpEL)

Spring 表達語言&#xff08;SpEL&#xff09;&#xff0c;支持在運行時查詢和操作對象圖&#xff0c;可以用于數據綁定、屬性訪問、方法調用等。使用SpEL可以簡化代碼并提高應用程序的可維護性。 1 概覽 SpelExpressionParser是SpEL的一個核心組件&#xff0c;負責解析和編譯…
閱讀更多...
CentOS安裝編譯Python3.11.6

CentOS安裝編譯Python3.11.6

CentOs自帶python2版本太低&#xff0c;項目需要python3&#xff0c;于是自己安裝python 操作指南&#xff1a; 重新下載源代碼&#xff1a; # 刪除舊的 Python 源代碼文件&#xff08;如果有&#xff09; rm -rf Python-3.11.6.tar.xz # 下載 Python 3.11.6 的源代碼文件 wget…
閱讀更多...
Java泛型簡介

Java泛型簡介

Java泛型簡介 Java泛型是在Java 5中引入的一個特性&#xff0c;它允許程序員在編譯時指定類、接口或方法能夠接受的類型。泛型的主要目的是提供編譯時類型安全檢查&#xff0c;避免在運行時因為類型轉換錯誤而導致的ClassCastException。 在沒有泛型之前&#xff0c;Java中的集…
閱讀更多...
如何利用動態靜態代理IP實現跨地域網絡營銷與市場研究

如何利用動態靜態代理IP實現跨地域網絡營銷與市場研究

動態代理IP和靜態代理IP都可以在跨地域網絡營銷與市場研究中發揮關鍵作用&#xff0c;具體實現方式如下&#xff1a; ### 動態代理IP的應用&#xff1a; 1. 跨地域營銷活動測試&#xff1a; - 在進行網絡營銷時&#xff0c;尤其是要驗證廣告投放、SEO效果或A/B測試不同地區用戶…
閱讀更多...
最新文章

Copyright @ 2022~2023