derek.holloway
7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
53 lines
2.3 KiB
MLIR
53 lines
2.3 KiB
MLIR
// RUN: mlir-opt -allow-unregistered-dialect -pass-pipeline="builtin.module(func.func(test-mapping-to-processing-elements))" %s | FileCheck %s
|
|
|
|
// CHECK: #[[mul_map:.+]] = affine_map<()[s0, s1] -> (s0 * s1)>
|
|
// CHECK: #[[add_map:.+]] = affine_map<()[s0, s1] -> (s0 + s1)>
|
|
|
|
// CHECK: func @map1d
|
|
// CHECK-SAME: (%[[lb:.*]]: index, %[[ub:.*]]: index, %[[step:.*]]: index)
|
|
func.func @map1d(%lb: index, %ub: index, %step: index) {
|
|
// CHECK: %[[threads:.*]]:2 = "new_processor_id_and_range"() : () -> (index, index)
|
|
%0:2 = "new_processor_id_and_range"() : () -> (index, index)
|
|
|
|
// CHECK: %[[thread_offset:.+]] = affine.apply #[[mul_map]]()[%[[threads]]#0, %[[step]]]
|
|
// CHECK: %[[new_lb:.+]] = affine.apply #[[add_map]]()[%[[thread_offset]], %[[lb]]]
|
|
// CHECK: %[[new_step:.+]] = affine.apply #[[mul_map]]()[%[[threads]]#1, %[[step]]]
|
|
|
|
// CHECK: scf.for %{{.*}} = %[[new_lb]] to %[[ub]] step %[[new_step]] {
|
|
scf.for %i = %lb to %ub step %step {}
|
|
return
|
|
}
|
|
|
|
// CHECK: func @map2d
|
|
// CHECK-SAME: (%[[lb:.*]]: index, %[[ub:.*]]: index, %[[step:.*]]: index)
|
|
func.func @map2d(%lb : index, %ub : index, %step : index) {
|
|
// CHECK: %[[blocks:.*]]:2 = "new_processor_id_and_range"() : () -> (index, index)
|
|
%0:2 = "new_processor_id_and_range"() : () -> (index, index)
|
|
|
|
// CHECK: %[[threads:.*]]:2 = "new_processor_id_and_range"() : () -> (index, index)
|
|
%1:2 = "new_processor_id_and_range"() : () -> (index, index)
|
|
|
|
// blockIdx.x * blockDim.x
|
|
// CHECK: %[[bidxXbdimx:.+]] = affine.apply #[[mul_map]]()[%[[blocks]]#0, %[[threads]]#1]
|
|
//
|
|
// threadIdx.x + blockIdx.x * blockDim.x
|
|
// CHECK: %[[tidxpbidxXbdimx:.+]] = affine.apply #[[add_map]]()[%[[bidxXbdimx]], %[[threads]]#0]
|
|
//
|
|
// thread_offset = step * (threadIdx.x + blockIdx.x * blockDim.x)
|
|
// CHECK: %[[thread_offset:.+]] = affine.apply #[[mul_map]]()[%[[tidxpbidxXbdimx]], %[[step]]]
|
|
//
|
|
// new_lb = lb + thread_offset
|
|
// CHECK: %[[new_lb:.+]] = affine.apply #[[add_map]]()[%[[thread_offset]], %[[lb]]]
|
|
//
|
|
// stepXgdimx = step * gridDim.x
|
|
// CHECK: %[[stepXgdimx:.+]] = affine.apply #[[mul_map]]()[%[[blocks]]#1, %[[step]]]
|
|
//
|
|
// new_step = step * gridDim.x * blockDim.x
|
|
// CHECK: %[[new_step:.+]] = affine.apply #[[mul_map]]()[%[[threads]]#1, %[[stepXgdimx]]]
|
|
//
|
|
// CHECK: scf.for %{{.*}} = %[[new_lb]] to %[[ub]] step %[[new_step]] {
|
|
|
|
scf.for %i = %lb to %ub step %step {}
|
|
return
|
|
}
|