diff options
Diffstat (limited to 'clang/test/CodeGenCUDA/builtins-amdgcn.cu')
| -rw-r--r-- | clang/test/CodeGenCUDA/builtins-amdgcn.cu | 18 |
1 files changed, 18 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/builtins-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-amdgcn.cu new file mode 100644 index 00000000000..82a666717ac --- /dev/null +++ b/clang/test/CodeGenCUDA/builtins-amdgcn.cu @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s +#include "Inputs/cuda.h" + +// CHECK-LABEL: @_Z16use_dispatch_ptrPi( +// CHECK: %2 = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() +// CHECK: %3 = addrspacecast i8 addrspace(4)* %2 to i8 addrspace(4)** +__global__ void use_dispatch_ptr(int* out) { + const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr(); + *out = *dispatch_ptr; +} + +// CHECK-LABEL: @_Z12test_ds_fmaxf( +// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %2, i32 0, i32 0, i1 false) +__global__ +void test_ds_fmax(float src) { + __shared__ float shared; + volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false); +} |
