Skip to content

Commit d4fbab3

Browse files
authored
1 parent 56d2262 commit d4fbab3

File tree

3 files changed

+13
-0
lines changed

3 files changed

+13
-0
lines changed

clang/lib/CIR/CodeGen/CIRGenModule.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -547,6 +547,7 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() {
547547
return *(theTargetCIRGenInfo = createSPIRVTargetCIRGenInfo(genTypes));
548548
}
549549

550+
case llvm::Triple::nvptx:
550551
case llvm::Triple::nvptx64: {
551552
return *(theTargetCIRGenInfo = createNVPTXTargetCIRGenInfo(genTypes));
552553
}

clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -81,8 +81,11 @@ createTargetLoweringInfo(LowerModule &LM) {
8181
}
8282
case llvm::Triple::spirv64:
8383
return createSPIRVTargetLoweringInfo(LM);
84+
85+
case llvm::Triple::nvptx:
8486
case llvm::Triple::nvptx64:
8587
return createNVPTXTargetLoweringInfo(LM);
88+
8689
default:
8790
cir_cconv_unreachable("ABI NYI");
8891
}
Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
#include "../Inputs/cuda.h"
2+
3+
// RUN: %clang_cc1 -triple nvptx -fclangir \
4+
// RUN: -fcuda-is-device -emit-cir -target-sdk-version=12.3 \
5+
// RUN: %s -o %t.cir
6+
// RUN: FileCheck --input-file=%t.cir %s
7+
8+
__device__ void device_fn(int* a, double b, float c) {}
9+
// CHECK: cir.func @_Z9device_fnPidf

0 commit comments

Comments
 (0)