1 // REQUIRES: amdgpu-registered-target 2 3 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc 4 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s 5 // expected-no-diagnostics 6 #ifndef HEADER 7 #define HEADER 8 9 #define N 1000 10 11 int test_amdgcn_target_tid_threads() { 12 // CHECK-LABEL: define weak void @{{.*}}test_amdgcn_target_tid_threads 13 14 int arr[N]; 15 16 // CHECK: [[NUM_THREADS:%.+]] = call i64 @__ockl_get_local_size(i32 0) 17 // CHECK-NEXT: [[VAR:%.+]] = trunc i64 [[NUM_THREADS]] to i32 18 // CHECK-NEXT: sub nuw i32 [[VAR]], 64 19 // CHECK: call i32 @llvm.amdgcn.workitem.id.x() 20 #pragma omp target 21 for (int i = 0; i < N; i++) { 22 arr[i] = 1; 23 } 24 25 return arr[0]; 26 } 27 28 int test_amdgcn_target_tid_threads_simd() { 29 // CHECK-LABEL: define weak void @{{.*}}test_amdgcn_target_tid_threads_simd 30 31 int arr[N]; 32 33 // CHECK: [[NUM_THREADS:%.+]] = call i64 @__ockl_get_local_size(i32 0) 34 // CHECK-NEXT: [[VAR:%.+]] = trunc i64 [[NUM_THREADS]] to i32 35 // CHECK-NEXT: call void @__kmpc_spmd_kernel_init(i32 [[VAR]], i16 0, i16 0) 36 #pragma omp target simd 37 for (int i = 0; i < N; i++) { 38 arr[i] = 1; 39 } 40 return arr[0]; 41 } 42 43 #endif 44