记录一个编译的LLVM 含clang 和 PTX 来支持 HIPIFY 的构建配置

news2024/11/18 1:23:39

llvm 18 debug 版本

build llvmorg-18.1rc4 debug

$ cd llvm-project

$ git checkout llvmorg-18.1.0-rc4

$ mkdir build_d

$ cd build_d

$ mkdir -p ../../local_d

cmake \
-DCMAKE_INSTALL_PREFIX=../../local_d \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_ENABLE_PROJECTS="bolt;clang;clang-tools-extra;lld;mlir"  \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX"  \
-DLLVM_INCLUDE_TESTS=OFF \
-DCMAKE_BUILD_TYPE=Debug \
../llvm

其余部分拆出来了：

cross-project-tests;libclc;lldb;polly;flang

-DLLVM_ENABLE_RUNTIMES="libunwind;libcxxabi;pstl;libcxx;openmp"      \
libc;compiler-rt;

$ make -j34

$make install

llvm 18 release版本

cd llvm-project

mkdir build_r

cd build_r

mkdir -p ../../local_r

cmake \
-DCMAKE_INSTALL_PREFIX=../../local_r \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_ENABLE_PROJECTS="bolt;clang;clang-tools-extra;lld;mlir"  \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX"  \
-DLLVM_INCLUDE_TESTS=OFF \
-DCMAKE_BUILD_TYPE=Release \
../llvm

$ make -j34

效果：

$make install

build HIPIFY debug

$ mkdir /home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/hipify

cmake  \
-DCMAKE_INSTALL_PREFIX=/home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/hipify  \
-DCMAKE_BUILD_TYPE=Debug  \
-DCMAKE_PREFIX_PATH=/home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d  \
..

还有一种更多配置的编译配置方法，其实用不到：

cmake
 -DHIPIFY_CLANG_TESTS=ON \
 -DCMAKE_BUILD_TYPE=Release \
 -DCMAKE_INSTALL_PREFIX=../dist \
 -DCMAKE_PREFIX_PATH=/usr/llvm/17.0.6/dist \
 -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-12.3.2 \
 -DCUDA_DNN_ROOT_DIR=/usr/local/cudnn-8.9.7 \
 -DCUDA_CUB_ROOT_DIR=/usr/local/cub-2.1.0 \
 -DLLVM_EXTERNAL_LIT=/usr/llvm/17.0.6/build/bin/llvm-lit \
 ..

using hipify-clang
 

hipify-clang intro.cu --cuda-path="/usr/local/cuda-12.3" --print-stats-csv

$ /home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/hipify/bin/hipify-clang vectorAdd.cu --cuda-path="/usr/local/cuda-12.3" --clang-resource-directory="/home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/lib/clang/18"

写成Makefile：

EXE := vectorAdd_hip

all: $(EXE)


$(EXE): vectorAdd.cu.hip
	hipcc $< -o $@

%.hip: %
	/home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/hipify/bin/hipify-clang $< --cuda-path=/usr/local/cuda-12.3 --clang-resource-directory=/home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/lib/clang/18

.PHONY: clean
clean:
	${RM} $(EXE) *.hip

效果：

源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;
  }
  if(i==7)printf("Hello kernel threadID=%d\n", i);
}

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;
  }
  if(i==7)printf("Hello kernel threadID=%d\n", i);
}

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;
}

其他参考选项示例：

指示头文件文件夹

./hipify-clang square.cu --cuda-path=/usr/local/cuda-12.3 -I /usr/local/cuda-12.3/samples/common/inc

指示C++标准

./hipify-clang cpp17.cu --cuda-path=/usr/local/cuda-12.3 -- -std=c++17

多个 .cu 文件一起编译

./hipify-clang cpp17.cu ../../square.cu /home/user/cuda/intro.cu --cuda-path=/usr/local/cuda-12.3 -- -std=c++17

统计修改的信息

$ /home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/hipify/bin/hipify-clang vectorAdd.cu --cuda-path=/usr/local/cuda-12.3 --clang-resource-directory=/home/hipper/llvm_3_4_0_ex/browse_llvm_17/local_d/lib/clang/18 --print-stats

将 统计信息存入 .csv文件中

 --print-stats
改成 
 --print-stats-csv

遗留问题

llvmorg-18.1.rc release 配置有问题：

cmake \
-DCMAKE_INSTALL_PREFIX=../../local \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_ENABLE_PROJECTS="bolt;clang;clang-tools-extra;cross-project-tests;libclc;lld;mlir;polly;flang"  \
-DLLVM_ENABLE_RUNTIMES="libc;libunwind;libcxxabi;pstl;libcxx;compiler-rt;openmp"      \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX"  \
-DLLVM_INCLUDE_TESTS=OFF \
-DCMAKE_BUILD_TYPE=Release \
../llvm

