Skip to content

Commit 69f2099

Browse files
bcardosolopesbcardosolopes
committed
[CIR] Fix CUDA tests and comment some unsupported bits
1 parent 34f8c17 commit 69f2099

File tree

4 files changed

+72
-74
lines changed

4 files changed

+72
-74
lines changed

‎clang/test/CIR/CodeGen/CUDA/builtin-functions.cu

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,3 @@
1-
// XFAIL: *
21
#include "../Inputs/cuda.h"
32

43
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
@@ -59,7 +58,8 @@ __device__ void builtins() {
5958
// CIR: cir.llvm.intrinsic "nvvm.membar.sys"
6059
// LLVM: call void @llvm.nvvm.membar.sys()
6160
__nvvm_membar_sys();
62-
// CIR: cir.llvm.intrinsic "nvvm.barrier0"
63-
// LLVM: call void @llvm.nvvm.barrier0()
64-
__syncthreads();
61+
62+
// TODO-CIR: cir.llvm.intrinsic "nvvm.barrier0"
63+
// TODO-LLVM: call void @llvm.nvvm.barrier0()
64+
// __syncthreads();
6565
}

‎clang/test/CIR/CodeGen/CUDA/cuda-builtin-vars.cu

Lines changed: 55 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,3 @@
1-
// XFAIL: *
2-
31
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
42
// RUN: -fcuda-is-device -emit-llvm -o - %s \
53
// RUN: | FileCheck --check-prefix=LLVM %s
@@ -11,80 +9,81 @@
119
#include "__clang_cuda_builtin_vars.h"
1210

1311
// LLVM: define{{.*}} void @_Z6kernelPi(ptr %0)
12+
// CIR-LABEL: @_Z6kernelPi
1413
__attribute__((global))
1514
void kernel(int *out) {
1615
int i = 0;
1716

18-
out[i++] = threadIdx.x;
19-
// CIR: cir.func linkonce_odr @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_xEv()
20-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.tid.x"
21-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.tid.x()
17+
// out[i++] = threadIdx.x;
18+
// CIR-DISABLED: cir.func linkonce_odr @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_xEv()
19+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.tid.x"
20+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.tid.x()
2221

23-
out[i++] = threadIdx.y;
24-
// CIR: cir.func linkonce_odr @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_yEv()
25-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.tid.y"
26-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.tid.y()
22+
// out[i++] = threadIdx.y;
23+
// CIR-DISABLED: cir.func linkonce_odr @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_yEv()
24+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.tid.y"
25+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.tid.y()
2726

28-
out[i++] = threadIdx.z;
29-
// CIR: cir.func linkonce_odr @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_zEv()
30-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.tid.z"
31-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.tid.z()
27+
// out[i++] = threadIdx.z;
28+
// CIR-DISABLED: cir.func linkonce_odr @_ZN26__cuda_builtin_threadIdx_t17__fetch_builtin_zEv()
29+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.tid.z"
30+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.tid.z()
3231

3332

34-
out[i++] = blockIdx.x;
35-
// CIR: cir.func linkonce_odr @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_xEv()
36-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ctaid.x"
37-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
33+
// out[i++] = blockIdx.x;
34+
// CIR-DISABLED: cir.func linkonce_odr @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_xEv()
35+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ctaid.x"
36+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
3837

39-
out[i++] = blockIdx.y;
40-
// CIR: cir.func linkonce_odr @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_yEv()
41-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ctaid.y"
42-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
38+
// out[i++] = blockIdx.y;
39+
// CIR-DISABLED: cir.func linkonce_odr @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_yEv()
40+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ctaid.y"
41+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
4342

44-
out[i++] = blockIdx.z;
45-
// CIR: cir.func linkonce_odr @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_zEv()
46-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ctaid.z"
47-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
43+
// out[i++] = blockIdx.z;
44+
// CIR-DISABLED: cir.func linkonce_odr @_ZN25__cuda_builtin_blockIdx_t17__fetch_builtin_zEv()
45+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ctaid.z"
46+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
4847

4948

50-
out[i++] = blockDim.x;
51-
// CIR: cir.func linkonce_odr @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_xEv()
52-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ntid.x"
53-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
49+
// out[i++] = blockDim.x;
50+
// CIR-DISABLED: cir.func linkonce_odr @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_xEv()
51+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ntid.x"
52+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
5453

55-
out[i++] = blockDim.y;
56-
// CIR: cir.func linkonce_odr @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_yEv()
57-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ntid.y"
58-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
54+
// out[i++] = blockDim.y;
55+
// CIR-DISABLED: cir.func linkonce_odr @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_yEv()
56+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ntid.y"
57+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
5958

60-
out[i++] = blockDim.z;
61-
// CIR: cir.func linkonce_odr @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_zEv()
62-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ntid.z"
63-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
59+
// out[i++] = blockDim.z;
60+
// CIR-DISABLED: cir.func linkonce_odr @_ZN25__cuda_builtin_blockDim_t17__fetch_builtin_zEv()
61+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.ntid.z"
62+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
6463

6564

66-
out[i++] = gridDim.x;
67-
// CIR: cir.func linkonce_odr @_ZN24__cuda_builtin_gridDim_t17__fetch_builtin_xEv()
68-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.nctaid.x"
69-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
65+
// out[i++] = gridDim.x;
66+
// CIR-DISABLED: cir.func linkonce_odr @_ZN24__cuda_builtin_gridDim_t17__fetch_builtin_xEv()
67+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.nctaid.x"
68+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
7069

71-
out[i++] = gridDim.y;
72-
// CIR: cir.func linkonce_odr @_ZN24__cuda_builtin_gridDim_t17__fetch_builtin_yEv()
73-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.nctaid.y"
74-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
70+
// out[i++] = gridDim.y;
71+
// CIR-DISABLED: cir.func linkonce_odr @_ZN24__cuda_builtin_gridDim_t17__fetch_builtin_yEv()
72+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.nctaid.y"
73+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
7574

76-
out[i++] = gridDim.z;
77-
// CIR: cir.func linkonce_odr @_ZN24__cuda_builtin_gridDim_t17__fetch_builtin_zEv()
78-
// CIR: cir.llvm.intrinsic "nvvm.read.ptx.sreg.nctaid.z"
79-
// LLVM: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
75+
// out[i++] = gridDim.z;
76+
// CIR-DISABLED: cir.func linkonce_odr @_ZN24__cuda_builtin_gridDim_t17__fetch_builtin_zEv()
77+
// CIR-DISABLED: cir.llvm.intrinsic "nvvm.read.ptx.sreg.nctaid.z"
78+
// LLVM-DISABLED: call{{.*}} i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
8079

8180

82-
out[i++] = warpSize;
83-
// CIR: [[REGISTER:%.*]] = cir.const #cir.int<32>
84-
// CIR: cir.store{{.*}} [[REGISTER]]
85-
// LLVM: store i32 32,
81+
// out[i++] = warpSize;
82+
// CIR-DISABLED: [[REGISTER:%.*]] = cir.const #cir.int<32>
83+
// CIR-DISABLED: cir.store{{.*}} [[REGISTER]]
84+
// LLVM-DISABLED: store i32 32,
8685

8786

88-
// CIR: cir.return loc
89-
// LLVM: ret void
87+
// CIR-DISABLED: cir.return loc
88+
// LLVM-DISABLED: ret void
9089
}

‎clang/test/CIR/CodeGen/CUDA/destructor.cu

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,5 @@
11
#include "../Inputs/cuda.h"
22

3-
// XFAIL: *
4-
53
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
64
// RUN: -x cuda -emit-cir -target-sdk-version=12.3 \
75
// RUN: %s -o %t.cir
@@ -19,11 +17,13 @@ template<typename T> struct A {
1917
~A() { f<<<1, 1>>>(T()); }
2018
};
2119

22-
// CIR-DEVICE: cir.func dso_local @_Z1fIiEvT_
20+
// CIR-HOST: module
21+
// CIR-DEVICE: module
22+
// CIR-DEVICE-DISABLED: cir.func dso_local @_Z1fIiEvT_
2323

24-
// CIR-HOST: cir.func {{.*}} @_ZN1AIiED2Ev{{.*}} {
25-
// CIR-HOST: cir.call @__cudaPushCallConfiguration
26-
// CIR-HOST: cir.call @_Z16__device_stub__fIiEvT_
27-
// CIR-HOST: }
24+
// CIR-HOST-DISABLED: cir.func {{.*}} @_ZN1AIiED2Ev{{.*}} {
25+
// CIR-HOST-DISABLED: cir.call @__cudaPushCallConfiguration
26+
// CIR-HOST-DISABLED: cir.call @_Z16__device_stub__fIiEvT_
27+
// CIR-HOST-DISABLED: }
2828

29-
A<int> a;
29+
// A<int> a;

‎clang/test/CIR/CodeGen/CUDA/simple.cu

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,10 @@
11
#include "../Inputs/cuda.h"
22

3-
// XFAIL: *
4-
5-
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
6-
// RUN: -x cuda -emit-cir -target-sdk-version=12.3 \
7-
// RUN: %s -o %t.cir
8-
// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
3+
// TODO: host build is currently crashing.
4+
// RUN-DISABLE: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
5+
// RUN-DISABLE: -x cuda -emit-cir -target-sdk-version=12.3 \
6+
// RUN-DISABLE: %s -o %t.cir
7+
// RUN-DISABLE: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
98

109
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
1110
// RUN: -fcuda-is-device -emit-cir -target-sdk-version=12.3 \

0 commit comments

Comments
 (0)