[AMDGPU] Add amdgcn builtin for s_prefetch_inst (#211642)
This commit is contained in:
Gang Chen
GitHub
parent
3d20ba7135
commit
92cdb9dc83
11 changed files with 147 additions and 6 deletions
|
|
@ -761,8 +761,12 @@ def __builtin_amdgcn_s_barrier_join : AMDGPUBuiltin<"void(void *)", [], "gfx12-i
|
|||
def __builtin_amdgcn_s_barrier_leave : AMDGPUBuiltin<"void(_Constant short)", [], "gfx12-insts">;
|
||||
def __builtin_amdgcn_s_get_barrier_state : AMDGPUBuiltin<"unsigned int(int)", [], "gfx12-insts">;
|
||||
def __builtin_amdgcn_s_get_named_barrier_state : AMDGPUBuiltin<"unsigned int(void *)", [], "gfx12-insts">;
|
||||
def __builtin_amdgcn_s_prefetch_data : AMDGPUBuiltin<"void(void const *, unsigned int)", [Const], "gfx12-insts">;
|
||||
def __builtin_amdgcn_s_buffer_prefetch_data : AMDGPUBuiltin<"void(__amdgpu_buffer_rsrc_t, _Constant int, unsigned int)", [Const], "gfx12-insts">;
|
||||
def __builtin_amdgcn_s_prefetch_data : AMDGPUBuiltin<"void(void const *, unsigned int)", [Const], "smem-prefetch-insts">;
|
||||
def __builtin_amdgcn_s_prefetch_inst : AMDGPUBuiltin<"void(void const *, unsigned int)", [Const], "smem-prefetch-insts"> {
|
||||
let Documentation = [DocPrefetchInst_GFX12];
|
||||
let ArgNames = ["ptr", "len"];
|
||||
}
|
||||
def __builtin_amdgcn_s_buffer_prefetch_data : AMDGPUBuiltin<"void(__amdgpu_buffer_rsrc_t, _Constant int, unsigned int)", [Const], "smem-prefetch-insts">;
|
||||
|
||||
def __builtin_amdgcn_global_load_tr_b64_v2i32 : AMDGPUBuiltin<"_ExtVector<2, int>(_ExtVector<2, int> address_space<1> *)", [Const], "gfx12-insts,wavefrontsize32">;
|
||||
def __builtin_amdgcn_global_load_tr_b128_v8i16 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<8, short> address_space<1> *)", [Const], "gfx12-insts,wavefrontsize32">;
|
||||
|
|
|
|||
|
|
@ -78,6 +78,13 @@ between standard floating-point types, integer types and packed low-precision fo
|
|||
}];
|
||||
}
|
||||
|
||||
def DocCatAMDGPUPrefetch : DocumentationCategory<"Prefetch Builtins"> {
|
||||
let Content = [{
|
||||
These builtins provide access to AMDGPU prefetch instructions, including instructions
|
||||
for prefetching data and instructions.
|
||||
}];
|
||||
}
|
||||
|
||||
//===----------------------------------------------------------------------===//
|
||||
// ABI / Special Register Builtins — Documentation records
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
|
@ -793,3 +800,16 @@ using the scale factor ``scale``, then converts the values to a packed
|
|||
integers.
|
||||
}];
|
||||
}
|
||||
|
||||
//===----------------------------------------------------------------------===//
|
||||
// Prefetch Builtins
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
||||
def DocPrefetchInst_GFX12 : Documentation {
|
||||
let Category = DocCatAMDGPUPrefetch;
|
||||
let Content = [{
|
||||
Prefetch instructions into I-Cache from the address given by ``ptr`` pointing
|
||||
to instruction memory, with the length specified by ``len``, which should be
|
||||
0-31 (1-32 chunks, units of 128 bytes).
|
||||
}];
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1004,6 +1004,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
|
|||
getContext().BuiltinInfo.getName(builtinId));
|
||||
return mlir::Value{};
|
||||
}
|
||||
case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst: {
|
||||
cgm.errorNYI(expr->getSourceRange(),
|
||||
std::string("unimplemented AMDGPU builtin call: ") +
|
||||
getContext().BuiltinInfo.getName(builtinId));
|
||||
return mlir::Value{};
|
||||
}
|
||||
case Builtin::BIlogbf:
|
||||
case Builtin::BI__builtin_logbf:
|
||||
return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle());
|
||||
|
|
|
|||
|
|
@ -2197,6 +2197,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
|
|||
case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
|
||||
return emitBuiltinWithOneOverloadedType<2>(
|
||||
*this, E, Intrinsic::amdgcn_s_prefetch_data);
|
||||
case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst:
|
||||
return emitBuiltinWithOneOverloadedType<2>(
|
||||
*this, E, Intrinsic::amdgcn_s_prefetch_inst);
|
||||
case Builtin::BIlogbf:
|
||||
case Builtin::BI__builtin_logbf: {
|
||||
Value *Src0 = EmitScalarExpr(E->getArg(0));
|
||||
|
|
|
|||
|
|
@ -0,0 +1,10 @@
|
|||
// REQUIRES: amdgpu-registered-target
|
||||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
|
||||
// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-cir %s -verify -o %t.cir
|
||||
|
||||
#define __device__ __attribute__((device))
|
||||
|
||||
// expected-error@+2 {{ClangIR code gen Not Yet Implemented: unimplemented AMDGPU builtin call: __builtin_amdgcn_s_prefetch_inst}}
|
||||
__device__ void test_s_prefetch_inst(const void *p, unsigned int len) {
|
||||
__builtin_amdgcn_s_prefetch_inst(p, len);
|
||||
}
|
||||
|
|
@ -59,7 +59,7 @@ void foo() {
|
|||
// AMDGCN-GFX1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
|
||||
// AMDGCN-GFX1010: attributes #[[ATTR2]] = { noreturn nounwind }
|
||||
//.
|
||||
// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind }
|
||||
|
|
|
|||
|
|
@ -68,7 +68,7 @@ void foo() {
|
|||
//.
|
||||
// AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
//.
|
||||
// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
|
||||
// AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind }
|
||||
|
|
|
|||
|
|
@ -107,9 +107,9 @@ const B& f(A *a) {
|
|||
// CHECK: attributes #[[ATTR3]] = { nounwind }
|
||||
// CHECK: attributes #[[ATTR4]] = { noreturn }
|
||||
//.
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) }
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+async-load-to-lds-insts,+async-store-from-lds-insts,+asynccnt,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+f32-to-fp6bf6-cvt-scale-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx13-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+smem-prefetch-insts,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind }
|
||||
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn }
|
||||
//.
|
||||
|
|
|
|||
71
clang/test/CodeGenHIP/builtins-amdgcn-prefetch.hip
Normal file
71
clang/test/CodeGenHIP/builtins-amdgcn-prefetch.hip
Normal file
|
|
@ -0,0 +1,71 @@
|
|||
// 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);
|
||||
}
|
||||
|
|
@ -247,6 +247,32 @@ void test_s_prefetch_data(int *fp, global float *gp, constant char *cp, unsigned
|
|||
__builtin_amdgcn_s_prefetch_data(cp, 31);
|
||||
}
|
||||
|
||||
// CHECK-LABEL: @test_s_prefetch_inst(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[FP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
||||
// CHECK-NEXT: [[GP_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
|
||||
// CHECK-NEXT: [[CP_ADDR:%.*]] = alloca ptr addrspace(4), align 8, addrspace(5)
|
||||
// CHECK-NEXT: [[LEN_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
||||
// CHECK-NEXT: store ptr [[FP:%.*]], ptr addrspace(5) [[FP_ADDR]], align 8
|
||||
// CHECK-NEXT: store ptr addrspace(1) [[GP:%.*]], ptr addrspace(5) [[GP_ADDR]], align 8
|
||||
// CHECK-NEXT: store ptr addrspace(4) [[CP:%.*]], ptr addrspace(5) [[CP_ADDR]], align 8
|
||||
// CHECK-NEXT: store i32 [[LEN:%.*]], ptr addrspace(5) [[LEN_ADDR]], align 4
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[FP_ADDR]], align 8
|
||||
// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.inst.p0(ptr [[TMP0]], i32 0)
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[GP_ADDR]], align 8
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[LEN_ADDR]], align 4
|
||||
// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) [[TMP1]], i32 [[TMP2]])
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(5) [[CP_ADDR]], align 8
|
||||
// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) [[TMP3]], i32 31)
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
void test_s_prefetch_inst(int *fp, global float *gp, constant char *cp, unsigned int len)
|
||||
{
|
||||
__builtin_amdgcn_s_prefetch_inst(fp, 0);
|
||||
__builtin_amdgcn_s_prefetch_inst(gp, len);
|
||||
__builtin_amdgcn_s_prefetch_inst(cp, 31);
|
||||
}
|
||||
|
||||
// CHECK-LABEL: @test_s_buffer_prefetch_data(
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[RSRC_ADDR:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)
|
||||
|
|
|
|||
|
|
@ -639,6 +639,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T,
|
|||
Features["wmma-128b-insts"] = true;
|
||||
Features["swmmac-gfx1200-insts"] = true;
|
||||
Features["atomic-fmin-fmax-global-f32"] = true;
|
||||
Features["smem-prefetch-insts"] = true;
|
||||
break;
|
||||
case GK_GFX1170:
|
||||
case GK_GFX1171:
|
||||
|
|
|
|||
Loading…
Reference in a new issue