Skip to content

Commit 749786d

Browse files
ahueckahueck
committed
Checked GPU allocations and wrapper parser extension
1 parent 0ea8e83 commit 749786d

17 files changed

Lines changed: 147 additions & 65 deletions

‎README.md‎

Lines changed: 11 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -164,6 +164,7 @@ Specifically, `TYPEART_OPTIONS` can globally modify the TypeART pass (stack/heap
164164
| `TYPEART_STACK` | `stack` | `false` | Instrument stack and global allocations. Enables instrumentation of global allocations. |
165165
| `TYPEART_STACK_LIFETIME` | `stack-lifetime` | `true` | Instrument stack `llvm.lifetime.start` instead of `alloca` directly |
166166
| `TYPEART_GLOBAL` | `global` | `false` | Instrument global allocations (see stack). |
167+
| `TYPEART_GPU` | `gpu` | `false` | Instrument GPU allocation/free instructions (HIP and CUDA). |
167168
| `TYPEART_TYPEGEN` | `typegen` | `dimeta` | Values: `dimeta`, `ir`. How serializing of type information is done, see [Section 2.2](#22-serialized-type-information). |
168169
| `TYPEART_TYPE_SERIALIZATION` | `type-serialization` | `hybrid` | Values: `file`, `hybrid`, `inline`. How type information are stored (in the executable or externally), see [Section 2.2](#22-serialized-type-information). |
169170
| `TYPEART_STATS` | `stats` | `false` | Show instrumentation statistic counters |
@@ -185,6 +186,11 @@ Additionally, there are two debug environment flags for dumping the LLVM IR per
185186

186187
<!--- @formatter:on --->
187188

189+
#### 2.1.1 Passing options via compiler wrapper
190+
191+
The compiler wrappers support passing TypeART options directly via the command line using the `--typeart-<option name>=<value>` syntax. For boolean flags, no assignment value is required to enable them. These options are transformed into the corresponding `TYPEART_<OPTION>=<value>` environment variables by the wrapper. For example, invoking, e.g., `typeart-clang` with `--typeart-gpu` is equivalent to setting `TYPEART_GPU=true` in the environment.
192+
193+
188194

189195
### 2.2 Serialized type information
190196

@@ -284,7 +290,7 @@ void foo() {
284290

285291
## 3. Building TypeART
286292

287-
TypeART supports LLVM version 14, 18-21, and CMake version >= 3.20.
293+
TypeART supports LLVM version 14, 18-22, and CMake version >= 3.20.
288294

289295
### 3.1 Optional software requirements
290296

@@ -316,10 +322,10 @@ $> cmake --build build --target install --parallel
316322

317323
<!--- @formatter:off --->
318324

319-
| Option | Default | Description |
320-
|------------------------------|:-------:|----------------------------------------------------------------------------------|
321-
| `TYPEART_MPI_WRAPPER` | `ON` | Install TypeART MPI wrapper (mpic, mpic++). Requires MPI. |
322-
| `TYPEART_USE_LEGACY_WRAPPER` | `OFF` | Use legacy wrapper invoking opt/llc directly instead of Clang's `-fpass-plugin`. |
325+
| Option | Default | Description |
326+
|------------------------------|:-------:|-----------------------------------------------------------------------------------------------|
327+
| `TYPEART_MPI_WRAPPER` | `ON` | Install TypeART MPI wrapper (mpic, mpic++). Requires MPI. |
328+
| `TYPEART_USE_LEGACY_WRAPPER` | `OFF` | (Deprecated) Use legacy wrapper invoking opt/llc directly instead of Clang's `-fpass-plugin`. |
323329
324330
325331
<!--- @formatter:on --->

‎lib/passes/instrumentation/MemOpInstrumentation.cpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,7 @@
4141
#include "llvm/IR/Type.h"
4242
#include "llvm/Support/Casting.h"
4343
#include "llvm/Support/raw_ostream.h"
44+
#include "llvm/Transforms/Utils/BasicBlockUtils.h"
4445
#include "llvm/Transforms/Utils/ModuleUtils.h"
4546

4647
#include <llvm/IR/InstrTypes.h>
@@ -152,6 +153,10 @@ InstrCount MemOpInstrumentation::instrumentHeap(const HeapArgList& heap) {
152153
case MemOpKind::CudaMallocLike:
153154
[[fallthrough]];
154155
case MemOpKind::HipMallocLike: {
156+
auto* is_success = IRB.CreateICmpEQ(malloc_call, llvm::ConstantInt::get(malloc_call->getType(), 0));
157+
auto* then_term = llvm::SplitBlockAndInsertIfThen(is_success, insertBefore, false);
158+
IRB.SetInsertPoint(then_term);
159+
155160
auto* runtime_ptr_type = instrumentation_helper->getTypeFor(IType::ptr);
156161
#if LLVM_VERSION_MAJOR >= 15
157162
auto* loaded_ptr = IRB.CreateLoad(runtime_ptr_type, pointer_value);

‎scripts/typeart-wrapperv2.in‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -117,6 +117,11 @@ function typeart_parse_cmd_line_fn() {
117117
shift
118118
;;
119119
--typeart-*)
120+
case "$typeart_parse_mode" in
121+
ON)
122+
typeart_parse_typeart_option_fn "${arg}=true" || return 1
123+
;;
124+
esac
120125
shift
121126
;;
122127
*)

‎test/cuda/pass/01_cudamalloc.c‎

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2,15 +2,20 @@
22

33
// REQUIRES: cuda
44

5-
// LLVM: call i32 @cudaMalloc(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
5+
// LLVM: [[RET:%[0-9a-z_]+]] = call i32 @cudaMalloc(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
6+
// LLVM-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
7+
// LLVM-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
8+
// LLVM: [[LABEL]]:
69
// LLVM-NEXT: [[CUDA_PTR:%[0-9a-z_]+]] = load {{.*}}, {{.*}}[[CU_POINTER]]
710
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr {{.*}}[[CUDA_PTR]], i32 23, i64 20)
811

9-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
10-
// LLVM_LEGACY: call i32 @cudaMalloc(i8** {{.*}}[[CAST1]],
11-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
12-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
13-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
12+
// LLVM_LEGACY: [[RET:%[0-9a-z_]+]] = call i32 @cudaMalloc(i8** {{.*}}[[CAST1:%[0-9a-z_]+]],
13+
// LLVM_LEGACY-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
14+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
15+
// LLVM_LEGACY: [[LABEL]]:
16+
// LLVM_LEGACY-NEXT: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
17+
// LLVM_LEGACY-NEXT: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
18+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
1419

1520
int main() {
1621
const int N = 20;

‎test/cuda/pass/03_cudahostalloc.c‎

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2,15 +2,20 @@
22

33
// REQUIRES: cuda
44

5-
// LLVM: call i32 @cudaHostAlloc(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
5+
// LLVM: [[RET:%[0-9a-z_]+]] = call i32 @cudaHostAlloc(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
6+
// LLVM-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
7+
// LLVM-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
8+
// LLVM: [[LABEL]]:
69
// LLVM-NEXT: [[CUDA_PTR:%[0-9a-z_]+]] = load ptr, ptr [[CU_POINTER]]
710
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr {{.*}}[[CUDA_PTR]],
811

9-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
10-
// LLVM_LEGACY: call i32 @cudaHostAlloc(i8** {{.*}}[[CAST1]],
11-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
12-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
13-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
12+
// LLVM_LEGACY: [[RET:%[0-9a-z_]+]] = call i32 @cudaHostAlloc(i8** {{.*}}[[CAST1:%[0-9a-z_]+]],
13+
// LLVM_LEGACY-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
14+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
15+
// LLVM_LEGACY: [[LABEL]]:
16+
// LLVM_LEGACY-NEXT: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
17+
// LLVM_LEGACY-NEXT: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
18+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
1419

1520
int main() {
1621
const int N = 20;

‎test/cuda/pass/04_cudamalloc_host.c‎

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2,15 +2,20 @@
22

33
// REQUIRES: cuda
44

5-
// LLVM: call i32 @cudaMallocHost(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
5+
// LLVM: [[RET:%[0-9a-z_]+]] = call i32 @cudaMallocHost(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
6+
// LLVM-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
7+
// LLVM-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
8+
// LLVM: [[LABEL]]:
69
// LLVM-NEXT: [[CUDA_PTR:%[0-9a-z_]+]] = load {{.*}}, {{.*}}[[CU_POINTER]]
710
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr {{.*}}[[CUDA_PTR]], i32 23, i64 20)
811

9-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
10-
// LLVM_LEGACY: call i32 @cudaMallocHost(i8** {{.*}}[[CAST1]],
11-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
12-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
13-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
12+
// LLVM_LEGACY: [[RET:%[0-9a-z_]+]] = call i32 @cudaMallocHost(i8** {{.*}}[[CAST1:%[0-9a-z_]+]],
13+
// LLVM_LEGACY-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
14+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
15+
// LLVM_LEGACY: [[LABEL]]:
16+
// LLVM_LEGACY-NEXT: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
17+
// LLVM_LEGACY-NEXT: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
18+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
1419

1520
int main() {
1621
const int N = 20;

‎test/cuda/pass/05_cudamalloc_managed.c‎

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2,15 +2,20 @@
22

33
// REQUIRES: cuda
44

5-
// LLVM: call i32 @cudaMallocManaged(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
5+
// LLVM: [[RET:%[0-9a-z_]+]] = call i32 @cudaMallocManaged(ptr {{.*}}[[CU_POINTER:%[_0-9a-z]+]],
6+
// LLVM-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
7+
// LLVM-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
8+
// LLVM: [[LABEL]]:
69
// LLVM-NEXT: [[CUDA_PTR:%[0-9a-z_]+]] = load {{.*}}, {{.*}}[[CU_POINTER]]
710
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr {{.*}}[[CUDA_PTR]], i32 23, i64 20)
811

9-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
10-
// LLVM_LEGACY: call i32 @cudaMallocManaged(i8** {{.*}}[[CAST1]],
11-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
12-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
13-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
12+
// LLVM_LEGACY: [[RET:%[0-9a-z_]+]] = call i32 @cudaMallocManaged(i8** {{.*}}[[CAST1:%[0-9a-z_]+]],
13+
// LLVM_LEGACY-NEXT: [[SUCCESS:%[0-9a-z_]+]] = icmp eq i32 [[RET]], 0
14+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS]], label %[[LABEL:[0-9a-z_.]+]], label %[[SKIP:[0-9a-z_.]+]]
15+
// LLVM_LEGACY: [[LABEL]]:
16+
// LLVM_LEGACY-NEXT: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
17+
// LLVM_LEGACY-NEXT: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
18+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
1419

1520
int main() {
1621
const int N = 20;

‎test/cuda/pass/07_cudamalloc_async.c‎

Lines changed: 23 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -3,25 +3,35 @@
33
// REQUIRES: cuda
44

55
// clang-format off
6-
// LLVM: call i32 @cudaMallocAsync(ptr {{.*}}[[CU_POINTER_X:%[_0-9a-z]+]], i64{{.*}} 80, ptr {{.*}})
6+
// LLVM: [[RET_X:%[0-9a-z_]+]] = call i32 @cudaMallocAsync(ptr {{.*}}[[CU_POINTER_X:%[_0-9a-z]+]], i64{{.*}} 80, ptr {{.*}})
7+
// LLVM-NEXT: [[SUCCESS_X:%[0-9a-z_]+]] = icmp eq i32 [[RET_X]], 0
8+
// LLVM-NEXT: br i1 [[SUCCESS_X]], label %[[LABEL_X:[0-9a-z_.]+]], label %[[SKIP_X:[0-9a-z_.]+]]
9+
// LLVM: [[LABEL_X]]:
710
// LLVM-NEXT: [[CUDA_PTR_X:%[0-9a-z_]+]] = load ptr, ptr [[CU_POINTER_X]]
811
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr [[CUDA_PTR_X]], i32 23, i64 20)
912

10-
// LLVM: call i32 @cudaMallocFromPoolAsync(ptr {{.*}}[[CU_POINTER_Y:%[_0-9a-z]+]], i64{{.*}} 80, ptr {{.*}}, ptr {{.*}})
13+
// LLVM: [[RET_Y:%[0-9a-z_]+]] = call i32 @cudaMallocFromPoolAsync(ptr {{.*}}[[CU_POINTER_Y:%[_0-9a-z]+]], i64{{.*}} 80, ptr {{.*}}, ptr {{.*}})
14+
// LLVM-NEXT: [[SUCCESS_Y:%[0-9a-z_]+]] = icmp eq i32 [[RET_Y]], 0
15+
// LLVM-NEXT: br i1 [[SUCCESS_Y]], label %[[LABEL_Y:[0-9a-z_.]+]], label %[[SKIP_Y:[0-9a-z_.]+]]
16+
// LLVM: [[LABEL_Y]]:
1117
// LLVM-NEXT: [[CUDA_PTR_Y:%[0-9a-z_]+]] = load ptr, ptr [[CU_POINTER_Y]]
1218
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr [[CUDA_PTR_Y]], i32 23, i64 20)
1319

14-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
15-
// LLVM_LEGACY: call i32 @cudaMallocAsync(i8** {{.*}}[[CAST1]],
16-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
17-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
18-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
19-
20-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
21-
// LLVM_LEGACY: call i32 @cudaMallocFromPoolAsync(i8** {{.*}}[[CAST1]],
22-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
23-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
24-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
20+
// LLVM_LEGACY: [[RET_X:%[0-9a-z_]+]] = call i32 @cudaMallocAsync(i8** {{.*}}[[CAST1_X:%[0-9a-z_]+]],
21+
// LLVM_LEGACY-NEXT: [[SUCCESS_X:%[0-9a-z_]+]] = icmp eq i32 [[RET_X]], 0
22+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS_X]], label %[[LABEL_X:[0-9a-z_.]+]], label %[[SKIP_X:[0-9a-z_.]+]]
23+
// LLVM_LEGACY: [[LABEL_X]]:
24+
// LLVM_LEGACY-NEXT: [[CAST2_X:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR_X:%[0-9a-zA-Z_]+]] to i8**
25+
// LLVM_LEGACY-NEXT: [[LOADED_PTR_X:%[0-9a-z_]+]] = load i8*, i8** [[CAST2_X]]
26+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR_X]], i32 23, i64 20)
27+
28+
// LLVM_LEGACY: [[RET_Y:%[0-9a-z_]+]] = call i32 @cudaMallocFromPoolAsync(i8** {{.*}}[[CAST1_Y:%[0-9a-z_]+]],
29+
// LLVM_LEGACY-NEXT: [[SUCCESS_Y:%[0-9a-z_]+]] = icmp eq i32 [[RET_Y]], 0
30+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS_Y]], label %[[LABEL_Y:[0-9a-z_.]+]], label %[[SKIP_Y:[0-9a-z_.]+]]
31+
// LLVM_LEGACY: [[LABEL_Y]]:
32+
// LLVM_LEGACY-NEXT: [[CAST2_Y:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR_Y:%[0-9a-zA-Z_]+]] to i8**
33+
// LLVM_LEGACY-NEXT: [[LOADED_PTR_Y:%[0-9a-z_]+]] = load i8*, i8** [[CAST2_Y]]
34+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR_Y]], i32 23, i64 20)
2535
// clang-format on
2636

