clang/test/CodeGenHIP/builtins-amdgcn-prefetch.hip

// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -emit-llvm -fcuda-is-device -o - %s | FileCheck %s --check-prefix=CHECK-GFX1200

#define __device__ __attribute__((device))
#define __shared__ __attribute__((shared))
#define __constant__ __attribute__((constant))

extern "C" __device__ int bar(const char *s);

// CHECK-GFX1200-LABEL: define dso_local noundef i32 @_Z4foo1v(
// CHECK-GFX1200-SAME: ) #[[ATTR0:[0-9]+]] {
// CHECK-GFX1200-NEXT:  [[ENTRY:.*:]]
// CHECK-GFX1200-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-GFX1200-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
// CHECK-GFX1200-NEXT:    call void @llvm.amdgcn.s.prefetch.inst.p0(ptr @bar, i32 0)
// CHECK-GFX1200-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8
// CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
// CHECK-GFX1200-NEXT:    [[CALL:%.*]] = call i32 @bar(ptr noundef [[TMP0]]) #[[ATTR3:[0-9]+]]
// CHECK-GFX1200-NEXT:    ret i32 [[CALL]]
//
__device__ int foo1() {
  __builtin_amdgcn_s_prefetch_inst((const void *)bar, 0);
  const char *s = "hello world";
  return bar(s);
}

// CHECK-GFX1200-LABEL: define dso_local noundef i32 @_Z4foo2i(
// CHECK-GFX1200-SAME: i32 noundef [[ID:%.*]]) #[[ATTR0]] {
// CHECK-GFX1200-NEXT:  [[ENTRY:.*:]]
// CHECK-GFX1200-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-GFX1200-NEXT:    [[ID_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-GFX1200-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-GFX1200-NEXT:    [[S2:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-GFX1200-NEXT:    [[ID_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[ID_ADDR]] to ptr
// CHECK-GFX1200-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
// CHECK-GFX1200-NEXT:    [[S2_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S2]] to ptr
// CHECK-GFX1200-NEXT:    store i32 [[ID]], ptr [[ID_ADDR_ASCAST]], align 4
// CHECK-GFX1200-NEXT:    call void @llvm.amdgcn.s.prefetch.inst.p0(ptr blockaddress(@_Z4foo2i, %[[NOBAR:.*]]), i32 0)
// CHECK-GFX1200-NEXT:    [[TMP0:%.*]] = load i32, ptr [[ID_ADDR_ASCAST]], align 4
// CHECK-GFX1200-NEXT:    [[CMP:%.*]] = icmp eq i32 [[TMP0]], 0
// CHECK-GFX1200-NEXT:    br i1 [[CMP]], label %[[IF_THEN:.*]], label %[[IF_END:.*]]
// CHECK-GFX1200:       [[IF_THEN]]:
// CHECK-GFX1200-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8
// CHECK-GFX1200-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
// CHECK-GFX1200-NEXT:    [[CALL:%.*]] = call i32 @bar(ptr noundef [[TMP1]]) #[[ATTR3]]
// CHECK-GFX1200-NEXT:    store i32 [[CALL]], ptr addrspace(5) [[RETVAL]], align 4
// CHECK-GFX1200-NEXT:    br label %[[RETURN:.*]]
// CHECK-GFX1200:       [[IF_END]]:
// CHECK-GFX1200-NEXT:    br label %[[NOBAR]]
// CHECK-GFX1200:       [[NOBAR]]:
// CHECK-GFX1200-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr), ptr [[S2_ASCAST]], align 8
// CHECK-GFX1200-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[S2_ASCAST]], align 8
// CHECK-GFX1200-NEXT:    [[CALL1:%.*]] = call i32 @bar(ptr noundef [[TMP2]]) #[[ATTR3]]
// CHECK-GFX1200-NEXT:    store i32 [[CALL1]], ptr addrspace(5) [[RETVAL]], align 4
// CHECK-GFX1200-NEXT:    br label %[[RETURN]]
// CHECK-GFX1200:       [[RETURN]]:
// CHECK-GFX1200-NEXT:    [[TMP3:%.*]] = load i32, ptr addrspace(5) [[RETVAL]], align 4
// CHECK-GFX1200-NEXT:    ret i32 [[TMP3]]
// CHECK-GFX1200:       [[INDIRECTGOTO:.*:]]
// CHECK-GFX1200-NEXT:    indirectbr ptr poison, [label %[[NOBAR]]]
//
__device__ int foo2(int id) {
  __builtin_amdgcn_s_prefetch_inst(&&NOBAR, 0);
  if (id == 0) {
    const char *s = "hello world";
    return bar(s);
  }
NOBAR:
  const char *s2 = "skip hello";
  return bar(s2);
}