diff options
| author | Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com> | 2025-08-12 15:07:07 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-08-12 15:07:07 -0700 |
| commit | d0ee82040cb22ae38c92eb83b0c9ae71ca51a517 (patch) | |
| tree | 6f46ad7f7dd4aa37e6e8e2abc00cbadeb786bfa3 /clang/test/CodeGenOpenCL | |
| parent | 8710571aba2bacf73bcde58270f9b166c66f32e5 (diff) | |
| download | llvm-d0ee82040cb22ae38c92eb83b0c9ae71ca51a517.zip llvm-d0ee82040cb22ae38c92eb83b0c9ae71ca51a517.tar.gz llvm-d0ee82040cb22ae38c92eb83b0c9ae71ca51a517.tar.bz2 | |
[AMDGPU] Add s_barrier_init|join|leave instructions (#153296)
Diffstat (limited to 'clang/test/CodeGenOpenCL')
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl | 7 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl | 44 |
2 files changed, 51 insertions, 0 deletions
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl index 5d86a9b..1a50433 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl @@ -23,6 +23,13 @@ kernel void builtins_amdgcn_s_barrier_signal_isfirst_err(global int* in, global *out = *in; } +kernel void builtins_amdgcn_s_barrier_leave_err(global int* in, global int* out, int barrier) { + + __builtin_amdgcn_s_barrier_signal(-1); + __builtin_amdgcn_s_barrier_leave(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_leave' must be a constant integer}} + *out = *in; +} + void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, unsigned int off) { __builtin_amdgcn_s_buffer_prefetch_data(rsrc, off, 31); // expected-error {{'__builtin_amdgcn_s_buffer_prefetch_data' must be a constant integer}} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl index f764128..8c02616 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl @@ -139,6 +139,50 @@ void test_s_barrier_signal_isfirst(int* a, int* b, int *c) __builtin_amdgcn_s_barrier_wait(1); } +// CHECK-LABEL: @test_s_barrier_init( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr +// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr +// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) +// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4 +// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) [[TMP1]], i32 [[TMP2]]) +// CHECK-NEXT: ret void +// +void test_s_barrier_init(void *bar, int a) +{ + __builtin_amdgcn_s_barrier_init(bar, a); +} + +// CHECK-LABEL: @test_s_barrier_join( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr +// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3) +// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) [[TMP1]]) +// CHECK-NEXT: ret void +// +void test_s_barrier_join(void *bar) +{ + __builtin_amdgcn_s_barrier_join(bar); +} + +// CHECK-LABEL: @test_s_barrier_leave( +// CHECK-NEXT: entry: +// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.leave(i16 1) +// CHECK-NEXT: ret void +// +void test_s_barrier_leave() +{ + __builtin_amdgcn_s_barrier_leave(1); +} + // CHECK-LABEL: @test_s_get_barrier_state( // CHECK-NEXT: entry: // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5) |