lldb;

貌似拿掉 libc 就能行

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

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

相关文章

WorkPlus Meet提供高效、安全视频会议解决方案

WorkPlus Meet提供高效、安全视频会议解决方案

WorkPlus Meet是一款私有部署和定制化的视频会议解决方案&#xff0c;为企业提供高效、安全的远程协作平台。随着全球数字化转型的加速&#xff0c;视频会议已成为企业必不可少的工作工具&#xff0c;而WorkPlus Meet的私有部署和定制化功能&#xff0c;为企业提供了更大的控制…
阅读更多...
【leetcode热题】排序链表

【leetcode热题】排序链表

给你链表的头结点 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&#xff1a; 输入…
阅读更多...
14：00面试，14：10就出来了，问的问题有点变态。。。

14：00面试，14：10就出来了，问的问题有点变态。。。

从小厂出来&#xff0c;没想到在另一家公司又寄了。 到这家公司开始上班&#xff0c;加班是每天必不可少的&#xff0c;看在钱给的比较多的份上&#xff0c;就不太计较了。没想到3月一纸通知&#xff0c;所有人不准加班&#xff0c;加班费不仅没有了&#xff0c;薪资还要降40%,…
阅读更多...
有没有能用蓝牙的游泳耳机？四款好评如潮的游泳耳机速来围观

有没有能用蓝牙的游泳耳机？四款好评如潮的游泳耳机速来围观

在如今的快节奏生活中&#xff0c;音乐已经成为了我们日常休闲和运动时的重要伴侣。尤其对于水上运动爱好者来说&#xff0c;一款好的游泳耳机更是必不可少。它不仅能让我们在水中畅享音乐&#xff0c;还能有效隔绝外界噪音&#xff0c;让我们更加专注于自己的运动。 &#xf…
阅读更多...
鸿蒙开发为什么这么火，现在入行鸿蒙是否来的及？

鸿蒙开发为什么这么火，现在入行鸿蒙是否来的及？

鸿蒙开发是当前备受关注的技术领域之一&#xff0c;对于想要入门学习鸿蒙开发的初学者来说&#xff0c;需要掌握一定的基础知识和技能。鸿蒙开发又是否能为程序员们带来一片光明的未来呢&#xff1f;让我们一同探讨这些问题。 对于初学者来说&#xff0c;鸿蒙开发是否易于上手呢…
阅读更多...
300分钟吃透分布式缓存-22讲：怎么认识和应用Redis内部数据结构？

300分钟吃透分布式缓存-22讲：怎么认识和应用Redis内部数据结构？

Redis 内部数据结构 RdeisDb Redis 中所有数据都保存在 DB 中&#xff0c;一个 Redis 默认最多支持 16 个 DB。Redis 中的每个 DB 都对应一个 redisDb 结构&#xff0c;即每个 Redis 实例&#xff0c;默认有 16 个 redisDb。用户访问时&#xff0c;默认使用的是 0 号 DB&…
阅读更多...
C语言实现贪吃蛇

C语言实现贪吃蛇

前言&#xff1a;今天给大家详细介绍一下小游戏贪吃蛇的代码。 目录 一 .贪吃蛇实现的功能 二.贪吃蛇游戏设计与分析 1.贪吃蛇以及贪吃蛇所需要维护的数据 &#xff08;1&#xff09;贪吃蛇蛇体 &#xff08;2&#xff09;数据维护 2.地图设计 &#xff08;1&#x…
阅读更多...
【三十】springboot项目上高并发解决示例

【三十】springboot项目上高并发解决示例

互相交流入口地址 整体目录&#xff1a; 【一】springboot整合swagger 【二】springboot整合自定义swagger 【三】springboot整合token 【四】springboot整合mybatis-plus 【五】springboot整合mybatis-plus 【六】springboot整合redis 【七】springboot整合AOP实现日志操作 【…
阅读更多...
【Python】一文详细介绍plt.rcParams 在 Matplotlib 中的原理、作用、注意事项

【Python】一文详细介绍plt.rcParams 在 Matplotlib 中的原理、作用、注意事项

【Python】一文详细介绍plt.rcParams 在 Matplotlib 中的原理、作用、注意事项 &#x1f308; 个人主页&#xff1a;高斯小哥 &#x1f525; 高质量专栏&#xff1a;Matplotlib之旅&#xff1a;零基础精通数据可视化、Python基础【高质量合集】、PyTorch零基础入门教程&#x…
阅读更多...
导出谷歌浏览器收藏的网页，并查看网页保存的登录密码

导出谷歌浏览器收藏的网页，并查看网页保存的登录密码

