[clang] e88d53d - [HIP] Generate offloading entries for HIP with the new driver.
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Mon Jul 11 12:49:32 PDT 2022
Author: Joseph Huber
Date: 2022-07-11T15:49:21-04:00
New Revision: e88d53d25f3b48204615195615da35e575081903
URL: https://github.com/llvm/llvm-project/commit/e88d53d25f3b48204615195615da35e575081903
DIFF: https://github.com/llvm/llvm-project/commit/e88d53d25f3b48204615195615da35e575081903.diff
LOG: [HIP] Generate offloading entries for HIP with the new driver.
This patch adds the small change required to output offloading entried
for HIP instead of CUDA. These should be placed in different sections so
because they need to be distinct to the offloading toolchain, otherwise
we'd have HIP trying to register CUDA kernels or vice-versa. This patch will
precede support for HIP in the linker wrapper.
Reviewed By: yaxunl, tra
Differential Revision: https://reviews.llvm.org/D128850
Added:
Modified:
clang/lib/CodeGen/CGCUDANV.cpp
clang/test/CodeGenCUDA/offloading-entries.cu
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 6f2679cb15e4c..6a185c29d3cec 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -1116,7 +1116,8 @@ void CGNVCUDARuntime::createOffloadingEntries() {
llvm::OpenMPIRBuilder OMPBuilder(CGM.getModule());
OMPBuilder.initialize();
- StringRef Section = "cuda_offloading_entries";
+ StringRef Section = CGM.getLangOpts().HIP ? "hip_offloading_entries"
+ : "cuda_offloading_entries";
for (KernelInfo &I : EmittedKernels)
OMPBuilder.emitOffloadingEntry(KernelHandles[I.Kernel],
getDeviceSideName(cast<NamedDecl>(I.D)), 0,
diff --git a/clang/test/CodeGenCUDA/offloading-entries.cu b/clang/test/CodeGenCUDA/offloading-entries.cu
index f243028d84267..544dcec45dcb3 100644
--- a/clang/test/CodeGenCUDA/offloading-entries.cu
+++ b/clang/test/CodeGenCUDA/offloading-entries.cu
@@ -1,33 +1,57 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex ".omp_offloading.entry.*"
// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu \
// RUN: --offload-new-driver -emit-llvm -o - -x cuda %s | FileCheck \
-// RUN: --check-prefix=HOST %s
+// RUN: --check-prefix=CUDA %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu \
+// RUN: --offload-new-driver -emit-llvm -o - -x hip %s | FileCheck \
+// RUN: --check-prefix=HIP %s
#include "Inputs/cuda.h"
//.
-// HOST: @x = internal global i32 undef, align 4
-// HOST: @.omp_offloading.entry_name = internal unnamed_addr constant [8 x i8] c"_Z3foov\00"
-// HOST: @.omp_offloading.entry._Z3foov = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__foov, ptr @.omp_offloading.entry_name, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
-// HOST: @.omp_offloading.entry_name.1 = internal unnamed_addr constant [8 x i8] c"_Z3barv\00"
-// HOST: @.omp_offloading.entry._Z3barv = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__barv, ptr @.omp_offloading.entry_name.1, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
-// HOST: @.omp_offloading.entry_name.2 = internal unnamed_addr constant [2 x i8] c"x\00"
-// HOST: @.omp_offloading.entry.x = weak constant %struct.__tgt_offload_entry { ptr @x, ptr @.omp_offloading.entry_name.2, i64 4, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
+// CUDA: @.omp_offloading.entry_name = internal unnamed_addr constant [8 x i8] c"_Z3foov\00"
+// CUDA: @.omp_offloading.entry._Z3foov = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__foov, ptr @.omp_offloading.entry_name, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
+// CUDA: @.omp_offloading.entry_name.1 = internal unnamed_addr constant [8 x i8] c"_Z3barv\00"
+// CUDA: @.omp_offloading.entry._Z3barv = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__barv, ptr @.omp_offloading.entry_name.1, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
+// CUDA: @.omp_offloading.entry_name.2 = internal unnamed_addr constant [2 x i8] c"x\00"
+// CUDA: @.omp_offloading.entry.x = weak constant %struct.__tgt_offload_entry { ptr @x, ptr @.omp_offloading.entry_name.2, i64 4, i32 0, i32 0 }, section "cuda_offloading_entries", align 1
+//.
+// HIP: @.omp_offloading.entry_name = internal unnamed_addr constant [8 x i8] c"_Z3foov\00"
+// HIP: @.omp_offloading.entry._Z3foov = weak constant %struct.__tgt_offload_entry { ptr @_Z3foov, ptr @.omp_offloading.entry_name, i64 0, i32 0, i32 0 }, section "hip_offloading_entries", align 1
+// HIP: @.omp_offloading.entry_name.1 = internal unnamed_addr constant [8 x i8] c"_Z3barv\00"
+// HIP: @.omp_offloading.entry._Z3barv = weak constant %struct.__tgt_offload_entry { ptr @_Z3barv, ptr @.omp_offloading.entry_name.1, i64 0, i32 0, i32 0 }, section "hip_offloading_entries", align 1
+// HIP: @.omp_offloading.entry_name.2 = internal unnamed_addr constant [2 x i8] c"x\00"
+// HIP: @.omp_offloading.entry.x = weak constant %struct.__tgt_offload_entry { ptr @x, ptr @.omp_offloading.entry_name.2, i64 4, i32 0, i32 0 }, section "hip_offloading_entries", align 1
//.
-// HOST-LABEL: @_Z18__device_stub__foov(
-// HOST-NEXT: entry:
-// HOST-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__foov)
-// HOST-NEXT: br label [[SETUP_END:%.*]]
-// HOST: setup.end:
-// HOST-NEXT: ret void
+// CUDA-LABEL: @_Z18__device_stub__foov(
+// CUDA-NEXT: entry:
+// CUDA-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__foov)
+// CUDA-NEXT: br label [[SETUP_END:%.*]]
+// CUDA: setup.end:
+// CUDA-NEXT: ret void
+//
+// HIP-LABEL: @_Z18__device_stub__foov(
+// HIP-NEXT: entry:
+// HIP-NEXT: [[TMP0:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z3foov)
+// HIP-NEXT: br label [[SETUP_END:%.*]]
+// HIP: setup.end:
+// HIP-NEXT: ret void
//
__global__ void foo() {}
-// HOST-LABEL: @_Z18__device_stub__barv(
-// HOST-NEXT: entry:
-// HOST-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__barv)
-// HOST-NEXT: br label [[SETUP_END:%.*]]
-// HOST: setup.end:
-// HOST-NEXT: ret void
+
+// CUDA-LABEL: @_Z18__device_stub__barv(
+// CUDA-NEXT: entry:
+// CUDA-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__barv)
+// CUDA-NEXT: br label [[SETUP_END:%.*]]
+// CUDA: setup.end:
+// CUDA-NEXT: ret void
+//
+// HIP-LABEL: @_Z18__device_stub__barv(
+// HIP-NEXT: entry:
+// HIP-NEXT: [[TMP0:%.*]] = call i32 @hipLaunchByPtr(ptr @_Z3barv)
+// HIP-NEXT: br label [[SETUP_END:%.*]]
+// HIP: setup.end:
+// HIP-NEXT: ret void
//
__global__ void bar() {}
__device__ int x = 1;
More information about the cfe-commits
mailing list