173 lines · plain
1// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s2// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO3 4kernel void foo(global int * globalintp, global int * restrict globalintrestrictp,5 global const int * globalconstintp,6 global const int * restrict globalconstintrestrictp,7 constant int * constantintp, constant int * restrict constantintrestrictp,8 global const volatile int * globalconstvolatileintp,9 global const volatile int * restrict globalconstvolatileintrestrictp,10 global volatile int * globalvolatileintp,11 global volatile int * restrict globalvolatileintrestrictp,12 local int * localintp, local int * restrict localintrestrictp,13 local const int * localconstintp,14 local const int * restrict localconstintrestrictp,15 local const volatile int * localconstvolatileintp,16 local const volatile int * restrict localconstvolatileintrestrictp,17 local volatile int * localvolatileintp,18 local volatile int * restrict localvolatileintrestrictp,19 int X, const int constint, const volatile int constvolatileint,20 volatile int volatileint) {21 *globalintrestrictp = constint + volatileint;22}23// CHECK: define{{.*}} spir_kernel void @foo{{[^!]+}}24// ARGINFO: define{{.*}} spir_kernel void @foo{{[^!]+}}25// CHECK: !kernel_arg_addr_space ![[MD11:[0-9]+]]26// CHECK: !kernel_arg_access_qual ![[MD12:[0-9]+]]27// CHECK: !kernel_arg_type ![[MD13:[0-9]+]]28// CHECK: !kernel_arg_base_type ![[MD13]]29// CHECK: !kernel_arg_type_qual ![[MD14:[0-9]+]]30// CHECK-NOT: !kernel_arg_name31// ARGINFO: !kernel_arg_name ![[MD15:[0-9]+]]32 33kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3, read_write image1d_t img4) {34}35// CHECK: define{{.*}} spir_kernel void @foo2{{[^!]+}}36// ARGINFO: define{{.*}} spir_kernel void @foo2{{[^!]+}}37// CHECK: !kernel_arg_addr_space ![[MD21:[0-9]+]]38// CHECK: !kernel_arg_access_qual ![[MD22:[0-9]+]]39// CHECK: !kernel_arg_type ![[MD23:[0-9]+]]40// CHECK: !kernel_arg_base_type ![[MD23]]41// CHECK: !kernel_arg_type_qual ![[MD24:[0-9]+]]42// CHECK-NOT: !kernel_arg_name43// ARGINFO: !kernel_arg_name ![[MD25:[0-9]+]]44 45kernel void foo3(__global half * X) {46}47// CHECK: define{{.*}} spir_kernel void @foo3{{[^!]+}}48// ARGINFO: define{{.*}} spir_kernel void @foo3{{[^!]+}}49// CHECK: !kernel_arg_addr_space ![[MD31:[0-9]+]]50// CHECK: !kernel_arg_access_qual ![[MD32:[0-9]+]]51// CHECK: !kernel_arg_type ![[MD33:[0-9]+]]52// CHECK: !kernel_arg_base_type ![[MD33]]53// CHECK: !kernel_arg_type_qual ![[MD34:[0-9]+]]54// CHECK-NOT: !kernel_arg_name55// ARGINFO: !kernel_arg_name ![[MD35:[0-9]+]]56 57typedef unsigned int myunsignedint;58kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) {59}60// CHECK: define{{.*}} spir_kernel void @foo4{{[^!]+}}61// ARGINFO: define{{.*}} spir_kernel void @foo4{{[^!]+}}62// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]63// CHECK: !kernel_arg_access_qual ![[MD42:[0-9]+]]64// CHECK: !kernel_arg_type ![[MD43:[0-9]+]]65// CHECK: !kernel_arg_base_type ![[MD44:[0-9]+]]66// CHECK: !kernel_arg_type_qual ![[MD45:[0-9]+]]67// CHECK-NOT: !kernel_arg_name68// ARGINFO: !kernel_arg_name ![[MD46:[0-9]+]]69 70typedef image1d_t myImage;71kernel void foo5(myImage img1, write_only image1d_t img2) {72}73// CHECK: define{{.*}} spir_kernel void @foo5{{[^!]+}}74// ARGINFO: define{{.*}} spir_kernel void @foo5{{[^!]+}}75// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]76// CHECK: !kernel_arg_access_qual ![[MD51:[0-9]+]]77// CHECK: !kernel_arg_type ![[MD52:[0-9]+]]78// CHECK: !kernel_arg_base_type ![[MD53:[0-9]+]]79// CHECK: !kernel_arg_type_qual ![[MD45]]80// CHECK-NOT: !kernel_arg_name81// ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]]82 83typedef char char16 __attribute__((ext_vector_type(16)));84__kernel void foo6(__global char16 arg[]) {}85// CHECK: define{{.*}} spir_kernel void @foo6{{[^!]+}}86// ARGINFO: define{{.*}} spir_kernel void @foo6{{[^!]+}}87// CHECK: !kernel_arg_type ![[MD61:[0-9]+]]88// ARGINFO: !kernel_arg_name ![[MD62:[0-9]+]]89 90typedef read_only image1d_t ROImage;91typedef write_only image1d_t WOImage;92typedef read_write image1d_t RWImage;93kernel void foo7(ROImage ro, WOImage wo, RWImage rw) {94}95 96// CHECK: define{{.*}} spir_kernel void @foo7{{[^!]+}}97// ARGINFO: define{{.*}} spir_kernel void @foo7{{[^!]+}}98// CHECK: !kernel_arg_addr_space ![[MD71:[0-9]+]]99// CHECK: !kernel_arg_access_qual ![[MD72:[0-9]+]]100// CHECK: !kernel_arg_type ![[MD73:[0-9]+]]101// CHECK: !kernel_arg_base_type ![[MD74:[0-9]+]]102// CHECK: !kernel_arg_type_qual ![[MD75:[0-9]+]]103// CHECK-NOT: !kernel_arg_name104// ARGINFO: !kernel_arg_name ![[MD76:[0-9]+]]105 106typedef unsigned char uchar;107typedef uchar uchar2 __attribute__((ext_vector_type(2)));108kernel void foo8(pipe int p1, pipe uchar p2, pipe uchar2 p3, const pipe uchar p4, write_only pipe uchar p5) {}109// CHECK: define{{.*}} spir_kernel void @foo8{{[^!]+}}110// ARGINFO: define{{.*}} spir_kernel void @foo8{{[^!]+}}111// CHECK: !kernel_arg_addr_space ![[PIPE_AS_QUAL:[0-9]+]]112// CHECK: !kernel_arg_access_qual ![[PIPE_ACCESS_QUAL:[0-9]+]]113// CHECK: !kernel_arg_type ![[PIPE_TY:[0-9]+]]114// CHECK: !kernel_arg_base_type ![[PIPE_BASE_TY:[0-9]+]]115// CHECK: !kernel_arg_type_qual ![[PIPE_QUAL:[0-9]+]]116// CHECK-NOT: !kernel_arg_name117// ARGINFO: !kernel_arg_name ![[PIPE_ARG_NAMES:[0-9]+]]118 119kernel void foo9(signed char sc1, global const signed char* sc2) {}120// CHECK: define{{.*}} spir_kernel void @foo9{{[^!]+}}121// ARGINFO: define{{.*}} spir_kernel void @foo9{{[^!]+}}122// CHECK: !kernel_arg_addr_space ![[SCHAR_AS_QUAL:[0-9]+]]123// CHECK: !kernel_arg_access_qual ![[MD42]]124// CHECK: !kernel_arg_type ![[SCHAR_TY:[0-9]+]]125// CHECK: !kernel_arg_base_type ![[SCHAR_TY]]126// CHECK: !kernel_arg_type_qual ![[SCHAR_QUAL:[0-9]+]]127// CHECK-NOT: !kernel_arg_name128// ARGINFO: !kernel_arg_name ![[SCHAR_ARG_NAMES:[0-9]+]]129 130// CHECK: ![[MD11]] = !{i32 1, i32 1, i32 1, i32 1, i32 2, i32 2, i32 1, i32 1, i32 1, i32 1, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 0, i32 0, i32 0, i32 0}131// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none"}132// CHECK: ![[MD13]] = !{!"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int", !"int", !"int", !"int"}133// CHECK: ![[MD14]] = !{!"", !"restrict", !"const", !"restrict const", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"restrict", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"", !"", !""}134// ARGINFO: ![[MD15]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"constantintp", !"constantintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp", !"localintp", !"localintrestrictp", !"localconstintp", !"localconstintrestrictp", !"localconstvolatileintp", !"localconstvolatileintrestrictp", !"localvolatileintp", !"localvolatileintrestrictp", !"X", !"constint", !"constvolatileint", !"volatileint"}135// CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1, i32 1}136// CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only", !"read_write"}137// CHECK: ![[MD23]] = !{!"image1d_t", !"image2d_t", !"image2d_array_t", !"image1d_t"}138// CHECK: ![[MD24]] = !{!"", !"", !"", !""}139// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3", !"img4"}140// CHECK: ![[MD31]] = !{i32 1}141// CHECK: ![[MD32]] = !{!"none"}142// CHECK: ![[MD33]] = !{!"half*"}143// CHECK: ![[MD34]] = !{!""}144// ARGINFO: ![[MD35]] = !{!"X"}145// CHECK: ![[MD41]] = !{i32 1, i32 1}146// CHECK: ![[MD42]] = !{!"none", !"none"}147// CHECK: ![[MD43]] = !{!"uint*", !"myunsignedint*"}148// CHECK: ![[MD44]] = !{!"uint*", !"uint*"}149// CHECK: ![[MD45]] = !{!"", !""}150// ARGINFO: ![[MD46]] = !{!"X", !"Y"}151// CHECK: ![[MD51]] = !{!"read_only", !"write_only"}152// CHECK: ![[MD52]] = !{!"myImage", !"image1d_t"}153// CHECK: ![[MD53]] = !{!"image1d_t", !"image1d_t"}154// ARGINFO: ![[MD54]] = !{!"img1", !"img2"}155// CHECK: ![[MD61]] = !{!"char16*"}156// ARGINFO: ![[MD62]] = !{!"arg"}157// CHECK: ![[MD71]] = !{i32 1, i32 1, i32 1}158// CHECK: ![[MD72]] = !{!"read_only", !"write_only", !"read_write"}159// CHECK: ![[MD73]] = !{!"ROImage", !"WOImage", !"RWImage"}160// CHECK: ![[MD74]] = !{!"image1d_t", !"image1d_t", !"image1d_t"}161// CHECK: ![[MD75]] = !{!"", !"", !""}162// ARGINFO: ![[MD76]] = !{!"ro", !"wo", !"rw"}163// CHECK: ![[PIPE_AS_QUAL]] = !{i32 1, i32 1, i32 1, i32 1, i32 1}164// CHECK: ![[PIPE_ACCESS_QUAL]] = !{!"read_only", !"read_only", !"read_only", !"read_only", !"write_only"}165// CHECK: ![[PIPE_TY]] = !{!"int", !"uchar", !"uchar2", !"uchar", !"uchar"}166// CHECK: ![[PIPE_BASE_TY]] = !{!"int", !"uchar", !"uchar __attribute__((ext_vector_type(2)))", !"uchar", !"uchar"}167// CHECK: ![[PIPE_QUAL]] = !{!"pipe", !"pipe", !"pipe", !"pipe", !"pipe"}168// ARGINFO: ![[PIPE_ARG_NAMES]] = !{!"p1", !"p2", !"p3", !"p4", !"p5"}169// CHECK: ![[SCHAR_AS_QUAL]] = !{i32 0, i32 1}170// CHECK: ![[SCHAR_TY]] = !{!"char", !"char*"}171// CHECK: ![[SCHAR_QUAL]] = !{!"", !"const"}172// ARGINFO: ![[SCHAR_ARG_NAMES]] = !{!"sc1", !"sc2"}173