导出谷歌浏览器&#xff08;Chrome&#xff09;收藏的网页&#xff08;书签&#xff09;&#xff1a; 打开谷歌浏览器。在浏览器右上角找到并点击三个垂直排列的小点&#xff08;或称汉堡菜单&#xff09;以打开主菜单。在下拉菜单中选择“书签” > “书签管理器”。在书签…
阅读更多...
NodeJS实现线性查找算法

NodeJS实现线性查找算法

NodeJS实现线性查找算法 以下是使用 Node.js 实现线性搜索算法的示例代码&#xff1a; function linearSearch(arr, target) {for (let i 0; i < arr.length; i) {if (arr[i] target) {return i; // 如果找到目标&#xff0c;返回索引}}return -1; // 如果未找到目标&am…
阅读更多...
Android随手记

Android随手记

activity的生命周期 创建时 onCreate() - onStart() - onResume() - onPause() - onStop() - onDestroy() 切换时 a切换到b a.onCreate() - a.onStart() - a.onResume - a.onPause - b.onCreate() - b.onStart() - b.onResume() - a.onStop() b切换回a b.onPause() - a.onR…
阅读更多...
uniapp:小程序数字键盘功能样式实现

uniapp:小程序数字键盘功能样式实现

代码如下&#xff1a; <template><view><view><view class"money-input"><view class"input-container" click"toggleBox"><view class"input-wrapper"><view class"input-iconone"…
阅读更多...
信息系统项目管理师--成本管理

信息系统项目管理师--成本管理

项⽬成本管理重点关注完成项⽬活动所需资源的成本&#xff0c;但同时也考虑项⽬决策对项⽬产品、服务或成果的使⽤成本、维护成本和⽀持成本的影响。不同的⼲系⼈会在不同的时间&#xff0c;⽤不同的⽅法 测算项⽬成本。 就某些项⽬&#xff0c;特别是⼩项⽬⽽⾔&#xff0c;成…
阅读更多...
鸿蒙操作系统 HarmonyOS 3.2 API 9 Stage模型通过ArkTS接入高德地图

鸿蒙操作系统 HarmonyOS 3.2 API 9 Stage模型通过ArkTS接入高德地图

用鸿蒙ArkTS语言开发地图APP应用时&#xff0c;很多地图厂商只接入了鸿蒙Java&#xff0c;ArkTS版本陆续接入中&#xff0c;等一段时间才能面世&#xff0c;当前使用地图只能通过鸿蒙的Web组件&#xff0c;将HTML页面嵌入到鸿蒙APP中。具体方法如下&#xff1a;编写HTML <!…
阅读更多...
【Linux】Shell及Linux权限

【Linux】Shell及Linux权限

Shell Shell的定义 Shell最简单的定义是&#xff1a;命令行解释器。 Shell的主要任务&#xff1a;1. 将使用者的命令翻译给核心进行处理。2.将核心的处理结果翻译给使用者 为什么要有Shell? 使用者和内核的关系就相当于两个完全陌生的外国人之间的关系&#xff0c;他们要进…
阅读更多...
酒店客房管理系统|基于Springboot的酒店客房管理系统设计与实现（源码+数据库+文档）

酒店客房管理系统|基于Springboot的酒店客房管理系统设计与实现（源码+数据库+文档）

酒店客房管理系统目录 目录 基于Springboot的酒店客房管理系统设计与实现 一、前言 二、系统功能设计 三、系统实现 1、 用户信息管理 2、会员信息管理 3、 客房信息管理 4、收藏客房管理 四、数据库设计 1、实体ER图 五、核心代码 六、论文参考 七、最新计算机…
阅读更多...
【Vue】.sync 修饰符作用

【Vue】.sync 修饰符作用

文章目录 基本用法 基本用法 官方文档是这样介绍的&#xff1a;.sync 修饰符 简单来说就是实现父子组件数据之间的双向绑定&#xff0c;当子组件修改了一个 props 的值时&#xff0c;也会同步到父组件中&#xff0c;实现子组件同步修改父组件&#xff0c;与v-model类似。类别在…
阅读更多...
Redis分段锁，如何设计？

Redis分段锁，如何设计？

问题场景&#xff1a;热点库存扣减问题 秒杀场景&#xff0c;有一个难度的问题&#xff1a;热点库存扣减问题。 既要保证不发生超卖 又要保证高并发 如果解决这个高难度的问题呢&#xff1f; 答案就是使用redis 分段锁。 什么是分布式锁&#xff1f; 一个分布式系统中&am…
阅读更多...
SSM整合项目（使用Vue3 + Element-Plus创建项目基础页面）

SSM整合项目（使用Vue3 + Element-Plus创建项目基础页面）

1.配置Vue启动端口 1.修改vue.config.js const {defineConfig} require(vue/cli-service) module.exports defineConfig({transpileDependencies: true }) module.exports {devServer: {port: 9999 //启动端口} }2.启动 2.安装Element Plus 命令行输入 npm install eleme…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023