Skip to content

Commit a275939

Browse files
authored
clang_delta: support HIP (.hip) inputs on LLVM 17+ (#489)
clang_delta cannot parse files with .hip extension, the PR extends the parsing logic. Otherwise, clang_delta fails with an Unsupported file type! error message.
1 parent 6f0dea6 commit a275939

File tree

5 files changed

+42
-0
lines changed

5 files changed

+42
-0
lines changed

clang_delta/CMakeLists.txt

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -87,6 +87,8 @@ set(SOURCE_FILES
8787
"/tests/empty-struct-to-int/test2.output"
8888
"/tests/empty-struct-to-int/test3.c"
8989
"/tests/empty-struct-to-int/test3.output"
90+
"/tests/hip/kernel.hip"
91+
"/tests/hip/kernel.output"
9092
"/tests/instantiate-template-param/default_param.cc"
9193
"/tests/instantiate-template-param/default_param.output"
9294
"/tests/instantiate-template-param/test1.cc"

clang_delta/TransformationManager.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -187,6 +187,18 @@ bool TransformationManager::initializeCompilerInstance(std::string &ErrorMsg)
187187
Language::OpenCL,
188188
T, includes);
189189
}
190+
#if LLVM_VERSION_MAJOR >= 17
191+
else if (IK.getLanguage() == Language::HIP) {
192+
std::vector<const char *> Args;
193+
Args.push_back("-x");
194+
Args.push_back("hip");
195+
196+
CompilerInvocation::CreateFromArgs(Invocation, Args,
197+
ClangInstance->getDiagnostics());
198+
LangOptions::setLangDefaults(ClangInstance->getLangOpts(), Language::HIP,
199+
T, includes, LSTD);
200+
}
201+
#endif
190202
else {
191203
ErrorMsg = "Unsupported file type!";
192204
return false;

clang_delta/tests/hip/kernel.hip

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
int a1[10];
2+
3+
void foo() {
4+
char a;
5+
int t = 1;
6+
((char (*)[t]) a)[0][0] = 0;
7+
a1[1] = 1;
8+
}
9+
10+
__global__ void k() {}
Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
int a1[10];
2+
int a1_1;
3+
4+
void foo() {
5+
char a;
6+
int t = 1;
7+
((char (*)[t]) a)[0][0] = 0;
8+
a1_1 = 1;
9+
}
10+
11+
__global__ void k() {}

clang_delta/tests/test_clang_delta.py

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1308,6 +1308,13 @@ def test_simplify_simple_recursive_template(self):
13081308
'--transformation=simplify-recursive-template-instantiation --counter=1',
13091309
)
13101310

1311+
@unittest.skipIf(get_clang_version() < 17, 'HIP input requires LLVM 17+')
1312+
def test_hip_kernel_aggregate_to_scalar(self):
1313+
self.check_clang_delta(
1314+
'hip/kernel.hip',
1315+
'--transformation=aggregate-to-scalar --counter=1',
1316+
)
1317+
13111318
def test_template_arg_to_int_not_valid_5(self):
13121319
self.check_query_instances(
13131320
'template-arg-to-int/not_valid5.cc',

0 commit comments

Comments
 (0)