From 57d56a9ca2a0ce8e44b32eabf2b247b52f8e94db Mon Sep 17 00:00:00 2001 From: haonanya Date: Fri, 24 Dec 2021 13:52:48 +0800 Subject: [PATCH] Align cl_khr_mipmap_image_writes support with community and remove outdated patch 1. clang/0001-Add-support-for-cl_khr_mipmap_image_writes.patch is backporting of https://reviews.llvm.org/D71460 2. Remove spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch, since SPIRV has it Signed-off-by: haonanya --- ...pport-for-cl_khr_mipmap_image_writes.patch | 37 +++--- ...of-work-item-builtin-translation-745.patch | 114 ------------------ 2 files changed, 18 insertions(+), 133 deletions(-) delete mode 100644 patches/spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch diff --git a/patches/clang/0001-Add-support-for-cl_khr_mipmap_image_writes.patch b/patches/clang/0001-Add-support-for-cl_khr_mipmap_image_writes.patch index fcdfc71c..78ebfb13 100644 --- a/patches/clang/0001-Add-support-for-cl_khr_mipmap_image_writes.patch +++ b/patches/clang/0001-Add-support-for-cl_khr_mipmap_image_writes.patch @@ -1,16 +1,19 @@ -From 1ed55ebdc8181c67b9cf1cfa3a977f069413dd89 Mon Sep 17 00:00:00 2001 +From 9625ef3fea5f587558965a2a40dd51f8fee9eb05 Mon Sep 17 00:00:00 2001 From: Ilya Mashkov Date: Thu, 12 Dec 2019 14:33:41 +0300 Subject: [PATCH] Add support for cl_khr_mipmap_image_writes +This is backporting of https://reviews.llvm.org/D71460 + +Signed-off-by: haonanya --- clang/include/clang/Basic/OpenCLExtensions.def | 1 + - clang/lib/Headers/opencl-c.h | 22 ++++++++++++++-------- + clang/lib/Headers/opencl-c.h | 18 ++++++++++-------- clang/test/SemaOpenCL/extension-version.cl | 12 ++++++++++++ - 3 files changed, 27 insertions(+), 8 deletions(-) + 3 files changed, 23 insertions(+), 8 deletions(-) diff --git a/clang/include/clang/Basic/OpenCLExtensions.def b/clang/include/clang/Basic/OpenCLExtensions.def -index 5536a6e..5174815 100644 +index 5536a6e8e4df..517481584313 100644 --- a/clang/include/clang/Basic/OpenCLExtensions.def +++ b/clang/include/clang/Basic/OpenCLExtensions.def @@ -70,6 +70,7 @@ OPENCLEXT_INTERNAL(cl_khr_spir, 120, ~0U) @@ -22,20 +25,19 @@ index 5536a6e..5174815 100644 OPENCLEXT_INTERNAL(cl_khr_subgroups, 200, ~0U) OPENCLEXT_INTERNAL(cl_khr_terminate_context, 200, ~0U) diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h -index 06c5ab6..d3ae0dd 100644 +index 06c5ab6a72f0..1985a361168b 100644 --- a/clang/lib/Headers/opencl-c.h +++ b/clang/lib/Headers/opencl-c.h -@@ -14682,7 +14682,8 @@ void __ovld write_imagef(write_only image2d_array_depth_t image, int4 coord, flo +@@ -14682,7 +14682,7 @@ void __ovld write_imagef(write_only image2d_array_depth_t image, int4 coord, flo // OpenCL Extension v2.0 s9.18 - Mipmaps #if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0) -#ifdef cl_khr_mipmap_image +#if defined(cl_khr_mipmap_image_writes) -+#pragma OPENCL EXTENSION cl_khr_mipmap_image_writes : begin void __ovld write_imagef(write_only image1d_t image, int coord, int lod, float4 color); void __ovld write_imagei(write_only image1d_t image, int coord, int lod, int4 color); void __ovld write_imageui(write_only image1d_t image, int coord, int lod, uint4 color); -@@ -14699,15 +14700,17 @@ void __ovld write_imagef(write_only image2d_array_t image_array, int4 coord, int +@@ -14699,15 +14699,16 @@ void __ovld write_imagef(write_only image2d_array_t image_array, int4 coord, int void __ovld write_imagei(write_only image2d_array_t image_array, int4 coord, int lod, int4 color); void __ovld write_imageui(write_only image2d_array_t image_array, int4 coord, int lod, uint4 color); @@ -52,22 +54,20 @@ index 06c5ab6..d3ae0dd 100644 -#endif //cl_khr_mipmap_image +#endif //cl_khr_3d_image_writes + -+#pragma OPENCL EXTENSION cl_khr_mipmap_image_writes : end +#endif //defined(cl_khr_mipmap_image_writes) #endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0) // Image write functions for half4 type -@@ -14756,7 +14759,8 @@ void __ovld write_imagef(read_write image2d_array_depth_t image, int4 coord, flo +@@ -14756,7 +14757,7 @@ void __ovld write_imagef(read_write image2d_array_depth_t image, int4 coord, flo #endif //cl_khr_depth_images #if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0) -#ifdef cl_khr_mipmap_image +#ifdef cl_khr_mipmap_image_writes -+#pragma OPENCL EXTENSION cl_khr_mipmap_image_writes : begin void __ovld write_imagef(read_write image1d_t image, int coord, int lod, float4 color); void __ovld write_imagei(read_write image1d_t image, int coord, int lod, int4 color); void __ovld write_imageui(read_write image1d_t image, int coord, int lod, uint4 color); -@@ -14780,8 +14784,10 @@ void __ovld write_imagef(read_write image2d_array_depth_t image, int4 coord, int +@@ -14780,8 +14781,9 @@ void __ovld write_imagef(read_write image2d_array_depth_t image, int4 coord, int void __ovld write_imagef(read_write image3d_t image, int4 coord, int lod, float4 color); void __ovld write_imagei(read_write image3d_t image, int4 coord, int lod, int4 color); void __ovld write_imageui(read_write image3d_t image, int4 coord, int lod, uint4 color); @@ -75,19 +75,19 @@ index 06c5ab6..d3ae0dd 100644 -#endif //cl_khr_mipmap_image +#endif //cl_khr_3d_image_writes + -+#pragma OPENCL EXTENSION cl_khr_mipmap_image_writes : end +#endif //cl_khr_mipmap_image_writes #endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0) // Image write functions for half4 type diff --git a/clang/test/SemaOpenCL/extension-version.cl b/clang/test/SemaOpenCL/extension-version.cl -index 19d0884..0e6bbb7 100644 +index 19d088495350..0e6bbb7d3bcd 100644 --- a/clang/test/SemaOpenCL/extension-version.cl +++ b/clang/test/SemaOpenCL/extension-version.cl -@@ -243,6 +243,18 @@ +@@ -242,6 +242,18 @@ + #endif #pragma OPENCL EXTENSION cl_khr_mipmap_image : enable - #if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) ++#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) +#ifndef cl_khr_mipmap_image_writes +#error "Missing cl_khr_mipmap_image_writes define" +#endif @@ -99,10 +99,9 @@ index 19d0884..0e6bbb7 100644 +#endif +#pragma OPENCL EXTENSION cl_khr_mipmap_image_writes : enable + -+#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) + #if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) #ifndef cl_khr_srgb_image_writes #error "Missing cl_khr_srgb_image_writes define" - #endif -- -2.7.4 +2.17.1 diff --git a/patches/spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch b/patches/spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch deleted file mode 100644 index 1e48d9e6..00000000 --- a/patches/spirv/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch +++ /dev/null @@ -1,114 +0,0 @@ -From a1029c9c1f768ea3a4ddf13d47fe97e2342cac0e Mon Sep 17 00:00:00 2001 -From: haonanya -Date: Wed, 11 Aug 2021 18:47:42 +0800 -Subject: [PATCH] Fix debug info of work-item builtin translation - -Signed-off-by: haonanya ---- - lib/SPIRV/OCL20ToSPIRV.cpp | 5 ++- - lib/SPIRV/SPIRVReader.cpp | 1 + - test/DebugInfo/builtin-get-global-id.ll | 60 +++++++++++++++++++++++++ - 3 files changed, 65 insertions(+), 1 deletion(-) - create mode 100644 test/DebugInfo/builtin-get-global-id.ll - -diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp -index 1262c48c..a742c8cf 100644 ---- a/lib/SPIRV/OCL20ToSPIRV.cpp -+++ b/lib/SPIRV/OCL20ToSPIRV.cpp -@@ -1297,11 +1297,14 @@ void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() { - for (auto UI = I.user_begin(), UE = I.user_end(); UI != UE; ++UI) { - auto CI = dyn_cast(*UI); - assert(CI && "invalid instruction"); -- Value *NewValue = new LoadInst(BV, "", CI); -+ const DebugLoc &DLoc = CI->getDebugLoc(); -+ Instruction *NewValue = new LoadInst(BV, "", CI); -+ NewValue->setDebugLoc(DLoc); - LLVM_DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n'); - if (IsVec) { - NewValue = - ExtractElementInst::Create(NewValue, CI->getArgOperand(0), "", CI); -+ NewValue->setDebugLoc(DLoc); - LLVM_DEBUG(dbgs() << *NewValue << '\n'); - } - NewValue->takeName(CI); -diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp -index 97cf9410..429d1584 100644 ---- a/lib/SPIRV/SPIRVReader.cpp -+++ b/lib/SPIRV/SPIRVReader.cpp -@@ -305,6 +305,7 @@ bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV, - auto Replace = [&](std::vector Arg, Instruction *I) { - auto Call = CallInst::Create(Func, Arg, "", I); - Call->takeName(I); -+ Call->setDebugLoc(I->getDebugLoc()); - setAttrByCalledFunc(Call); - SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << *Call - << '\n';) -diff --git a/test/DebugInfo/builtin-get-global-id.ll b/test/DebugInfo/builtin-get-global-id.ll -new file mode 100644 -index 00000000..a4a00e63 ---- /dev/null -+++ b/test/DebugInfo/builtin-get-global-id.ll -@@ -0,0 +1,60 @@ -+; Check debug info of builtin get_global_id is preserved from LLVM IR to spirv -+; and spirv to LLVM IR translation. -+ -+; Original .cl source: -+; kernel void test() { -+; size_t gid = get_global_id(0); -+; } -+ -+; Command line: -+; ./clang -cc1 1.cl -triple spir64 -cl-std=cl2.0 -emit-llvm -finclude-default-header -debug-info-kind=line-tables-only -O0 -+ -+; RUN: llvm-as %s -o %t.bc -+; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix CHECK-SPIRV -+; RUN: llvm-spirv %t.bc -o %t.spv -+; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s -+ -+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" -+target triple = "spir64" -+ -+; CHECK-SPIRV: ExtInst {{.*}} DebugScope -+; CHECK-SPIRV-NEXT: Line {{[0-9]+}} 2 16 -+; CHECK-SPIRV-NEXT: Load {{[0-9]+}} [[LoadRes:[0-9]+]] -+; CHECK-SPIRV-NEXT: CompositeExtract {{[0-9]+}} {{[0-9]+}} [[LoadRes]] 0 -+ -+; Function Attrs: convergent noinline norecurse nounwind optnone -+define spir_kernel void @test() #0 !dbg !7 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 { -+entry: -+ %gid = alloca i64, align 8 -+ %call = call spir_func i64 @_Z13get_global_idj(i32 0) #2, !dbg !10 -+; CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0) #1, !dbg [[DBG:![0-9]+]] -+ store i64 %call, i64* %gid, align 8, !dbg !11 -+ ret void, !dbg !12 -+} -+ -+; Function Attrs: convergent nounwind readnone -+declare spir_func i64 @_Z13get_global_idj(i32) #1 -+ -+attributes #0 = { convergent noinline norecurse nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #1 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #2 = { convergent nounwind readnone } -+ -+!llvm.dbg.cu = !{!0} -+!llvm.module.flags = !{!3, !4} -+!opencl.ocl.version = !{!5} -+!opencl.spir.version = !{!5} -+!llvm.ident = !{!6} -+ -+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2, nameTableKind: None) -+!1 = !DIFile(filename: "", directory: "") -+!2 = !{} -+!3 = !{i32 2, !"Debug Info Version", i32 3} -+!4 = !{i32 1, !"wchar_size", i32 4} -+!5 = !{i32 2, i32 0} -+!6 = !{!"clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)"} -+!7 = distinct !DISubprogram(name: "test", scope: !8, file: !8, line: 1, type: !9, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0, retainedNodes: !2) -+!8 = !DIFile(filename: "1.cl", directory: "") -+!9 = !DISubroutineType(types: !2) -+!10 = !DILocation(line: 2, column: 16, scope: !7) -+!11 = !DILocation(line: 2, column: 10, scope: !7) -+!12 = !DILocation(line: 3, column: 1, scope: !7) --- -2.17.1 -