Mercurial > hg > CbC > CbC_llvm

comparison clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl @ 236:c4bab56944e8 llvm-original

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
LLVM 16
author kono
date Wed, 09 Nov 2022 17:45:10 +0900
parents
children 1f2b6ac9f198
comparison
equal deleted inserted replaced
232:70dce7da266c 236:c4bab56944e8
1 // REQUIRES: amdgpu-registered-target
2 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -emit-llvm -o - %s | FileCheck %s
3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1101 -S -emit-llvm -o - %s | FileCheck %s
4 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1102 -S -emit-llvm -o - %s | FileCheck %s
5 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1103 -S -emit-llvm -o - %s | FileCheck %s
6
7 typedef unsigned int uint;
8 typedef unsigned long ulong;
9 typedef uint uint2 __attribute__((ext_vector_type(2)));
10 typedef uint uint4 __attribute__((ext_vector_type(4)));
11
12 // CHECK-LABEL: @test_s_sendmsg_rtn(
13 // CHECK: call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 0)
14 void test_s_sendmsg_rtn(global uint* out) {
15 *out = __builtin_amdgcn_s_sendmsg_rtn(0);
16 }
17
18 // CHECK-LABEL: @test_s_sendmsg_rtnl(
19 // CHECK: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0)
20 void test_s_sendmsg_rtnl(global ulong* out) {
21 *out = __builtin_amdgcn_s_sendmsg_rtnl(0);
22 }
23
24 // CHECK-LABEL: @test_ds_bvh_stack_rtn(
25 // CHECK: %0 = tail call { i32, i32 } @llvm.amdgcn.ds.bvh.stack.rtn(i32 %addr, i32 %data, <4 x i32> %data1, i32 128)
26 // CHECK: %1 = extractvalue { i32, i32 } %0, 0
27 // CHECK: %2 = extractvalue { i32, i32 } %0, 1
28 // CHECK: %3 = insertelement <2 x i32> poison, i32 %1, i64 0
29 // CHECK: %4 = insertelement <2 x i32> %3, i32 %2, i64 1
30 void test_ds_bvh_stack_rtn(global uint2* out, uint addr, uint data, uint4 data1)
31 {
32 *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, 128);
33 }
34
35 // CHECK-LABEL: @test_permlane64(
36 // CHECK: call i32 @llvm.amdgcn.permlane64(i32 %a)
37 void test_permlane64(global uint* out, uint a) {
38 *out = __builtin_amdgcn_permlane64(a);
39 }