Skip to content

Commit

Permalink
Enhance the robustness of multiple LIT tests
Browse files Browse the repository at this point in the history
Enhance the robustness of multiple LIT tests to minimize susceptibility
to trivial modifications
  • Loading branch information
aratajew authored and igcbot committed Feb 7, 2025
1 parent 8f34836 commit 16ee9ec
Show file tree
Hide file tree
Showing 4 changed files with 9 additions and 25 deletions.
4 changes: 2 additions & 2 deletions IGC/ocloc_tests/Builtins/intel_sub_group_shuffle.cl
Original file line number Diff line number Diff line change
Expand Up @@ -17,10 +17,10 @@ kernel void test_intel_sub_group_shuffle_immediate_index_simd32(global int* in,
size_t gid = get_global_id(0);
int x = in[gid];

// CHECK: mov (M5_NM, 1) simdShuffle(0,0)<1> V0039(1,15)<0;1,0>
// CHECK: mov (M5_NM, 1) simdShuffle(0,0)<1> {{V[0-9]+}}(1,15)<0;1,0>

// CHECK: mov (M1, 32) simdShuffleBroadcast(0,0)<1> simdShuffle(0,0)<0;1,0>
// CHECK: lsc_store.ugm (M1, 32) flat[V0041]:a64 simdShuffleBroadcast:d32
// CHECK: lsc_store.ugm (M1, 32) flat[{{V[0-9]+}}]:a64 simdShuffleBroadcast:d32
out[gid] = intel_sub_group_shuffle(x, 31);
}

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ kernel void test_sub_group_non_uniform_broadcast_non_immediate_sub_group_local_i
// CHECK: mov (M1, 16) simdShuffle(0,0)<1> r[A0(0),0]<1,0>:d
// CHECK: addr_add (M5, 16) A0(0)<1> &{{V[0-9]+}} ShuffleTmp(0,16)<1;1,0>
// CHECK: mov (M5, 16) simdShuffle(1,0)<1> r[A0(0),0]<1,0>:d
// CHECK: lsc_store.ugm (M1, 32) flat[V0046]:a64 simdShuffle:d32
// CHECK: lsc_store.ugm (M1, 32) flat[{{.+}}]:a64 simdShuffle:d32
bool isOddLane = get_sub_group_local_id() % 2 == 1;
if (isOddLane)
{
Expand Down
12 changes: 2 additions & 10 deletions IGC/ocloc_tests/features/fp64_conv_emu/fp64_conv_emu.cl
Original file line number Diff line number Diff line change
Expand Up @@ -14,18 +14,10 @@ SPDX-License-Identifier: MIT
// CHECK-LABEL: @conversion_kernel(
// CHECK-BASE: entry:
// CHECK-BASE: [[DPEmuFlag:%.*]] = alloca i32, align 4
// CHECK-BASE: [[TMP0:%.*]] = extractelement <8 x i32> %payloadHeader, i64 0
// CHECK-BASE: [[TMP1:%.*]] = extractelement <3 x i32> %enqueuedLocalSize, i64 0
// CHECK-BASE: [[TMP2:%.*]] = extractelement <8 x i32> %r0, i64 1
// CHECK-BASE: [[MUL:%.*]] = mul i32 [[TMP1]], [[TMP2]]
// CHECK-BASE: [[LOCAL_ID_X:%.*]] = zext i16 %localIdX to i32
// CHECK-BASE: [[ADD0:%.*]] = add i32 [[MUL]], [[LOCAL_ID_X]]
// CHECK-BASE: [[ADD1:%.*]] = add i32 [[ADD0]], [[TMP0]]
// CHECK-BASE: [[CONV0:%.*]] = zext i32 [[ADD1]] to i64
// CHECK-BASE: [[ARRAY_IDX0:%.*]] = getelementptr inbounds double, double addrspace(1)* %inA, i64 [[CONV0]]
// CHECK-BASE: [[ARRAY_IDX0:%.*]] = getelementptr inbounds double, double addrspace(1)* %inA, i64 %{{.*}}
// CHECK-BASE: [[TMP3:%.*]] = load double, double addrspace(1)* [[ARRAY_IDX0]], align 8
// CHECK-BASE: [[CALL_FTMP:%.*]] = call i32 @__igcbuiltin_dp_to_int32(double [[TMP3]], i32 3, i32 0, i32* [[DPEmuFlag]])
// CHECK-BASE: [[ARRAY_IDX2:%.*]] = getelementptr inbounds i32, i32 addrspace(1)* %out, i64 [[CONV0]]
// CHECK-BASE: [[ARRAY_IDX2:%.*]] = getelementptr inbounds i32, i32 addrspace(1)* %out, i64 %{{.*}}
// CHECK-BASE: store i32 [[CALL_FTMP]], i32 addrspace(1)* [[ARRAY_IDX2]], align 4
// CHECK-BASE: ret void

Expand Down
16 changes: 4 additions & 12 deletions IGC/ocloc_tests/features/fp64_conv_emu/fp64_conv_emu_fcmp.cl
Original file line number Diff line number Diff line change
Expand Up @@ -13,29 +13,21 @@ SPDX-License-Identifier: MIT

// CHECK-LABEL: @fcmp_kernel(
// CHECK-BASE: entry:
// CHECK-BASE: [[TMP0:%.*]] = extractelement <8 x i32> %payloadHeader, i64 0
// CHECK-BASE: [[TMP1:%.*]] = extractelement <3 x i32> %enqueuedLocalSize, i64 0
// CHECK-BASE: [[TMP2:%.*]] = extractelement <8 x i32> %r0, i64 1
// CHECK-BASE: [[MUL:%.*]] = mul i32 [[TMP1]], [[TMP2]]
// CHECK-BASE: [[LOCAL_ID_X:%.*]] = zext i16 %localIdX to i32
// CHECK-BASE: [[ADD0:%.*]] = add i32 [[MUL]], [[LOCAL_ID_X]]
// CHECK-BASE: [[ADD1:%.*]] = add i32 [[ADD0]], [[TMP0]]
// CHECK-BASE: [[CONV0:%.*]] = zext i32 [[ADD1]] to i64
// CHECK-BASE: [[ARRAY_IDX0:%.*]] = getelementptr inbounds double, double addrspace(1)* %inA, i64 [[CONV0]]
// CHECK-BASE: [[ARRAY_IDX0:%.*]] = getelementptr inbounds double, double addrspace(1)* %inA, i64 %{{.*}}
// CHECK-BASE: [[TMP3:%.*]] = load double, double addrspace(1)* [[ARRAY_IDX0]], align 8
// CHECK-BASE: [[ARRAY_IDX1:%.*]] = getelementptr inbounds double, double addrspace(1)* %inB, i64 [[CONV0]]
// CHECK-BASE: [[ARRAY_IDX1:%.*]] = getelementptr inbounds double, double addrspace(1)* %inB, i64 %{{.*}}
// CHECK-BASE: [[TMP4:%.*]] = load double, double addrspace(1)* [[ARRAY_IDX1]], align 8
// CHECK-BASE: [[CALL_FTMP:%.*]] = call i32 @__igcbuiltin_dp_cmp(double [[TMP3]], double [[TMP4]], i32 0)
// CHECK-BASE: [[SHL:%.*]] = shl i32 1, [[CALL_FTMP]]
// CHECK-BASE: [[AND:%.*]] = and i32 4, [[SHL]]
// CHECK-BASE: [[DPEmuCmp:%.*]] = icmp ne i32 [[AND]], 0
// CHECK-BASE: br i1 [[DPEmuCmp]], label %if.then, label %if.else
// CHECK-BASE: if.then:
// CHECK-BASE: [[ARRAY_IDX2:%.*]] = getelementptr inbounds double, double addrspace(1)* %out, i64 [[CONV0]]
// CHECK-BASE: [[ARRAY_IDX2:%.*]] = getelementptr inbounds double, double addrspace(1)* %out, i64 %{{.*}}
// CHECK-BASE: store double [[TMP3]], double addrspace(1)* [[ARRAY_IDX2]], align 8
// CHECK-BASE: br label %if.end
// CHECK-BASE: if.else:
// CHECK-BASE: [[ARRAY_IDX3:%.*]] = getelementptr inbounds double, double addrspace(1)* %out, i64 [[CONV0]]
// CHECK-BASE: [[ARRAY_IDX3:%.*]] = getelementptr inbounds double, double addrspace(1)* %out, i64 %{{.*}}
// CHECK-BASE: store double [[TMP4]], double addrspace(1)* [[ARRAY_IDX3]], align 8
// CHECK-BASE: br label %if.end
// CHECK-BASE: if.end:
Expand Down

0 comments on commit 16ee9ec

Please sign in to comment.