Skip to content

Commit

Permalink
Merge pull request #101 from frasercrmck/combine-lit-test
Browse files Browse the repository at this point in the history 
[compiler] Combine two lit tests into one
  • Loading branch information
@frasercrmck
frasercrmck authored Aug 24, 2023
2 parents 38ec5f8 + 933c90d commit 1ca64b9
Show file tree
Hide file tree
Showing 3 changed files with 51 additions and 80 deletions.
8 changes: 4 additions & 4 deletions modules/compiler/test/lit/passes/barriers-cfg-reduce.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
;
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

; RUN: muxc --passes barriers-pass -S %s | FileCheck %s
; RUN: muxc --passes barriers-pass,verify < %s | FileCheck %s

target triple = "spir64-unknown-unknown"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
Expand All @@ -28,9 +28,9 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:
; CHECK: [[REDUCE_LOOP]]:
; CHECK: %[[IDX:.+]] = phi i64 [ 0, %sw.bb2 ], [ %[[IDX_NEXT:.+]], %[[REDUCE_LOOP]] ]
; CHECK: %[[ACCUM:.+]] = phi i32 [ 0, %sw.bb2 ], [ %[[ACCUM_NEXT:.+]], %[[REDUCE_LOOP]] ]
; CHECK: %[[ITEM:.+]] = getelementptr inbounds %reduction_live_mem_info, {{ptr|.+\*}} %live_variables, i64 %[[IDX]]
; CHECK: %[[VAL:.+]] = getelementptr inbounds %reduction_live_mem_info, {{ptr|.+\*}} %[[ITEM]], i32 0, i32 0
; CHECK: %[[LD:.+]] = load i32, {{ptr|.+\*}} %[[VAL]], align 4
; CHECK: %[[ITEM:.+]] = getelementptr inbounds %reduction_live_mem_info, ptr %live_variables, i64 %[[IDX]]
; CHECK: %[[VAL:.+]] = getelementptr inbounds %reduction_live_mem_info, ptr %[[ITEM]], i32 0, i32 0
; CHECK: %[[LD:.+]] = load i32, ptr %[[VAL]], align 4
; CHECK: %[[ACCUM_NEXT]] = add i32 %[[ACCUM]], %[[LD]]
; CHECK: %[[IDX_NEXT]] = add i64 %[[IDX]], 1
; CHECK: %[[LOOP_COND:.+]] = icmp ult i64 %24, 262144
Expand Down
40 changes: 0 additions & 40 deletions modules/compiler/test/lit/passes/replace-wgc-scans-only.ll
View file

This file was deleted.

83 changes: 47 additions & 36 deletions modules/compiler/test/lit/passes/replace-wgc.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,57 +14,66 @@
;
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

; RUN: muxc --passes replace-wgc,verify -S %s | FileCheck %s
; RUN: muxc --passes replace-wgc,verify < %s \
; RUN: | FileCheck %s --check-prefixes CHECK,CHECK-ALL
; RUN: muxc --passes "replace-wgc<scans-only>,verify" < %s \
; RUN: | FileCheck %s --check-prefixes CHECK,CHECK-SCANS-ONLY

; Check that the replace-wgc correctly defines the work-group collective functions
; Check that the replace-wgc correctly defines the work-group collective
; functions, optionally only defining the scans and leaving others intact.

target triple = "spir64-unknown-unknown"
target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"

; CHECK: @[[MINI_ACCUM:.+]] = internal addrspace(3) global i32 undef
; CHECK-ALL: @__mux_work_group_reduce_smin_i32.accumulator = internal addrspace(3) global i32 undef
; CHECK: @__mux_work_group_scan_inclusive_umax_i32.accumulator = internal addrspace(3) global i32 undef
; CHECK: @__mux_work_group_scan_exclusive_fadd_f32.accumulator = internal addrspace(3) global float undef
; CHECK: @__mux_work_group_broadcast_i32.accumulator = internal addrspace(3) global i32 undef
; CHECK-ALL: @__mux_work_group_broadcast_i32.accumulator = internal addrspace(3) global i32 undef

