diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-06-18 14:10:01 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2019-06-18 14:10:01 +0000 |
| commit | 2acc71762710bcaaea77541a6c1877c9593b9d0f (patch) | |
| tree | 54d3d3d8a5e03c3856fecdec0d1a1598163ae8d6 /clang/test | |
| parent | 677423997d6a31e9037e41dca990508d561ad42d (diff) | |
| download | bcm5719-llvm-2acc71762710bcaaea77541a6c1877c9593b9d0f.tar.gz bcm5719-llvm-2acc71762710bcaaea77541a6c1877c9593b9d0f.zip | |
AMDGPU: Add GWS instruction builtins
llvm-svn: 363684
Diffstat (limited to 'clang/test')
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 12 |
1 files changed, 12 insertions, 0 deletions
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index bd7fe78fc96..27a91a18cc9 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -548,6 +548,18 @@ kernel void test_ds_consume_lds(global int* out, local int* ptr) { *out = __builtin_amdgcn_ds_consume(ptr); } +// CHECK-LABEL: @test_gws_init( +// CHECK: call void @llvm.amdgcn.ds.gws.init(i32 %value, i32 %id) +kernel void test_gws_init(uint value, uint id) { + __builtin_amdgcn_ds_gws_init(value, id); +} + +// CHECK-LABEL: @test_gws_barrier( +// CHECK: call void @llvm.amdgcn.ds.gws.barrier(i32 %value, i32 %id) +kernel void test_gws_barrier(uint value, uint id) { + __builtin_amdgcn_ds_gws_barrier(value, id); +} + // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent } |
