int main(int argc, char *argv[]) {
// int n=2500;
// double a[n];
//
// #pragma omp target device(1)
// //#pragma omp teams num_teams(1) thread_limit(1)
// //#pragma omp distribute
// #pragma omp parallel for
// for (int i = 0; i < 32; ++i) {
// a[i] += 1;
// }
#pragma omp target device(1)
{
static __attribute__((address_space(3))) int A;
A = 123;
#pragma omp parallel num_threads(2)
{
printf("%d\n",A);
}
}
return 0;
}
// CHECK-COMMON: %[[PTRINT:.*]] = ptrtoint
// CHECK-COMMON-NEXT: %[[MISALIGN:.*]] = and i64 %[[PTRINT]], 3
// CHECK-COMMON-NEXT: icmp eq i64 %[[MISALIGN]], 0
// CHECK-UBSAN: %[[ARG:.*]] = ptrtoint
// CHECK-UBSAN-NEXT: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_200]] to i8*), i64 %[[ARG]])
// CHECK-TRAP: call void @llvm.trap() [[NR_NUW]]
// CHECK-TRAP-NEXT: unreachable
#line 200
return *a;
}
// CHECK-UBSAN-LABEL: @addr_space
int addr_space(int __attribute__((address_space(256))) *a) {
// CHECK-UBSAN-NOT: __ubsan
return *a;
}
// CHECK-COMMON-LABEL: @lsh_overflow
int lsh_overflow(int a, int b) {
// CHECK-COMMON: %[[RHS_INBOUNDS:.*]] = icmp ule i32 %[[RHS:.*]], 31
// CHECK-COMMON-NEXT: br i1 %[[RHS_INBOUNDS]], label %[[CHECK_BB:.*]], label %[[CONT_BB:.*]],
// CHECK-COMMON: [[CHECK_BB]]:
// CHECK-COMMON-NEXT: %[[SHIFTED_OUT_WIDTH:.*]] = sub nuw nsw i32 31, %[[RHS]]
// CHECK-COMMON-NEXT: %[[SHIFTED_OUT:.*]] = lshr i32 %[[LHS:.*]], %[[SHIFTED_OUT_WIDTH]]
// CHECK-COMMON-NEXT: %[[NO_OVERFLOW:.*]] = icmp eq i32 %[[SHIFTED_OUT]], 0
// CHECK-COMMON-NEXT: br label %[[CONT_BB]]
#include <pocl_device.h>
void _pocl_launcher_lud_internal_workgroup(void** args, struct pocl_context*);
void _pocl_launcher_lud_internal_workgroup_fast(void** args, struct pocl_context*);
__attribute__((address_space(3))) __kernel_metadata _lud_internal_md = {
"lud_internal", /* name */
5, /* num_args */
0, /* num_locals */
_pocl_launcher_lud_internal_workgroup_fast
};
#define CHECKER 2
typedef struct {
Vector diffuse, specular;
real reflection, refraction;
int specularRoughness;
} Surface;
typedef struct {
Vector from, ray;
} Ray;
typedef struct {
int type;
real radius;
Vector center, light, normal;
Surface surface;
} Item;
typedef Item __attribute__((address_space(2))) * ItemPtr;
//typedef Item * ItemPtr;
typedef struct {
int depth;
Ray ray;
real magnitude, refractiveIndex;
} StackEntry;
typedef struct {
StackEntry entries[100];
int size;
} Stack;
typedef struct {
Ray ray;
Vector vVector, hVector;
} Viewport;
#define MAX_LIGHTS (8)
// RUN: clang-cc -emit-llvm < %s | grep '@foo.*global.*addrspace(1)'
// RUN: clang-cc -emit-llvm < %s | grep '@ban.*global.*addrspace(1)'
// RUN: clang-cc -emit-llvm < %s | grep 'load.*addrspace(1)' | count 2
// RUN: clang-cc -emit-llvm < %s | grep 'load.*addrspace(2).. @A'
// RUN: clang-cc -emit-llvm < %s | grep 'load.*addrspace(2).. @B'
int foo __attribute__((address_space(1)));
int ban[10] __attribute__((address_space(1)));
int bar() { return foo; }
int baz(int i) { return ban[i]; }
// Both A and B point into addrspace(2).
__attribute__((address_space(2))) int *A, *B;
void test3() {
*A = *B;
}
// RUN: %clang_cc1 -fsyntax-only -verify %s
typedef int __attribute__((address_space(1))) int_1;
typedef int __attribute__((address_space(2))) int_2;
void f0(int_1 &); // expected-note{{candidate function not viable: address space mismatch in 1st argument ('int'), parameter type must be 'int_1 &' (aka '__attribute__((address_space(1))) int &')}} \
// expected-note{{candidate function not viable: address space mismatch in 1st argument ('int_2' (aka '__attribute__((address_space(2))) int')), parameter type must be 'int_1 &' (aka '__attribute__((address_space(1))) int &')}}
void f0(const int_1 &); // expected-note{{candidate function not viable: address space mismatch in 1st argument ('int'), parameter type must be 'const int_1 &' (aka 'const __attribute__((address_space(1))) int &')}} \
// expected-note{{candidate function not viable: address space mismatch in 1st argument ('int_2' (aka '__attribute__((address_space(2))) int')), parameter type must be 'const int_1 &' (aka 'const __attribute__((address_space(1))) int &')}}
void test_f0() {
int i;
static int_1 i1;
static int_2 i2;
f0(i); // expected-error{{no matching function for call to 'f0'}}
f0(i1);
f0(i2); // expected-error{{no matching function for call to 'f0'}}
}
// RUN: %clang_cc1 -triple x86_64-apple-darwin -emit-llvm < %s | FileCheck %s
// CHECK: @foo = common addrspace(1) global
int foo __attribute__((address_space(1)));
// CHECK: @ban = common addrspace(1) global
int ban[10] __attribute__((address_space(1)));
// CHECK-LABEL: define i32 @test1()
// CHECK: load i32 addrspace(1)* @foo
int test1() { return foo; }
// CHECK-LABEL: define i32 @test2(i32 %i)
// CHECK: load i32 addrspace(1)*
// CHECK-NEXT: ret i32
int test2(int i) { return ban[i]; }
// Both A and B point into addrspace(2).
__attribute__((address_space(2))) int *A, *B;
// CHECK-LABEL: define void @test3()
// CHECK: load i32 addrspace(2)** @B
// CHECK: load i32 addrspace(2)*
// CHECK: load i32 addrspace(2)** @A
// CHECK: store i32 {{.*}}, i32 addrspace(2)*
void test3() {
*A = *B;
}
// PR7437
typedef struct {
// RUN: %clang_cc1 -fsyntax-only -verify %s
template<typename T, typename U>
struct is_same {
static const bool value = false;
};
template<typename T>
struct is_same<T, T> {
static const bool value = true;
};
typedef int __attribute__((address_space(1))) int_1;;
typedef int __attribute__((address_space(2))) int_2;;
typedef int __attribute__((address_space(1))) *int_1_ptr;
typedef int_2 *int_2_ptr;
// Check that we maintain address spaces through template argument
// deduction from a type.
template<typename T>
struct remove_pointer {
typedef T type;
};
template<typename T>
struct remove_pointer<T *> {
typedef T type;
};
int check_remove0[is_same<remove_pointer<int_1_ptr>::type, int_1>::value? 1 : -1];
int check_remove1[is_same<remove_pointer<int_2_ptr>::type, int_2>::value? 1 : -1];
// Test this without pch.
// RUN: %clang_cc1 -fblocks -include %S/types.h -fsyntax-only -verify %s
// Test with pch.
// RUN: %clang_cc1 -emit-pch -fblocks -o %t %S/types.h
// RUN: %clang_cc1 -fblocks -include-pch %t -fsyntax-only -verify %s -ast-print
typedef int INT;
INT int_value;
__attribute__((address_space(1))) int int_as_one;
// TYPE_EXT_QUAL
ASInt *as_int_ptr1 = &int_value; // expected-error{{changes address space of pointer}}
ASInt *as_int_ptr2 = &int_as_one;
// TYPE_COMPLEX
_Complex float Cfloat_val;
Cfloat *Cfloat_ptr = &Cfloat_val;
// TYPE_ATOMIC
_Atomic(int) AtomicInt_val;
AtomicInt *AtomicInt_ptr = &AtomicInt_val;
// TYPE_POINTER
int_ptr int_value_ptr = &int_value;
// TYPE_BLOCK_POINTER
void test_block_ptr(Block *bl) {
*bl = ^(int x, float f) { return x; };
}
