diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index 050f110c6e365..56d4631f7845e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -355,6 +355,13 @@ class OpenACCRoutineClauseEmitter final curValue.getZExtValue()); } } + + void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) { + lastDeviceTypeValues.clear(); + + for (const DeviceTypeArgument &arg : clause.getArchitectures()) + lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo())); + } }; } // namespace diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index 2d4ed23a46d1c..8e7384ae66d8e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -112,19 +112,6 @@ class OpenACCClauseCIREmitter final return createConstantInt(cgf.cgm.getLoc(loc), width, value); } - mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) { - // '*' case leaves no identifier-info, just a nullptr. - if (!ii) - return mlir::acc::DeviceType::Star; - return llvm::StringSwitch(ii->getName()) - .CaseLower("default", mlir::acc::DeviceType::Default) - .CaseLower("host", mlir::acc::DeviceType::Host) - .CaseLower("multicore", mlir::acc::DeviceType::Multicore) - .CasesLower({"nvidia", "acc_device_nvidia"}, - mlir::acc::DeviceType::Nvidia) - .CaseLower("radeon", mlir::acc::DeviceType::Radeon); - } - mlir::acc::GangArgType decodeGangType(OpenACCGangKind gk) { switch (gk) { case OpenACCGangKind::Num: diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h index 5bcc9f57d67b1..639d14804087e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h @@ -40,4 +40,17 @@ convertOpenACCModifiers(OpenACCModifierKind modifiers) { mlirModifiers = mlirModifiers | static_cast(modifiers); return mlirModifiers; } + +inline mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) { + // '*' case leaves no identifier-info, just a nullptr. + if (!ii) + return mlir::acc::DeviceType::Star; + return llvm::StringSwitch(ii->getName()) + .CaseLower("default", mlir::acc::DeviceType::Default) + .CaseLower("host", mlir::acc::DeviceType::Host) + .CaseLower("multicore", mlir::acc::DeviceType::Multicore) + .CasesLower({"nvidia", "acc_device_nvidia"}, + mlir::acc::DeviceType::Nvidia) + .CaseLower("radeon", mlir::acc::DeviceType::Radeon); +} } // namespace clang::CIRGen diff --git a/clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp b/clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp new file mode 100644 index 0000000000000..61c985bd81f56 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp @@ -0,0 +1,79 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +#pragma acc routine nohost device_type(nvidia, radeon) seq +void Func1() {} +void Func2() {} +#pragma acc routine(Func2) device_type(radeon) seq + +#pragma acc routine device_type(multicore) worker device_type(nvidia, radeon) seq +void Func3() {} +void Func4() {} +#pragma acc routine(Func4) device_type(nvidia) seq device_type(radeon) vector + +#pragma acc routine device_type(multicore) gang device_type(nvidia, radeon) gang +void Func5() {} +void Func6() {} +#pragma acc routine(Func6) device_type(multicore) gang(dim:1) device_type(radeon) gang + +#pragma acc routine device_type(host) gang device_type(nvidia, radeon) gang(dim:1) +void Func7() {} +void Func8() {} +#pragma acc routine(Func8) device_type(radeon) gang(dim:2) + +#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3) +void Func9() {} +void Func10() {} +#pragma acc routine(Func10) device_type(nvidia) gang device_type(radeon) gang(dim:3) + +#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3) device_type(multicore) gang +void Func11() {} +void Func12() {} +#pragma acc routine(Func12) device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3) + +#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang +void Func13() {} +void Func14() {} +#pragma acc routine(Func14) device_type(nvidia) gang(dim:2) device_type(radeon) gang + +// CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq ([#acc.device_type, #acc.device_type]) nohost + +// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>} + +// CHECK: cir.func{{.*}} @[[F3_NAME:.*Func3[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F3_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F3_R_NAME]] func(@[[F3_NAME]]) worker ([#acc.device_type]) seq ([#acc.device_type, #acc.device_type]) + +// CHECK: cir.func{{.*}} @[[F4_NAME:.*Func4[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F4_R_NAME:.*]]]>} + +// CHECK: cir.func{{.*}} @[[F5_NAME:.*Func5[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F5_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F5_R_NAME]] func(@[[F5_NAME]]) gang([#acc.device_type, #acc.device_type, #acc.device_type]) + +// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>} +// +// CHECK: cir.func{{.*}} @[[F7_NAME:.*Func7[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F7_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F7_R_NAME]] func(@[[F7_NAME]]) gang([#acc.device_type], dim: 1 : i64 [#acc.device_type], dim: 1 : i64 [#acc.device_type]) + +// CHECK: cir.func{{.*}} @[[F8_NAME:.*Func8[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F8_R_NAME:.*]]]>} +// +// CHECK: cir.func{{.*}} @[[F9_NAME:.*Func9[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F9_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F9_R_NAME]] func(@[[F9_NAME]]) gang(dim: 2 : i64 [#acc.device_type], dim: 3 : i64 [#acc.device_type]) +// +// CHECK: cir.func{{.*}} @[[F10_NAME:.*Func10[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F10_R_NAME:.*]]]>} + +// CHECK: cir.func{{.*}} @[[F11_NAME:.*Func11[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F11_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F11_R_NAME]] func(@[[F11_NAME]]) gang([#acc.device_type], dim: 2 : i64 [#acc.device_type], dim: 3 : i64 [#acc.device_type]) +// +// CHECK: cir.func{{.*}} @[[F12_NAME:.*Func12[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F12_R_NAME:.*]]]>} +// +// CHECK: cir.func{{.*}} @[[F13_NAME:.*Func13[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F13_R_NAME:.*]]]>} +// CHECK: acc.routine @[[F13_R_NAME]] func(@[[F13_NAME]]) gang([#acc.device_type], dim: 2 : i64 [#acc.device_type]) +// +// CHECK: cir.func{{.*}} @[[F14_NAME:.*Func14[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F14_R_NAME:.*]]]>} + +// CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq ([#acc.device_type]) +// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) vector ([#acc.device_type]) seq ([#acc.device_type]) +// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) gang([#acc.device_type], dim: 1 : i64 [#acc.device_type]) +// CHECK: acc.routine @[[F8_R_NAME]] func(@[[F8_NAME]]) gang(dim: 2 : i64 [#acc.device_type]) +// CHECK: acc.routine @[[F10_R_NAME]] func(@[[F10_NAME]]) gang([#acc.device_type], dim: 3 : i64 [#acc.device_type]) +// CHECK: acc.routine @[[F12_R_NAME]] func(@[[F12_NAME]]) gang(dim: 2 : i64 [#acc.device_type], dim: 3 : i64 [#acc.device_type]) +// CHECK: acc.routine @[[F14_R_NAME]] func(@[[F14_NAME]]) gang([#acc.device_type], dim: 2 : i64 [#acc.device_type])