; CHECK: define spir_func i32 @__mux_work_group_reduce_smin_i32(i32 %id, i32 [[PARAM:%.*]])
; If this isn't a scan we shouldn't have defined it
; CHECK-SCANS-ONLY: declare spir_func i32 @__mux_work_group_reduce_smin_i32(i32, i32)
declare spir_func i32 @__mux_work_group_reduce_smin_i32(i32 %id, i32 %x)
; CHECK-LABEL: entry:
; CHECK: %[[SUBGROUP:.+]] = call i32 @__mux_sub_group_reduce_smin_i32(i32 %{{.+}})
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_ONCE:#[0-9]+]]
; CHECK: store i32 2147483647, [[PTR_i32:(i32 addrspace\(3\)\*)|(ptr addrspace\(3\))]] @[[MINI_ACCUM]]
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK: %[[CURRVAL:.+]] = load i32, [[PTR_i32]] @[[MINI_ACCUM]]
; CHECK: %[[ACCUM:.*]] = call i32 @llvm.smin.i32(i32 %[[CURRVAL]], i32 %[[SUBGROUP]])
; CHECK: store i32 %[[ACCUM]], [[PTR_i32]] @[[MINI_ACCUM]]
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK: %[[RESULT:.*]] = load i32, [[PTR_i32]] @[[MINI_ACCUM]]
; CHECK: ret i32 %[[RESULT]]
; CHECK-ALL: define spir_func i32 @__mux_work_group_reduce_smin_i32(i32 %id, i32 [[PARAM:%.*]])
; CHECK-ALL-LABEL: entry:
; CHECK-ALL: %[[SUBGROUP:.+]] = call i32 @__mux_sub_group_reduce_smin_i32(i32 %{{.+}})
; CHECK-ALL: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_ONCE:#[0-9]+]]
; CHECK-ALL: store i32 2147483647, ptr addrspace(3) @__mux_work_group_reduce_smin_i32.accumulator
; CHECK-ALL: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK-ALL: %[[CURRVAL:.+]] = load i32, ptr addrspace(3) @__mux_work_group_reduce_smin_i32.accumulator
; CHECK-ALL: %[[ACCUM:.*]] = call i32 @llvm.smin.i32(i32 %[[CURRVAL]], i32 %[[SUBGROUP]])
; CHECK-ALL: store i32 %[[ACCUM]], ptr addrspace(3) @__mux_work_group_reduce_smin_i32.accumulator
; CHECK-ALL: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK-ALL: %[[RESULT:.*]] = load i32, ptr addrspace(3) @__mux_work_group_reduce_smin_i32.accumulator
; CHECK-ALL: ret i32 %[[RESULT]]


; CHECK: define spir_func i32 @__mux_work_group_scan_inclusive_umax_i32(i32 %id, i32 [[PARAM:%.*]])
declare spir_func i32 @__mux_work_group_scan_inclusive_umax_i32(i32 %id, i32 %x)
; CHECK: define spir_func i32 @__mux_work_group_scan_inclusive_umax_i32(i32 %id, i32 [[PARAM:%.*]])
; CHECK-LABEL: entry:
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_ONCE]]
; CHECK: store i32 0, [[PTR_i32]] @__mux_work_group_scan_inclusive_umax_i32.accumulator
; This is just to ensure SCHEDULE_ONCE is defined on all paths...
; CHECK-ALL: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_ONCE]]
; CHECK-SCANS-ONLY: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_ONCE:#[0-9]+]]
; CHECK: store i32 0, ptr addrspace(3) @__mux_work_group_scan_inclusive_umax_i32.accumulator
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_LINEAR:#[0-9]+]]
; CHECK: %[[CURRVAL:.+]] = load i32, [[PTR_i32]] @__mux_work_group_scan_inclusive_umax_i32.accumulator
; CHECK: %[[CURRVAL:.+]] = load i32, ptr addrspace(3) @__mux_work_group_scan_inclusive_umax_i32.accumulator
; CHECK: %[[SCAN:.+]] = call i32 @__mux_sub_group_scan_inclusive_umax_i32(i32 %x)
; CHECK: %[[RESULT:.+]] = call i32 @llvm.umax.i32(i32 %[[CURRVAL]], i32 %[[SCAN]])
; CHECK: %[[SIZE:.+]] = call i32 @__mux_get_sub_group_size()
; CHECK: %[[LAST:.+]] = sub nuw i32 %[[SIZE]], 1
; CHECK: %[[TAIL:.+]] = call i32 @__mux_sub_group_broadcast_i32(i32 %[[SCAN]], i32 %[[LAST]])
; CHECK: %[[ACCUM:.+]] = call i32 @llvm.umax.i32(i32 %[[CURRVAL]], i32 %[[TAIL]])
; CHECK: store i32 %[[ACCUM]], [[PTR_i32]] @__mux_work_group_scan_inclusive_umax_i32.accumulator
; CHECK: store i32 %[[ACCUM]], ptr addrspace(3) @__mux_work_group_scan_inclusive_umax_i32.accumulator
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK: ret i32 %[[RESULT]]


