1 // RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s 2 #include "Inputs/cuda.h" 3 4 // CHECK-LABEL: @_Z16use_dispatch_ptrPi( 5 // CHECK: %[[PTR:.*]] = call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() 6 // CHECK: %{{.*}} = addrspacecast i8 addrspace(4)* %[[PTR]] to i8* use_dispatch_ptr(int * out)7__global__ void use_dispatch_ptr(int* out) { 8 const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr(); 9 *out = *dispatch_ptr; 10 } 11 12 // CHECK-LABEL: @_Z12test_ds_fmaxf( 13 // CHECK: call contract float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %{{[^,]*}}, i32 0, i32 0, i1 false) 14 __global__ test_ds_fmax(float src)15void test_ds_fmax(float src) { 16 __shared__ float shared; 17 volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false); 18 } 19 20 // CHECK-LABEL: @_Z6endpgmv( 21 // CHECK: call void @llvm.amdgcn.endpgm() endpgm()22__global__ void endpgm() { 23 __builtin_amdgcn_endpgm(); 24 } 25