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; }; }