; CHECK: define spir_func float @__mux_work_group_scan_exclusive_fadd_f32(i32 %id, float [[PARAM:%.*]])
declare spir_func float @__mux_work_group_scan_exclusive_fadd_f32(i32 %id, float %x)
; CHECK: define spir_func float @__mux_work_group_scan_exclusive_fadd_f32(i32 %id, float [[PARAM:%.*]])
; CHECK-LABEL: entry:
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_ONCE]]
; CHECK: store float -0.000000e+00, [[PTR_float:(float addrspace\(3\)\*)|(ptr addrspace\(3\))]] @__mux_work_group_scan_exclusive_fadd_f32.accumulator
; CHECK: store float -0.000000e+00, ptr addrspace(3) @__mux_work_group_scan_exclusive_fadd_f32.accumulator
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272) [[SCHEDULE_LINEAR]]
; CHECK: %[[CURRVAL:.+]] = load float, [[PTR_float]] @__mux_work_group_scan_exclusive_fadd_f32.accumulator
; CHECK: %[[CURRVAL:.+]] = load float, ptr addrspace(3) @__mux_work_group_scan_exclusive_fadd_f32.accumulator
; CHECK: %[[SGSCAN:.+]] = call float @__mux_sub_group_scan_exclusive_fadd_f32(float %x)
; CHECK: %[[SGID:.+]] = call i32 @__mux_get_sub_group_local_id()
; CHECK: %[[CMPID:.+]] = icmp eq i32 %[[SGID]], 0
Expand All @@ -76,7 +85,7 @@ declare spir_func float @__mux_work_group_scan_exclusive_fadd_f32(i32 %id, float
; CHECK: %[[TAIL:.+]] = call float @__mux_sub_group_broadcast_f32(float %x, i32 %[[LAST]])
; CHECK: %[[TOTAL:.+]] = fadd float %[[SCAN_TAIL]], %[[TAIL]]
; CHECK: %[[ACCUM:.+]] = fadd float %[[CURRVAL]], %[[TOTAL]]
; CHECK: store float %[[ACCUM]], [[PTR_float]] @__mux_work_group_scan_exclusive_fadd_f32.accumulator
; CHECK: store float %[[ACCUM]], ptr addrspace(3) @__mux_work_group_scan_exclusive_fadd_f32.accumulator
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK: %[[LIDX:.+]] = call spir_func i64 @__mux_get_local_id(i32 0)
; CHECK: %[[CMPX:.+]] = icmp eq i64 %[[LIDX]], 0
Expand All @@ -90,22 +99,24 @@ declare spir_func float @__mux_work_group_scan_exclusive_fadd_f32(i32 %id, float
; CHECK: ret float %[[RESULT]]


; CHECK: define spir_func i32 @__mux_work_group_broadcast_i32(i32 %barrier_id, i32 [[PARAM:%.*]], i64 {{%.*}}, i64 {{%.*}}, i64 {{%.*}})
; CHECK-SCANS-ONLY: declare spir_func i32 @__mux_work_group_broadcast_i32(i32, i32, i64, i64, i64)
declare spir_func i32 @__mux_work_group_broadcast_i32(i32 %barrier_id, i32 %x, i64 %idx, i64 %idy, i64 %idz)
; CHECK-LABEL: entry:
; CHECK: call i64 @__mux_get_local_id(i32 0)
; CHECK-ALL: define spir_func i32 @__mux_work_group_broadcast_i32(i32 %barrier_id, i32 [[PARAM:%.*]], i64 {{%.*}}, i64 {{%.*}}, i64 {{%.*}})
; CHECK-ALL-LABEL: entry:
; CHECK-ALL: call i64 @__mux_get_local_id(i32 0)

; CHECK-LABEL: broadcast:
; CHECK: store i32 [[PARAM]], ptr addrspace(3) @__mux_work_group_broadcast_i32.accumulator
; CHECK-ALL-LABEL: broadcast:
; CHECK-ALL: store i32 [[PARAM]], ptr addrspace(3) @__mux_work_group_broadcast_i32.accumulator

; CHECK-ALL-LABEL: exit:
; CHECK-ALL: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK-ALL: [[RESULT:%.*]] = load i32, ptr addrspace(3) @__mux_work_group_broadcast_i32.accumulator
; CHECK-ALL: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK-ALL: ret i32 [[RESULT]]

; CHECK-LABEL: exit:
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK: [[RESULT:%.*]] = load i32, ptr addrspace(3) @__mux_work_group_broadcast_i32.accumulator
; CHECK: call void @__mux_work_group_barrier(i32 0, i32 2, i32 272)
; CHECK: ret i32 [[RESULT]]

; CHECK: define spir_func half @__mux_work_group_scan_exclusive_fadd_f16(i32 %id, half [[PARAM:%.*]])
declare spir_func half @__mux_work_group_scan_exclusive_fadd_f16(i32 %id, half %x)
; CHECK: define spir_func half @__mux_work_group_scan_exclusive_fadd_f16(i32 %id, half [[PARAM:%.*]])

; CHECK-DAG: attributes [[SCHEDULE_ONCE]] = { "mux-barrier-schedule"="once" }
; CHECK-DAG: attributes [[SCHEDULE_LINEAR]] = { "mux-barrier-schedule"="linear" }
Expand Down

0 comments on commit 1ca64b9

Please sign in to comment.