From 222e2a7620e6520ffaf4fc4e69d79c18da31542e Mon Sep 17 00:00:00 2001 From: "Zancanaro; Carlo" Date: Mon, 24 Sep 2012 09:58:17 +1000 Subject: Add the clang library to the repo (with some of my changes, too). --- .../CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl | 11 +++++++++ clang/test/CodeGenOpenCL/address-spaces.cl | 27 ++++++++++++++++++++++ clang/test/CodeGenOpenCL/ext-vector-shuffle.cl | 17 ++++++++++++++ clang/test/CodeGenOpenCL/fpmath.cl | 25 ++++++++++++++++++++ clang/test/CodeGenOpenCL/kernel-metadata.cl | 10 ++++++++ clang/test/CodeGenOpenCL/local.cl | 7 ++++++ clang/test/CodeGenOpenCL/ptx-calls.cl | 12 ++++++++++ clang/test/CodeGenOpenCL/ptx-kernels.cl | 10 ++++++++ .../CodeGenOpenCL/single-precision-constant.cl | 7 ++++++ clang/test/CodeGenOpenCL/vector_literals_nested.cl | 23 ++++++++++++++++++ clang/test/CodeGenOpenCL/vector_literals_valid.cl | 22 ++++++++++++++++++ clang/test/CodeGenOpenCL/vector_logops.cl | 19 +++++++++++++++ 12 files changed, 190 insertions(+) create mode 100644 clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl create mode 100644 clang/test/CodeGenOpenCL/address-spaces.cl create mode 100644 clang/test/CodeGenOpenCL/ext-vector-shuffle.cl create mode 100644 clang/test/CodeGenOpenCL/fpmath.cl create mode 100644 clang/test/CodeGenOpenCL/kernel-metadata.cl create mode 100644 clang/test/CodeGenOpenCL/local.cl create mode 100644 clang/test/CodeGenOpenCL/ptx-calls.cl create mode 100644 clang/test/CodeGenOpenCL/ptx-kernels.cl create mode 100644 clang/test/CodeGenOpenCL/single-precision-constant.cl create mode 100644 clang/test/CodeGenOpenCL/vector_literals_nested.cl create mode 100644 clang/test/CodeGenOpenCL/vector_literals_valid.cl create mode 100644 clang/test/CodeGenOpenCL/vector_logops.cl (limited to 'clang/test/CodeGenOpenCL') diff --git a/clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl b/clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl new file mode 100644 index 0000000..fbe3d89 --- /dev/null +++ b/clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl @@ -0,0 +1,11 @@ +// RUN: %clang_cc1 %s -emit-llvm -o %t + +typedef __attribute__((ext_vector_type(4))) unsigned char uchar4; +typedef __attribute__((ext_vector_type(8))) unsigned char uchar8; + +// OpenCL allows vectors to be initialized by vectors Handle bug in +// VisitInitListExpr for this case below. +void foo( uchar8 x ) +{ + uchar4 val[4] = {{(uchar4){x.lo}}}; +} diff --git a/clang/test/CodeGenOpenCL/address-spaces.cl b/clang/test/CodeGenOpenCL/address-spaces.cl new file mode 100644 index 0000000..e030c77 --- /dev/null +++ b/clang/test/CodeGenOpenCL/address-spaces.cl @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s + +void f__p(__private int *arg) { } +// CHECK: i32* nocapture %arg + +void f__g(__global int *arg) { } +// CHECK: i32 addrspace(1)* nocapture %arg + +void f__l(__local int *arg) { } +// CHECK: i32 addrspace(2)* nocapture %arg + +void f__c(__constant int *arg) { } +// CHECK: i32 addrspace(3)* nocapture %arg + + +void fp(private int *arg) { } +// CHECK: i32* nocapture %arg + +void fg(global int *arg) { } +// CHECK: i32 addrspace(1)* nocapture %arg + +void fl(local int *arg) { } +// CHECK: i32 addrspace(2)* nocapture %arg + +void fc(constant int *arg) { } +// CHECK: i32 addrspace(3)* nocapture %arg + diff --git a/clang/test/CodeGenOpenCL/ext-vector-shuffle.cl b/clang/test/CodeGenOpenCL/ext-vector-shuffle.cl new file mode 100644 index 0000000..ee88ba3 --- /dev/null +++ b/clang/test/CodeGenOpenCL/ext-vector-shuffle.cl @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - | not grep 'extractelement' +// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - | not grep 'insertelement' +// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - | grep 'shufflevector' + +typedef __attribute__(( ext_vector_type(2) )) float float2; +typedef __attribute__(( ext_vector_type(4) )) float float4; + +float2 test1(float4 V) { + return V.xy + V.wz; +} + +float4 test2(float4 V) { + float2 W = V.ww; + return W.xyxy + W.yxyx; +} + +float4 test3(float4 V1, float4 V2) { return (float4)(V1.zw, V2.xy); } diff --git a/clang/test/CodeGenOpenCL/fpmath.cl b/clang/test/CodeGenOpenCL/fpmath.cl new file mode 100644 index 0000000..704fcd7 --- /dev/null +++ b/clang/test/CodeGenOpenCL/fpmath.cl @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s + +typedef __attribute__(( ext_vector_type(4) )) float float4; + +float spscalardiv(float a, float b) { + // CHECK: @spscalardiv + // CHECK: fdiv{{.*}}, !fpmath ![[MD:[0-9]+]] + return a / b; +} + +float4 spvectordiv(float4 a, float4 b) { + // CHECK: @spvectordiv + // CHECK: fdiv{{.*}}, !fpmath ![[MD]] + return a / b; +} + +#pragma OPENCL EXTENSION cl_khr_fp64 : enable + +double dpscalardiv(double a, double b) { + // CHECK: @dpscalardiv + // CHECK-NOT: !fpmath + return a / b; +} + +// CHECK: ![[MD]] = metadata !{float 2.500000e+00} diff --git a/clang/test/CodeGenOpenCL/kernel-metadata.cl b/clang/test/CodeGenOpenCL/kernel-metadata.cl new file mode 100644 index 0000000..3e10a11 --- /dev/null +++ b/clang/test/CodeGenOpenCL/kernel-metadata.cl @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s + +void normal_function() { +} + +__kernel void kernel_function() { +} + +// CHECK: !opencl.kernels = !{!0} +// CHECK: !0 = metadata !{void ()* @kernel_function} diff --git a/clang/test/CodeGenOpenCL/local.cl b/clang/test/CodeGenOpenCL/local.cl new file mode 100644 index 0000000..32fa7be --- /dev/null +++ b/clang/test/CodeGenOpenCL/local.cl @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s + +__kernel void foo(void) { + // CHECK: @foo.i = internal addrspace(2) + __local int i; + ++i; +} diff --git a/clang/test/CodeGenOpenCL/ptx-calls.cl b/clang/test/CodeGenOpenCL/ptx-calls.cl new file mode 100644 index 0000000..6f33640 --- /dev/null +++ b/clang/test/CodeGenOpenCL/ptx-calls.cl @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s + +void device_function() { +} +// CHECK: define ptx_device void @device_function() + +__kernel void kernel_function() { + device_function(); +} +// CHECK: define ptx_kernel void @kernel_function() +// CHECK: call ptx_device void @device_function() + diff --git a/clang/test/CodeGenOpenCL/ptx-kernels.cl b/clang/test/CodeGenOpenCL/ptx-kernels.cl new file mode 100644 index 0000000..4d6fa10 --- /dev/null +++ b/clang/test/CodeGenOpenCL/ptx-kernels.cl @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s + +void device_function() { +} +// CHECK: define ptx_device void @device_function() + +__kernel void kernel_function() { +} +// CHECK: define ptx_kernel void @kernel_function() + diff --git a/clang/test/CodeGenOpenCL/single-precision-constant.cl b/clang/test/CodeGenOpenCL/single-precision-constant.cl new file mode 100644 index 0000000..62b37c1 --- /dev/null +++ b/clang/test/CodeGenOpenCL/single-precision-constant.cl @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 %s -cl-single-precision-constant -emit-llvm -o - | FileCheck %s + +float fn(float f) { + // CHECK: fmul float + // CHECK: fadd float + return f*2. + 1.; +} diff --git a/clang/test/CodeGenOpenCL/vector_literals_nested.cl b/clang/test/CodeGenOpenCL/vector_literals_nested.cl new file mode 100644 index 0000000..b9013d0 --- /dev/null +++ b/clang/test/CodeGenOpenCL/vector_literals_nested.cl @@ -0,0 +1,23 @@ +// RUN: %clang_cc1 %s -emit-llvm -O3 -o - | FileCheck %s + +typedef int int2 __attribute((ext_vector_type(2))); +typedef int int4 __attribute((ext_vector_type(4))); + +__constant const int4 itest1 = (int4)(1, 2, ((int2)(3, 4))); +// CHECK: constant <4 x i32> +__constant const int4 itest2 = (int4)(1, 2, ((int2)(3))); +// CHECK: constant <4 x i32> + +typedef float float2 __attribute((ext_vector_type(2))); +typedef float float4 __attribute((ext_vector_type(4))); + +void ftest1(float4 *p) { + *p = (float4)(1.1f, 1.2f, ((float2)(1.3f, 1.4f))); +// CHECK: store <4 x float> +} + +float4 ftest2(float4 *p) { + *p = (float4)(1.1f, 1.2f, ((float2)(1.3f))); +// CHECK: store <4 x float> +} + diff --git a/clang/test/CodeGenOpenCL/vector_literals_valid.cl b/clang/test/CodeGenOpenCL/vector_literals_valid.cl new file mode 100644 index 0000000..bba5b23 --- /dev/null +++ b/clang/test/CodeGenOpenCL/vector_literals_valid.cl @@ -0,0 +1,22 @@ +// RUN: %clang_cc1 -emit-llvm %s -o %t + +typedef __attribute__(( ext_vector_type(2) )) int int2; +typedef __attribute__(( ext_vector_type(3) )) int int3; +typedef __attribute__(( ext_vector_type(4) )) int int4; +typedef __attribute__(( ext_vector_type(8) )) int int8; +typedef __attribute__(( ext_vector_type(4) )) float float4; + +void vector_literals_valid() { + int4 a_1_1_1_1 = (int4)(1,2,3,4); + int4 a_2_1_1 = (int4)((int2)(1,2),3,4); + int4 a_1_2_1 = (int4)(1,(int2)(2,3),4); + int4 a_1_1_2 = (int4)(1,2,(int2)(3,4)); + int4 a_2_2 = (int4)((int2)(1,2),(int2)(3,4)); + int4 a_3_1 = (int4)((int3)(1,2,3),4); + int4 a_1_3 = (int4)(1,(int3)(2,3,4)); + int4 a = (int4)(1); + int8 b = (int8)(1,2,a.xy,a); + float4 V2 = (float4) (1); +} + + diff --git a/clang/test/CodeGenOpenCL/vector_logops.cl b/clang/test/CodeGenOpenCL/vector_logops.cl new file mode 100644 index 0000000..388f1d7 --- /dev/null +++ b/clang/test/CodeGenOpenCL/vector_logops.cl @@ -0,0 +1,19 @@ +// RUN: %clang_cc1 -O3 %s -emit-llvm -o - | FileCheck %s + +typedef int int2 __attribute((ext_vector_type(2))); + +int test1() +{ + int2 a = (int2)(1,0); + int2 b = (int2)(1,1); + return (a&&b).x + (a||b).y; + // CHECK: ret i32 -2 +} + +int test2() +{ + int2 a = (int2)(1,0); + return (!a).y; + // CHECK: ret i32 -1 +} + -- cgit v1.2.3