68 lines · plain
1// RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv="use-64bit-index=true" -verify-diagnostics -split-input-file %s -o - | FileCheck %s2 3module attributes {4 gpu.container_module,5 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>6} {7 gpu.module @kernels {8 // CHECK-LABEL: spirv.module @{{.*}} Physical64 OpenCL9 // CHECK: spirv.func10 // CHECK-SAME: {{%.*}}: f3211 // CHECK-NOT: spirv.interface_var_abi12 // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup>13 // CHECK-NOT: spirv.interface_var_abi14 // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>15 // CHECK-LABEL: func.func @basic_module_structure16 // CHECK-SAME: attributes {gpu.kernel}17 gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel18 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {19 gpu.return20 }21 }22 23 func.func @main() {24 %0 = "op"() : () -> (f32)25 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)26 %cst = arith.constant 1 : index27 gpu.launch_func @kernels::@basic_module_structure28 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)29 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)30 return31 }32}33 34// -----35 36module attributes {37 gpu.container_module38} {39 gpu.module @kernels attributes {40 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>41 } {42 // CHECK-LABEL: spirv.module @{{.*}} Physical64 OpenCL43 // CHECK-SAME: spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>44 // CHECK: spirv.func45 // CHECK-SAME: {{%.*}}: f3246 // CHECK-NOT: spirv.interface_var_abi47 // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup>48 // CHECK-NOT: spirv.interface_var_abi49 // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>50 // CHECK-LABEL: func.func @basic_module_structure51 // CHECK-SAME: attributes {gpu.kernel}52 gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel53 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {54 gpu.return55 }56 }57 58 func.func @main() {59 %0 = "op"() : () -> (f32)60 %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)61 %cst = arith.constant 1 : index62 gpu.launch_func @kernels::@basic_module_structure63 blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)64 args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)65 return66 }67}68