diff options
| author | dim <dim@FreeBSD.org> | 2013-12-22 00:07:40 +0000 |
|---|---|---|
| committer | dim <dim@FreeBSD.org> | 2013-12-22 00:07:40 +0000 |
| commit | 952eddef9aff85b1e92626e89baaf7a360e2ac85 (patch) | |
| tree | df8df0b0067b381eab470a3b8f28d14a552a6340 /test/CodeGenOpenCL | |
| parent | ea266cad53e3d49771fa38103913d3ec7a166694 (diff) | |
| download | FreeBSD-src-952eddef9aff85b1e92626e89baaf7a360e2ac85.zip FreeBSD-src-952eddef9aff85b1e92626e89baaf7a360e2ac85.tar.gz | |
Vendor import of clang release_34 branch r197841 (effectively, 3.4 RC3):
https://llvm.org/svn/llvm-project/cfe/branches/release_34@197841
Diffstat (limited to 'test/CodeGenOpenCL')
| -rw-r--r-- | test/CodeGenOpenCL/address-spaces-mangling.cl | 30 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/kernel-attributes.cl | 2 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/local.cl | 4 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/opencl_types.cl | 2 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/ptx-calls.cl | 4 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/ptx-kernels.cl | 4 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/str_literals.cl | 9 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/vector_odd.cl | 17 |
8 files changed, 64 insertions, 8 deletions
diff --git a/test/CodeGenOpenCL/address-spaces-mangling.cl b/test/CodeGenOpenCL/address-spaces-mangling.cl new file mode 100644 index 0000000..3c7a518 --- /dev/null +++ b/test/CodeGenOpenCL/address-spaces-mangling.cl @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s +// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s + +// We can't name this f as private is equivalent to default +// no specifier given address space so we get multiple definition +// warnings, but we do want it for comparison purposes. +__attribute__((overloadable)) +void ff(int *arg) { } +// ASMANG: @_Z2ffPi +// NOASMANG: @_Z2ffPi + +__attribute__((overloadable)) +void f(private int *arg) { } +// ASMANG: @_Z1fPi +// NOASMANG: @_Z1fPi + +__attribute__((overloadable)) +void f(global int *arg) { } +// ASMANG: @_Z1fPU3AS1i +// NOASMANG: @_Z1fPU8CLglobali + +__attribute__((overloadable)) +void f(local int *arg) { } +// ASMANG: @_Z1fPU3AS2i +// NOASMANG: @_Z1fPU7CLlocali + +__attribute__((overloadable)) +void f(constant int *arg) { } +// ASMANG: @_Z1fPU3AS3i +// NOASMANG: @_Z1fPU10CLconstanti diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl index 1166f93..0825ffc 100644 --- a/test/CodeGenOpenCL/kernel-attributes.cl +++ b/test/CodeGenOpenCL/kernel-attributes.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -emit-llvm -O0 -o - %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s typedef unsigned int uint4 __attribute__((ext_vector_type(4))); diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl index b4bd008..b5c67d9 100644 --- a/test/CodeGenOpenCL/local.cl +++ b/test/CodeGenOpenCL/local.cl @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -emit-llvm -o - | FileCheck %s __kernel void foo(void) { // CHECK: @foo.i = internal addrspace(2) @@ -6,7 +6,7 @@ __kernel void foo(void) { ++i; } -// CHECK: define void @_Z3barPU3AS2i +// CHECK-LABEL: define void @_Z3barPU7CLlocali __kernel void __attribute__((__overloadable__)) bar(local int *x) { *x = 5; } diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl index b1e558d..7e99fc5 100644 --- a/test/CodeGenOpenCL/opencl_types.cl +++ b/test/CodeGenOpenCL/opencl_types.cl @@ -22,7 +22,7 @@ void fnc3(image3d_t img) {} // CHECK: @fnc3(%opencl.image3d_t* void fnc4smp(sampler_t s) {} -// CHECK: define void @fnc4smp(i32 +// CHECK-LABEL: define void @fnc4smp(i32 kernel void foo(image1d_t img) { sampler_t smp = 5; diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl index d990451..00f2a0e 100644 --- a/test/CodeGenOpenCL/ptx-calls.cl +++ b/test/CodeGenOpenCL/ptx-calls.cl @@ -2,12 +2,12 @@ void device_function() { } -// CHECK: define void @device_function() +// CHECK-LABEL: define void @device_function() __kernel void kernel_function() { device_function(); } -// CHECK: define void @kernel_function() +// CHECK-LABEL: define void @kernel_function() // CHECK: call void @device_function() // CHECK: !{{[0-9]+}} = metadata !{void ()* @kernel_function, metadata !"kernel", i32 1} diff --git a/test/CodeGenOpenCL/ptx-kernels.cl b/test/CodeGenOpenCL/ptx-kernels.cl index 07648e4..49d207f 100644 --- a/test/CodeGenOpenCL/ptx-kernels.cl +++ b/test/CodeGenOpenCL/ptx-kernels.cl @@ -2,10 +2,10 @@ void device_function() { } -// CHECK: define void @device_function() +// CHECK-LABEL: define void @device_function() __kernel void kernel_function() { } -// CHECK: define void @kernel_function() +// CHECK-LABEL: define void @kernel_function() // CHECK: !{{[0-9]+}} = metadata !{void ()* @kernel_function, metadata !"kernel", i32 1} diff --git a/test/CodeGenOpenCL/str_literals.cl b/test/CodeGenOpenCL/str_literals.cl new file mode 100644 index 0000000..78a9305 --- /dev/null +++ b/test/CodeGenOpenCL/str_literals.cl @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 %s -emit-llvm -o - -ffake-address-space-map | FileCheck %s + +__constant char * __constant x = "hello world"; +__constant char * __constant y = "hello world"; + +// CHECK: addrspace(3) unnamed_addr constant +// CHECK-NOT: addrspace(3) unnamed_addr constant +// CHECK: @x = addrspace(3) global i8 addrspace(3)* +// CHECK: @y = addrspace(3) global i8 addrspace(3)* diff --git a/test/CodeGenOpenCL/vector_odd.cl b/test/CodeGenOpenCL/vector_odd.cl new file mode 100644 index 0000000..c44328b --- /dev/null +++ b/test/CodeGenOpenCL/vector_odd.cl @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 %s -O0 -emit-llvm -o - | FileCheck %s + +typedef unsigned char __attribute__((ext_vector_type(3))) uchar3; + +//CHECK: {{%.*}} = shufflevector <3 x i8> {{%.*}}, <3 x i8> <i8 1, i8 1, i8 undef>, <3 x i32> <i32 0, i32 3, i32 2> + +kernel void test_odd_vector1 (uchar3 lhs) +{ + lhs.odd = 1; +} + +//CHECK: {{%.*}} = shufflevector <3 x i8> {{%.*}}, <3 x i8> <i8 2, i8 2, i8 undef>, <3 x i32> <i32 0, i32 1, i32 3> + +kernel void test_odd_vector2 (uchar3 lhs) +{ + lhs.hi = 2; +} |
