This revision was landed with ongoing or failed builds. This revision was automatically updated to reflect the committed changes. Closed by commit rG6a67e05a26eb: [HIP] Add __builtin_amdgcn_groupstaticsize (authored by ashi1). Herald added a project: clang. Herald added a subscriber: cfe-commits.
Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D102403/new/ https://reviews.llvm.org/D102403 Files: clang/include/clang/Basic/BuiltinsAMDGPU.def clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Index: clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -145,3 +145,10 @@ { *out = __builtin_amdgcn_perm(a, b, s); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) +{ + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl @@ -3,6 +3,7 @@ // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable +typedef unsigned int uint; typedef unsigned long ulong; // CHECK-LABEL: @test_fmed3_f16 @@ -18,3 +19,10 @@ { *out = __builtin_amdgcn_s_memtime(); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) +{ + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl @@ -30,3 +30,10 @@ { *out = __builtin_amdgcn_s_memtime(); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) +{ + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -62,3 +62,9 @@ int test_is_private_global(const global int* ptr) { return __builtin_amdgcn_is_private(ptr); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) { + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/include/clang/Basic/BuiltinsAMDGPU.def =================================================================== --- clang/include/clang/Basic/BuiltinsAMDGPU.def +++ clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -65,6 +65,7 @@ BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n") BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n") +BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n") BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n")
Index: clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -145,3 +145,10 @@ { *out = __builtin_amdgcn_perm(a, b, s); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) +{ + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl @@ -3,6 +3,7 @@ // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable +typedef unsigned int uint; typedef unsigned long ulong; // CHECK-LABEL: @test_fmed3_f16 @@ -18,3 +19,10 @@ { *out = __builtin_amdgcn_s_memtime(); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) +{ + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl @@ -30,3 +30,10 @@ { *out = __builtin_amdgcn_s_memtime(); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) +{ + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl =================================================================== --- clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -62,3 +62,9 @@ int test_is_private_global(const global int* ptr) { return __builtin_amdgcn_is_private(ptr); } + +// CHECK-LABEL: @test_groupstaticsize +// CHECK: call i32 @llvm.amdgcn.groupstaticsize() +void test_groupstaticsize(global uint* out) { + *out = __builtin_amdgcn_groupstaticsize(); +} Index: clang/include/clang/Basic/BuiltinsAMDGPU.def =================================================================== --- clang/include/clang/Basic/BuiltinsAMDGPU.def +++ clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -65,6 +65,7 @@ BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n") BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n") +BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n") BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n")
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits