diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-01-26 06:37:54 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-01-26 06:37:54 +0000 |
| commit | cf70cb9d00adfe5f0ba968b2944f4e28209a26aa (patch) | |
| tree | dfd4cd30439b8655e97ff6637110adad6717b0c5 | |
| parent | b9c932f26e24e4df776ff91a0efbcb6e8eb38ecc (diff) | |
| download | bcm5719-llvm-cf70cb9d00adfe5f0ba968b2944f4e28209a26aa.tar.gz bcm5719-llvm-cf70cb9d00adfe5f0ba968b2944f4e28209a26aa.zip | |
AMDGPU: Add amdgcn cube builtins
llvm-svn: 258794
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 24 |
2 files changed, 28 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 17d6547bb96..c0c7f08ebd3 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -33,6 +33,10 @@ BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc") BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc") BUILTIN(__builtin_amdgcn_class, "bdi", "nc") BUILTIN(__builtin_amdgcn_classf, "bfi", "nc") +BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc") +BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc") +BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc") +BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc") // Legacy names with amdgpu prefix BUILTIN(__builtin_amdgpu_rsq, "dd", "nc") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 6f9a8cf9216..641ec8e0fd2 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -148,6 +148,30 @@ void test_s_barrier() __builtin_amdgcn_s_barrier(); } +// CHECK-LABEL: @test_cubeid( +// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) +void test_cubeid(global float* out, float a, float b, float c) { + *out = __builtin_amdgcn_cubeid(a, b, c); +} + +// CHECK-LABEL: @test_cubesc( +// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c) +void test_cubesc(global float* out, float a, float b, float c) { + *out = __builtin_amdgcn_cubesc(a, b, c); +} + +// CHECK-LABEL: @test_cubetc( +// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c) +void test_cubetc(global float* out, float a, float b, float c) { + *out = __builtin_amdgcn_cubetc(a, b, c); +} + +// CHECK-LABEL: @test_cubema( +// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c) +void test_cubema(global float* out, float a, float b, float c) { + *out = __builtin_amdgcn_cubema(a, b, c); +} + // Legacy intrinsics with AMDGPU prefix // CHECK-LABEL: @test_legacy_rsq_f32 |
