Skip to content

Commit afd4578

Browse files
RiverDaveRiverDave
committed
rename AS asm format: clang_address_space -> language_address_space
1 parent d4df570 commit afd4578

File tree

16 files changed

+78
-78
lines changed

16 files changed

+78
-78
lines changed

‎clang/lib/CIR/Dialect/IR/CIRTypes.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -990,14 +990,14 @@ parseAddressSpaceValue(mlir::AsmParser &p,
990990
}
991991

992992
// Try to parse language specific address space.
993-
if (p.parseOptionalKeyword("clang_address_space").succeeded()) {
993+
if (p.parseOptionalKeyword("language_address_space").succeeded()) {
994994
if (p.parseLParen())
995995
return p.emitError(loc, "expected '(' after clang address space");
996996
mlir::FailureOr<cir::LanguageAddressSpace> result =
997997
mlir::FieldParser<cir::LanguageAddressSpace>::parse(p);
998998

999999
if (mlir::failed(result) || p.parseRParen())
1000-
return p.emitError(loc, "expected clang address space keyword");
1000+
return p.emitError(loc, "expected language address space keyword");
10011001

10021002
attr = cir::LanguageAddressSpaceAttr::get(p.getContext(), result.value());
10031003
return mlir::success();
@@ -1012,7 +1012,7 @@ void printAddressSpaceValue(mlir::AsmPrinter &p,
10121012
return;
10131013

10141014
if (auto logical = dyn_cast<cir::LanguageAddressSpaceAttr>(attr)) {
1015-
p << "clang_address_space("
1015+
p << "language_address_space("
10161016
<< cir::stringifyLanguageAddressSpace(logical.getValue()) << ')';
10171017

10181018
return;

‎clang/test/CIR/CodeGen/CUDA/address-spaces.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,9 +11,9 @@ __global__ void fn() {
1111
j = i;
1212
}
1313

14-
// CIR: cir.global "private" internal dso_local clang_address_space(offload_local) @_ZZ2fnvE1j : !s32i
14+
// CIR: cir.global "private" internal dso_local language_address_space(offload_local) @_ZZ2fnvE1j : !s32i
1515
// CIR: cir.func dso_local @_Z2fnv
1616
// CIR: [[Local:%[0-9]+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["i", init]
17-
// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i, clang_address_space(offload_local)>
17+
// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i, language_address_space(offload_local)>
1818
// CIR: [[Tmp:%[0-9]+]] = cir.load {{.*}} [[Local]] : !cir.ptr<!s32i>, !s32i
19-
// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr<!s32i, clang_address_space(offload_local)>
19+
// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr<!s32i, language_address_space(offload_local)>

‎clang/test/CIR/CodeGen/CUDA/global-vars.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -16,15 +16,15 @@
1616
// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
1717

1818
__device__ int a;
19-
// CIR-DEVICE: cir.global external clang_address_space(offload_global) @a = #cir.int<0>
19+
// CIR-DEVICE: cir.global external language_address_space(offload_global) @a = #cir.int<0>
2020
// LLVM-DEVICE: @a = addrspace(1) externally_initialized global i32 0, align 4
2121
// CIR-HOST: {{.*}}cir.global external @a = #cir.undef : !s32i {alignment = 4 : i64, cu.shadow_name = #cir.cu.shadow_name<a>}{{.*}}
2222

2323
__shared__ int shared;
24-
// CIR-DEVICE: cir.global external clang_address_space(offload_local) @shared = #cir.undef
24+
// CIR-DEVICE: cir.global external language_address_space(offload_local) @shared = #cir.undef
2525
// LLVM-DEVICE: @shared = addrspace(3) global i32 undef, align 4
2626

2727
__constant__ int b;
28-
// CIR-DEVICE: cir.global constant external clang_address_space(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
28+
// CIR-DEVICE: cir.global constant external language_address_space(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
2929
// LLVM-DEVICE: @b = addrspace(4) externally_initialized constant i32 0, align 4
3030
// CIR-HOST: {{.*}}cir.global external @b = #cir.undef : !s32i {alignment = 4 : i64, cu.shadow_name = #cir.cu.shadow_name<b>}{{.*}}

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,5 +22,5 @@ struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public
2222
surface<void, 2> surf;
2323

2424
// DEVICE-LLVM: @surf = addrspace(1) externally_initialized global i64 undef, align 4
25-
// DEVICE-CIR: cir.global external clang_address_space(offload_global) @surf = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
25+
// DEVICE-CIR: cir.global external language_address_space(offload_global) @surf = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
2626
// HOST: @surf = global %"struct.surface<void, 2>" zeroinitializer, align 4

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,4 +21,4 @@ struct __attribute__((device_builtin_texture_type)) texture : public textureRefe
2121
texture<float, 2, NormalizedFloat> tex;
2222

2323
// DEVICE-LLVM: @tex = addrspace(1) externally_initialized global i64 undef, align 4
24-
// DEVICE-CIR: cir.global external clang_address_space(offload_global) @tex = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}
24+
// DEVICE-CIR: cir.global external language_address_space(offload_global) @tex = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized}

‎clang/test/CIR/CodeGen/HIP/address-spaces.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,9 +11,9 @@ __global__ void fn() {
1111
j = i;
1212
}
1313

14-
// CIR: cir.global "private" internal dso_local clang_address_space(offload_local) @_ZZ2fnvE1j : !s32i
14+
// CIR: cir.global "private" internal dso_local language_address_space(offload_local) @_ZZ2fnvE1j : !s32i
1515
// CIR: cir.func dso_local @_Z2fnv
1616
// CIR: [[Local:%[0-9]+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["i", init]
17-
// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i, clang_address_space(offload_local)>
17+
// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr<!s32i, language_address_space(offload_local)>
1818
// CIR: [[Tmp:%[0-9]+]] = cir.load {{.*}} [[Local]] : !cir.ptr<!s32i>, !s32i
19-
// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr<!s32i, clang_address_space(offload_local)>
19+
// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr<!s32i, language_address_space(offload_local)>

‎clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -4,30 +4,30 @@
44
// RUN: FileCheck --input-file=%t.ll %s --check-prefix=LLVM
55

66

7-
// CIR: cir.func @func(%arg0: !cir.ptr<!s32i, clang_address_space(offload_local)>
7+
// CIR: cir.func @func(%arg0: !cir.ptr<!s32i, language_address_space(offload_local)>
88
// LLVM: @func(ptr addrspace(3)
99
kernel void func(local int *p) {
10-
// CIR-NEXT: %[[#ALLOCA_P:]] = cir.alloca !cir.ptr<!s32i, clang_address_space(offload_local)>, !cir.ptr<!cir.ptr<!s32i, clang_address_space(offload_local)>, clang_address_space(offload_private)>, ["p", init] {alignment = 8 : i64}
10+
// CIR-NEXT: %[[#ALLOCA_P:]] = cir.alloca !cir.ptr<!s32i, language_address_space(offload_local)>, !cir.ptr<!cir.ptr<!s32i, language_address_space(offload_local)>, language_address_space(offload_private)>, ["p", init] {alignment = 8 : i64}
1111
// LLVM-NEXT: %[[#ALLOCA_P:]] = alloca ptr addrspace(3), i64 1, align 8
1212

1313
int x;
14-
// CIR-NEXT: %[[#ALLOCA_X:]] = cir.alloca !s32i, !cir.ptr<!s32i, clang_address_space(offload_private)>, ["x"] {alignment = 4 : i64}
14+
// CIR-NEXT: %[[#ALLOCA_X:]] = cir.alloca !s32i, !cir.ptr<!s32i, language_address_space(offload_private)>, ["x"] {alignment = 4 : i64}
1515
// LLVM-NEXT: %[[#ALLOCA_X:]] = alloca i32, i64 1, align 4
1616

1717
global char *b;
18-
// CIR-NEXT: %[[#ALLOCA_B:]] = cir.alloca !cir.ptr<!s8i, clang_address_space(offload_global)>, !cir.ptr<!cir.ptr<!s8i, clang_address_space(offload_global)>, clang_address_space(offload_private)>, ["b"] {alignment = 8 : i64}
18+
// CIR-NEXT: %[[#ALLOCA_B:]] = cir.alloca !cir.ptr<!s8i, language_address_space(offload_global)>, !cir.ptr<!cir.ptr<!s8i, language_address_space(offload_global)>, language_address_space(offload_private)>, ["b"] {alignment = 8 : i64}
1919
// LLVM-NEXT: %[[#ALLOCA_B:]] = alloca ptr addrspace(1), i64 1, align 8
2020

2121
private int *ptr;
22-
// CIR-NEXT: %[[#ALLOCA_PTR:]] = cir.alloca !cir.ptr<!s32i, clang_address_space(offload_private)>, !cir.ptr<!cir.ptr<!s32i, clang_address_space(offload_private)>, clang_address_space(offload_private)>, ["ptr"] {alignment = 8 : i64}
22+
// CIR-NEXT: %[[#ALLOCA_PTR:]] = cir.alloca !cir.ptr<!s32i, language_address_space(offload_private)>, !cir.ptr<!cir.ptr<!s32i, language_address_space(offload_private)>, language_address_space(offload_private)>, ["ptr"] {alignment = 8 : i64}
2323
// LLVM-NEXT: %[[#ALLOCA_PTR:]] = alloca ptr, i64 1, align 8
2424

2525
// Store of the argument `p`
26-
// CIR-NEXT: cir.store{{.*}} %arg0, %[[#ALLOCA_P]] : !cir.ptr<!s32i, clang_address_space(offload_local)>, !cir.ptr<!cir.ptr<!s32i, clang_address_space(offload_local)>, clang_address_space(offload_private)>
26+
// CIR-NEXT: cir.store{{.*}} %arg0, %[[#ALLOCA_P]] : !cir.ptr<!s32i, language_address_space(offload_local)>, !cir.ptr<!cir.ptr<!s32i, language_address_space(offload_local)>, language_address_space(offload_private)>
2727
// LLVM-NEXT: store ptr addrspace(3) %{{[0-9]+}}, ptr %[[#ALLOCA_P]], align 8
2828

2929
ptr = &x;
30-
// CIR-NEXT: cir.store{{.*}} %[[#ALLOCA_X]], %[[#ALLOCA_PTR]] : !cir.ptr<!s32i, clang_address_space(offload_private)>, !cir.ptr<!cir.ptr<!s32i, clang_address_space(offload_private)>, clang_address_space(offload_private)>
30+
// CIR-NEXT: cir.store{{.*}} %[[#ALLOCA_X]], %[[#ALLOCA_PTR]] : !cir.ptr<!s32i, language_address_space(offload_private)>, !cir.ptr<!cir.ptr<!s32i, language_address_space(offload_private)>, language_address_space(offload_private)>
3131
// LLVM-NEXT: store ptr %[[#ALLOCA_X]], ptr %[[#ALLOCA_PTR]]
3232

3333
return;

‎clang/test/CIR/CodeGen/OpenCL/array-decay.cl‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,8 @@ kernel void func1(global int *data) {
99
local int arr[32];
1010

1111
local int *ptr = arr;
12-
// CIR: cir.cast array_to_ptrdecay %{{[0-9]+}} : !cir.ptr<!cir.array<!s32i x 32>, clang_address_space(offload_local)> -> !cir.ptr<!s32i, clang_address_space(offload_local)>
13-
// CIR-NEXT: cir.store{{.*}} %{{[0-9]+}}, %{{[0-9]+}} : !cir.ptr<!s32i, clang_address_space(offload_local)>, !cir.ptr<!cir.ptr<!s32i, clang_address_space(offload_local)>, clang_address_space(offload_private)>
12+
// CIR: cir.cast array_to_ptrdecay %{{[0-9]+}} : !cir.ptr<!cir.array<!s32i x 32>, language_address_space(offload_local)> -> !cir.ptr<!s32i, language_address_space(offload_local)>
13+
// CIR-NEXT: cir.store{{.*}} %{{[0-9]+}}, %{{[0-9]+}} : !cir.ptr<!s32i, language_address_space(offload_local)>, !cir.ptr<!cir.ptr<!s32i, language_address_space(offload_local)>, language_address_space(offload_private)>
1414

1515
// LLVM: store ptr addrspace(3) @func1.arr, ptr %{{[0-9]+}}
1616
}
@@ -19,7 +19,7 @@ kernel void func1(global int *data) {
1919
// LLVM: @func2
2020
kernel void func2(global int *data) {
2121
private int arr[32] = {data[2]};
22-
// CIR: %{{[0-9]+}} = cir.get_element %{{[0-9]+}}[%{{[0-9]+}}] : (!cir.ptr<!cir.array<!s32i x 32>, clang_address_space(offload_private)>, !s32i) -> !cir.ptr<!s32i, clang_address_space(offload_private)>
22+
// CIR: %{{[0-9]+}} = cir.get_element %{{[0-9]+}}[%{{[0-9]+}}] : (!cir.ptr<!cir.array<!s32i x 32>, language_address_space(offload_private)>, !s32i) -> !cir.ptr<!s32i, language_address_space(offload_private)>
2323

2424
// LLVM: %{{[0-9]+}} = getelementptr [32 x i32], ptr %3, i32 0, i64 0
2525
}

‎clang/test/CIR/CodeGen/OpenCL/global.cl‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -4,23 +4,23 @@
44
// RUN: FileCheck --input-file=%t.ll %s --check-prefix=LLVM
55

66
global int a = 13;
7-
// CIR-DAG: cir.global external clang_address_space(offload_global) @a = #cir.int<13> : !s32i
7+
// CIR-DAG: cir.global external language_address_space(offload_global) @a = #cir.int<13> : !s32i
88
// LLVM-DAG: @a = addrspace(1) global i32 13
99

1010
global int b = 15;
11-
// CIR-DAG: cir.global external clang_address_space(offload_global) @b = #cir.int<15> : !s32i
11+
// CIR-DAG: cir.global external language_address_space(offload_global) @b = #cir.int<15> : !s32i
1212
// LLVM-DAG: @b = addrspace(1) global i32 15
1313

1414
constant int c[2] = {18, 21};
15-
// CIR-DAG: cir.global constant {{.*}}clang_address_space(offload_constant) {{.*}}@c
15+
// CIR-DAG: cir.global constant {{.*}}language_address_space(offload_constant) {{.*}}@c
1616
// LLVM-DAG: @c = addrspace(2) constant
1717

1818
kernel void test_get_global() {
1919
a = b;
20-
// CIR: %[[#ADDRB:]] = cir.get_global @b : !cir.ptr<!s32i, clang_address_space(offload_global)>
21-
// CIR-NEXT: %[[#LOADB:]] = cir.load{{.*}} %[[#ADDRB]] : !cir.ptr<!s32i, clang_address_space(offload_global)>, !s32i
22-
// CIR-NEXT: %[[#ADDRA:]] = cir.get_global @a : !cir.ptr<!s32i, clang_address_space(offload_global)>
23-
// CIR-NEXT: cir.store{{.*}} %[[#LOADB]], %[[#ADDRA]] : !s32i, !cir.ptr<!s32i, clang_address_space(offload_global)>
20+
// CIR: %[[#ADDRB:]] = cir.get_global @b : !cir.ptr<!s32i, language_address_space(offload_global)>
21+
// CIR-NEXT: %[[#LOADB:]] = cir.load{{.*}} %[[#ADDRB]] : !cir.ptr<!s32i, language_address_space(offload_global)>, !s32i
22+
// CIR-NEXT: %[[#ADDRA:]] = cir.get_global @a : !cir.ptr<!s32i, language_address_space(offload_global)>
23+
// CIR-NEXT: cir.store{{.*}} %[[#LOADB]], %[[#ADDRA]] : !s32i, !cir.ptr<!s32i, language_address_space(offload_global)>
2424

2525
// LLVM: %[[#LOADB:]] = load i32, ptr addrspace(1) @b, align 4
2626
// LLVM-NEXT: store i32 %[[#LOADB]], ptr addrspace(1) @a, align 4

‎clang/test/CIR/CodeGen/OpenCL/printf.cl‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -28,8 +28,8 @@ kernel void test_printf_float2(float2 arg) {
2828
printf("%v2hlf", arg);
2929
}
3030
// CIR-ALL-LABEL: @test_printf_float2(
31-
// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, clang_address_space(offload_constant)>, !cir.vector<!cir.float x 2>) -> !s32i cc(spir_function)
32-
// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, clang_address_space(offload_constant)>, !cir.vector<!cir.float x 2>) -> !s32i cc(spir_function)
31+
// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, language_address_space(offload_constant)>, !cir.vector<!cir.float x 2>) -> !s32i cc(spir_function)
32+
// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, language_address_space(offload_constant)>, !cir.vector<!cir.float x 2>) -> !s32i cc(spir_function)
3333
// LLVM-ALL-LABEL: @test_printf_float2(
3434
// LLVM-FP64: %{{.+}} = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str, <2 x float> %{{.*}})
3535
// LLVM-NOFP64: call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str, <2 x float> %{{.*}})
@@ -38,8 +38,8 @@ kernel void test_printf_half2(half2 arg) {
3838
printf("%v2hf", arg);
3939
}
4040
// CIR-ALL-LABEL: @test_printf_half2(
41-
// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, clang_address_space(offload_constant)>, !cir.vector<!cir.f16 x 2>) -> !s32i cc(spir_function)
42-
// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, clang_address_space(offload_constant)>, !cir.vector<!cir.f16 x 2>) -> !s32i cc(spir_function)
41+
// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, language_address_space(offload_constant)>, !cir.vector<!cir.f16 x 2>) -> !s32i cc(spir_function)
42+
// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, language_address_space(offload_constant)>, !cir.vector<!cir.f16 x 2>) -> !s32i cc(spir_function)
4343
// LLVM-ALL-LABEL: @test_printf_half2(
4444
// LLVM-FP64: %{{.+}} = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.1, <2 x half> %{{.*}})
4545
// LLVM-NOFP64: %{{.+}} = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.1, <2 x half> %{{.*}})
@@ -49,7 +49,7 @@ kernel void test_printf_double2(double2 arg) {
4949
printf("%v2lf", arg);
5050
}
5151
// CIR-FP64-LABEL: @test_printf_double2(
52-
// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, clang_address_space(offload_constant)>, !cir.vector<!cir.double x 2>) -> !s32i cc(spir_function)
52+
// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i, language_address_space(offload_constant)>, !cir.vector<!cir.double x 2>) -> !s32i cc(spir_function)
5353
// LLVM-FP64-LABEL: @test_printf_double2(
5454
// LLVM-FP64: call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.2, <2 x double> %{{.*}})
5555
#endif

0 commit comments

Comments
 (0)