Skip to content

Commit 8d9130b

Browse files
easyonaaditeasyonaadit
committed
Using u32 inplace of i32 
1 parent 45ff803 commit 8d9130b

File tree

3 files changed

+48
-48
lines changed

3 files changed

+48
-48
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -357,17 +357,17 @@ BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi", "n")
357357

358358
//===----------------------------------------------------------------------===//
359359

360-
BUILTIN(__builtin_amdgcn_wave_reduce_add_i32, "ZiZiZi", "nc")
361-
BUILTIN(__builtin_amdgcn_wave_reduce_sub_i32, "ZiZiZi", "nc")
360+
BUILTIN(__builtin_amdgcn_wave_reduce_add_u32, "ZUiZUiZi", "nc")
361+
BUILTIN(__builtin_amdgcn_wave_reduce_sub_u32, "ZUiZUiZi", "nc")
362362
BUILTIN(__builtin_amdgcn_wave_reduce_min_i32, "ZiZiZi", "nc")
363363
BUILTIN(__builtin_amdgcn_wave_reduce_min_u32, "ZUiZUiZi", "nc")
364364
BUILTIN(__builtin_amdgcn_wave_reduce_max_i32, "ZiZiZi", "nc")
365365
BUILTIN(__builtin_amdgcn_wave_reduce_max_u32, "ZUiZUiZi", "nc")
366366
BUILTIN(__builtin_amdgcn_wave_reduce_and_b32, "ZiZiZi", "nc")
367367
BUILTIN(__builtin_amdgcn_wave_reduce_or_b32, "ZiZiZi", "nc")
368368
BUILTIN(__builtin_amdgcn_wave_reduce_xor_b32, "ZiZiZi", "nc")
369-
BUILTIN(__builtin_amdgcn_wave_reduce_add_i64, "WiWiZi", "nc")
370-
BUILTIN(__builtin_amdgcn_wave_reduce_sub_i64, "WiWiZi", "nc")
369+
BUILTIN(__builtin_amdgcn_wave_reduce_add_u64, "WUiWUiZi", "nc")
370+
BUILTIN(__builtin_amdgcn_wave_reduce_sub_u64, "WUiWUiZi", "nc")
371371
BUILTIN(__builtin_amdgcn_wave_reduce_min_i64, "WiWiZi", "nc")
372372
BUILTIN(__builtin_amdgcn_wave_reduce_min_u64, "WUiWUiZi", "nc")
373373
BUILTIN(__builtin_amdgcn_wave_reduce_max_i64, "WiWiZi", "nc")

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -299,11 +299,11 @@ static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID) {
299299
switch (BuiltinID) {
300300
default:
301301
llvm_unreachable("Unknown BuiltinID for wave reduction");
302-
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_i32:
303-
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_i64:
302+
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
303+
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
304304
return Intrinsic::amdgcn_wave_reduce_add;
305-
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_i32:
306-
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_i64:
305+
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
306+
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
307307
return Intrinsic::amdgcn_wave_reduce_sub;
308308
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
309309
case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
@@ -334,17 +334,17 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
334334
llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;
335335
llvm::SyncScope::ID SSID;
336336
switch (BuiltinID) {
337-
case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_i32:
338-
case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_i32:
337+
case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
338+
case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
339339
case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
340340
case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
341341
case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
342342
case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
343343
case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
344344
case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
345345
case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
346-
case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_i64:
347-
case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_i64:
346+
case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
347+
case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
348348
case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
349349
case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
350350
case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:

‎clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Lines changed: 36 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -398,88 +398,88 @@ void test_s_sendmsghalt_var(int in)
398398
__builtin_amdgcn_s_sendmsghalt(1, in);
399399
}
400400

401-
// CHECK-LABEL: @test_wave_reduce_add_i32_default
401+
// CHECK-LABEL: @test_wave_reduce_add_u32_default
402402
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(
403-
void test_wave_reduce_add_i32_default(global int* out, int in)
403+
void test_wave_reduce_add_u32_default(global int* out, int in)
404404
{
405-
*out = __builtin_amdgcn_wave_reduce_add_i32(in, 0);
405+
*out = __builtin_amdgcn_wave_reduce_add_u32(in, 0);
406406
}
407407

408-
// CHECK-LABEL: @test_wave_reduce_add_i64_default
408+
// CHECK-LABEL: @test_wave_reduce_add_u64_default
409409
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(
410-
void test_wave_reduce_add_i64_default(global int* out, long in)
410+
void test_wave_reduce_add_u64_default(global int* out, long in)
411411
{
412-
*out = __builtin_amdgcn_wave_reduce_add_i64(in, 0);
412+
*out = __builtin_amdgcn_wave_reduce_add_u64(in, 0);
413413
}
414414

415-
// CHECK-LABEL: @test_wave_reduce_add_i32_iterative
415+
// CHECK-LABEL: @test_wave_reduce_add_u32_iterative
416416
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(
417-
void test_wave_reduce_add_i32_iterative(global int* out, int in)
417+
void test_wave_reduce_add_u32_iterative(global int* out, int in)
418418
{
419-
*out = __builtin_amdgcn_wave_reduce_add_i32(in, 1);
419+
*out = __builtin_amdgcn_wave_reduce_add_u32(in, 1);
420420
}
421421

422-
// CHECK-LABEL: @test_wave_reduce_add_i64_iterative
422+
// CHECK-LABEL: @test_wave_reduce_add_u64_iterative
423423
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(
424-
void test_wave_reduce_add_i64_iterative(global int* out, long in)
424+
void test_wave_reduce_add_u64_iterative(global int* out, long in)
425425
{
426-
*out = __builtin_amdgcn_wave_reduce_add_i64(in, 1);
426+
*out = __builtin_amdgcn_wave_reduce_add_u64(in, 1);
427427
}
428428

429-
// CHECK-LABEL: @test_wave_reduce_add_i32_dpp
429+
// CHECK-LABEL: @test_wave_reduce_add_u32_dpp
430430
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(
431-
void test_wave_reduce_add_i32_dpp(global int* out, int in)
431+
void test_wave_reduce_add_u32_dpp(global int* out, int in)
432432
{
433-
*out = __builtin_amdgcn_wave_reduce_add_i32(in, 2);
433+
*out = __builtin_amdgcn_wave_reduce_add_u32(in, 2);
434434
}
435435

436-
// CHECK-LABEL: @test_wave_reduce_add_i64_dpp
436+
// CHECK-LABEL: @test_wave_reduce_add_u64_dpp
437437
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(
438-
void test_wave_reduce_add_i64_dpp(global int* out, long in)
438+
void test_wave_reduce_add_u64_dpp(global int* out, long in)
439439
{
440-
*out = __builtin_amdgcn_wave_reduce_add_i64(in, 2);
440+
*out = __builtin_amdgcn_wave_reduce_add_u64(in, 2);
441441
}
442442

443-
// CHECK-LABEL: @test_wave_reduce_sub_i32_default
443+
// CHECK-LABEL: @test_wave_reduce_sub_u32_default
444444
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(
445-
void test_wave_reduce_sub_i32_default(global int* out, int in)
445+
void test_wave_reduce_sub_u32_default(global int* out, int in)
446446
{
447-
*out = __builtin_amdgcn_wave_reduce_sub_i32(in, 0);
447+
*out = __builtin_amdgcn_wave_reduce_sub_u32(in, 0);
448448
}
449449

450-
// CHECK-LABEL: @test_wave_reduce_sub_i64_default
450+
// CHECK-LABEL: @test_wave_reduce_sub_u64_default
451451
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(
452-
void test_wave_reduce_sub_i64_default(global int* out, long in)
452+
void test_wave_reduce_sub_u64_default(global int* out, long in)
453453
{
454-
*out = __builtin_amdgcn_wave_reduce_sub_i64(in, 0);
454+
*out = __builtin_amdgcn_wave_reduce_sub_u64(in, 0);
455455
}
456456

457-
// CHECK-LABEL: @test_wave_reduce_sub_i32_iterative
457+
// CHECK-LABEL: @test_wave_reduce_sub_u32_iterative
458458
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(
459-
void test_wave_reduce_sub_i32_iterative(global int* out, int in)
459+
void test_wave_reduce_sub_u32_iterative(global int* out, int in)
460460
{
461-
*out = __builtin_amdgcn_wave_reduce_sub_i32(in, 1);
461+
*out = __builtin_amdgcn_wave_reduce_sub_u32(in, 1);
462462
}
463463

464-
// CHECK-LABEL: @test_wave_reduce_sub_i64_iterative
464+
// CHECK-LABEL: @test_wave_reduce_sub_u64_iterative
465465
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(
466-
void test_wave_reduce_sub_i64_iterative(global int* out, long in)
466+
void test_wave_reduce_sub_u64_iterative(global int* out, long in)
467467
{
468-
*out = __builtin_amdgcn_wave_reduce_sub_i64(in, 1);
468+
*out = __builtin_amdgcn_wave_reduce_sub_u64(in, 1);
469469
}
470470

471-
// CHECK-LABEL: @test_wave_reduce_sub_i32_dpp
471+
// CHECK-LABEL: @test_wave_reduce_sub_u32_dpp
472472
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(
473-
void test_wave_reduce_sub_i32_dpp(global int* out, int in)
473+
void test_wave_reduce_sub_u32_dpp(global int* out, int in)
474474
{
475-
*out = __builtin_amdgcn_wave_reduce_sub_i32(in, 2);
475+
*out = __builtin_amdgcn_wave_reduce_sub_u32(in, 2);
476476
}
477477

478-
// CHECK-LABEL: @test_wave_reduce_sub_i64_dpp
478+
// CHECK-LABEL: @test_wave_reduce_sub_u64_dpp
479479
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(
480-
void test_wave_reduce_sub_i64_dpp(global int* out, long in)
480+
void test_wave_reduce_sub_u64_dpp(global int* out, long in)
481481
{
482-
*out = __builtin_amdgcn_wave_reduce_sub_i64(in, 2);
482+
*out = __builtin_amdgcn_wave_reduce_sub_u64(in, 2);
483483
}
484484

485485
// CHECK-LABEL: @test_wave_reduce_and_b32_default

0 commit comments

Comments
 (0)