2737
int main() {

‎test/cuda/pass/09_cudamalloc_async.c‎

Lines changed: 23 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -3,25 +3,35 @@
33
// REQUIRES: cuda
44

55
// clang-format off
6-
// LLVM: call i32 @cudaMallocAsync(ptr {{.*}}[[CU_POINTER_X:%[_0-9a-z]+]], i64 {{.*}}80, ptr {{.*}})
6+
// LLVM: [[RET_X:%[0-9a-z_]+]] = call i32 @cudaMallocAsync(ptr {{.*}}[[CU_POINTER_X:%[_0-9a-z]+]], i64 {{.*}}80, ptr {{.*}})
7+
// LLVM-NEXT: [[SUCCESS_X:%[0-9a-z_]+]] = icmp eq i32 [[RET_X]], 0
8+
// LLVM-NEXT: br i1 [[SUCCESS_X]], label %[[LABEL_X:[0-9a-z_.]+]], label %[[SKIP_X:[0-9a-z_.]+]]
9+
// LLVM: [[LABEL_X]]:
710
// LLVM-NEXT: [[CUDA_PTR_X:%[0-9a-z_]+]] = load ptr, ptr [[CU_POINTER_X]]
811
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr [[CUDA_PTR_X]], i32 23, i64 20)
912

