[PROTON] Enable Proton dialect for XPU backend and run test_record.py in CI (#3942)

Signed-off-by: Anatoly Myachev <[email protected]>

Author: anmyachev

.github/workflows/build-test-reusable.yml

@@ -176,7 +176,7 @@ jobs:
         if: ${{ inputs.driver_version == 'rolling' && inputs.device == 'max1100' }}
         run: |
           cd third_party/proton/test
-          pytest test_api.py test_lib.py test_profile.py test_viewer.py -s -v
+          pytest test_api.py test_lib.py test_profile.py test_viewer.py test_record.py -s -v
           cd ..
 
       - name: Run unit tests

third_party/intel/lib/TritonIntelGPUToLLVM/CMakeLists.txt

@@ -34,4 +34,5 @@ add_triton_library(TritonIntelGPUToLLVM
     TritonGENToLLVM
     TritonIntelGPUIR
     TritonIntelUtils
+    TritonProtonToLLVM
 )

third_party/intel/lib/TritonIntelGPUToLLVM/PipelineManager.h

@@ -30,6 +30,8 @@
 
 #include "PatternTritonGPUOpToLLVM.h"
 
+#include "third_party/proton/dialect/include/TritonProtonToLLVM/PatternTritonProtonOpToLLVM.h"
+
 namespace mlir {
 
 FailureOr<LLVM::LLVMFuncOp>
@@ -231,6 +233,8 @@ class TritonGPUToLLVMPipelineManager {
                                     benefit);
       mlir::triton::populateMemoryOpToLLVMPatterns(typeConverter, targetInfo,
                                                    patterns, benefit);
+      mlir::triton::proton::populateRecordOpToLLVMPattern(
+          typeConverter, patterns, targetInfo, benefit);
       intel::populateControlFlowOpToLLVMPattern(typeConverter, patterns,
                                                 targetInfo, benefit);
       mlir::triton::populateMakeRangeOpToLLVMPattern(typeConverter, targetInfo,

third_party/proton/test/test_record.py

@@ -29,8 +29,8 @@ def add_kernel(
 
     torch.manual_seed(0)
     size = 2**12
-    x = torch.rand(size, device='cuda')
-    y = torch.rand(size, device='cuda')
+    x = torch.rand(size, device='xpu')
+    y = torch.rand(size, device='xpu')
     output = torch.empty_like(x)
     n_elements = output.numel()
     grid = (1, 1, 1)