| 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: %2 = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr() |
|---|
| 6 | // CHECK: %3 = addrspacecast i8 addrspace(4)* %2 to i8 addrspace(4)** |
|---|
| 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 float @llvm.amdgcn.ds.fmax(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %2, i32 0, i32 0, i1 false) |
|---|
| 14 | __global__ |
|---|
| 15 | void test_ds_fmax(float src) { |
|---|
| 16 | __shared__ float shared; |
|---|
| 17 | volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false); |
|---|
| 18 | } |
|---|
| 19 | |
|---|