10-
// LLVM: call i32 @cudaMallocFromPoolAsync(ptr {{.*}}[[CU_POINTER_Y:%[_0-9a-z]+]], i64 {{.*}}80, ptr {{.*}}, ptr {{.*}})
13+
// LLVM: [[RET_Y:%[0-9a-z_]+]] = call i32 @cudaMallocFromPoolAsync(ptr {{.*}}[[CU_POINTER_Y:%[_0-9a-z]+]], i64 {{.*}}80, ptr {{.*}}, ptr {{.*}})
14+
// LLVM-NEXT: [[SUCCESS_Y:%[0-9a-z_]+]] = icmp eq i32 [[RET_Y]], 0
15+
// LLVM-NEXT: br i1 [[SUCCESS_Y]], label %[[LABEL_Y:[0-9a-z_.]+]], label %[[SKIP_Y:[0-9a-z_.]+]]
16+
// LLVM: [[LABEL_Y]]:
1117
// LLVM-NEXT: [[CUDA_PTR_Y:%[0-9a-z_]+]] = load ptr, ptr [[CU_POINTER_Y]]
1218
// LLVM-NEXT: call void @__typeart_alloc_gpu(ptr [[CUDA_PTR_Y]], i32 23, i64 20)
1319

14-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
15-
// LLVM_LEGACY: call i32 @cudaMallocAsync(i8** {{.*}}[[CU_POINTER_X:%[_0-9a-z]+]], i64 {{.*}}80,
16-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
17-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
18-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
19-
20-
// LLVM_LEGACY: [[CAST1:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR:%[0-9a-zA-Z_]+]] to i8**
21-
// LLVM_LEGACY: call i32 @cudaMallocFromPoolAsync(i8** {{.*}}[[CU_POINTER_Y:%[_0-9a-z]+]], i64 {{.*}}80,
22-
// LLVM_LEGACY: [[CAST2:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR]] to i8**
23-
// LLVM_LEGACY: [[LOADED_PTR:%[0-9a-z_]+]] = load i8*, i8** [[CAST2]]
24-
// LLVM_LEGACY: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR]], i32 23, i64 20)
20+
// LLVM_LEGACY: [[RET_X:%[0-9a-z_]+]] = call i32 @cudaMallocAsync(i8** {{.*}}[[CU_POINTER_X:%[_0-9a-z]+]], i64 {{.*}}80,
21+
// LLVM_LEGACY-NEXT: [[SUCCESS_X:%[0-9a-z_]+]] = icmp eq i32 [[RET_X]], 0
22+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS_X]], label %[[LABEL_X:[0-9a-z_.]+]], label %[[SKIP_X:[0-9a-z_.]+]]
23+
// LLVM_LEGACY: [[LABEL_X]]:
24+
// LLVM_LEGACY-NEXT: [[CAST2_X:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR_X:%[0-9a-zA-Z_]+]] to i8**
25+
// LLVM_LEGACY-NEXT: [[LOADED_PTR_X:%[0-9a-z_]+]] = load i8*, i8** [[CAST2_X]]
26+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR_X]], i32 23, i64 20)
27+
28+
// LLVM_LEGACY: [[RET_Y:%[0-9a-z_]+]] = call i32 @cudaMallocFromPoolAsync(i8** {{.*}}[[CU_POINTER_Y:%[_0-9a-z]+]], i64 {{.*}}80,
29+
// LLVM_LEGACY-NEXT: [[SUCCESS_Y:%[0-9a-z_]+]] = icmp eq i32 [[RET_Y]], 0
30+
// LLVM_LEGACY-NEXT: br i1 [[SUCCESS_Y]], label %[[LABEL_Y:[0-9a-z_.]+]], label %[[SKIP_Y:[0-9a-z_.]+]]
31+
// LLVM_LEGACY: [[LABEL_Y]]:
32+
// LLVM_LEGACY-NEXT: [[CAST2_Y:%[0-9a-z_]+]] = bitcast float** [[SRC_VAR_Y:%[0-9a-zA-Z_]+]] to i8**
33+
// LLVM_LEGACY-NEXT: [[LOADED_PTR_Y:%[0-9a-z_]+]] = load i8*, i8** [[CAST2_Y]]
34+
// LLVM_LEGACY-NEXT: call void @__typeart_alloc_gpu(i8* [[LOADED_PTR_Y]], i32 23, i64 20)
2535
// clang-format on
2636

2737
int main() {

‎test/cuda/pass/11_wrapper_gpu_wrapper_option.c‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// RUN: %wrapper-cc -x cuda --cuda-host-only -nocudalib -S -emit-llvm -O1 --typeart-gpu=true %s -o - 2>&1 | %filecheck %s --check-prefix=GPU-ON
33
// RUN: %wrapper-cc -x cuda --cuda-host-only -nocudalib -S -emit-llvm -O1 --typeart-gpu=false %s -o - 2>&1 | %filecheck %s --check-prefix=GPU-OFF
44
// RUN: %wrapper-cc -x cuda --cuda-host-only -nocudalib -S -emit-llvm -O1 --typeart-gpu --typeart-gpu=true %s -o - 2>&1 | %filecheck %s --check-prefix=GPU-ON
5-
// RUN: %wrapper-cc -x cuda --cuda-host-only -nocudalib -S -emit-llvm -O1 --typeart-gpu %s -o - 2>&1 | %filecheck %s --check-prefix=GPU-OFF
5+
// RUN: %wrapper-cc -x cuda --cuda-host-only -nocudalib -S -emit-llvm -O1 --typeart-gpu %s -o - 2>&1 | %filecheck %s --check-prefix=GPU-ON
66
// RUN: TYPEART_WRAPPER=OFF %wrapper-cc -x cuda --cuda-host-only -nocudalib -S -emit-llvm -O1 --typeart-gpu=true --typeart-gpu %s -o - 2>&1 | %filecheck %s --check-prefix=WRAPPER-OFF
77
// clang-format on
88

0 commit comments

Comments
 (0)