Skip to content

Commit db0ef3a

Browse files
lamb-jlamb-j
committed
[Comgr][v3.0] Remove HIP_COMPILER tests and HIP dependencies from tests
Previously, specific tests were only run when HIP_COMPILER was set at build time. This led to them not being run by any CI jobs, and updating the CI to include a HIP would create a circular dependency between HIP and Comgr. Instead, with this change we remove the HIP header dependencies from the Comgr tests, and now run all tests during a regular test invocation. Change-Id: Icd4eabee74c0093910da4c6eccd6ca441955642e
1 parent 08d47ec commit db0ef3a

File tree

11 files changed

+43
-234
lines changed

11 files changed

+43
-234
lines changed

‎amd/comgr/docs/ReleaseNotes.md‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -199,6 +199,9 @@ where reporters provide Comgr logs.
199199
- Refactor nested kernel behavior into new test, as this behavior is less common
200200
and shouldn't be featured in the baseline tests
201201
- Add metadata parsing tests for code objects with multiple AMDGPU metadata note entries.
202+
- Updated Comgr HIP test to not rely on HIP\_COMPILER being set, or a valid HIP
203+
installation. We can test the functionality of Comgr HIP compilation without
204+
directly relying on HIP
202205

203206
New Targets
204207
-----------

‎amd/comgr/test/CMakeLists.txt‎

Lines changed: 11 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -119,6 +119,12 @@ add_test_input_binary(linking-empty source/linking/empty.cl source/linking/e
119119
add_test_input_linked(multiple-note-records source/linking/kernel0.o source/linking/kernel1.o source/multiple-note-records.out -w)
120120
add_test_input_linked(multiple-note-records-one-kernel source/linking/kernel0.o source/linking/empty.o source/multiple-note-records-one-kernel.out -w)
121121

122+
add_test_input_bitcode_bundle(square source/square.hip source/square.bc)
123+
add_test_input_object_file_bundle(double source/double.hip source/double.o)
124+
125+
add_test_input_bitcode_bundle(cube source/cube.hip source/cube.bc)
126+
add_test_archive(cube_archive cube source/cube.bc source/cube.a)
127+
122128
configure_file("source/linking/kernel0.cl" "source/linking/kernel0.cl" COPYONLY)
123129
configure_file("source/linking/kernel1.cl" "source/linking/kernel1.cl" COPYONLY)
124130
configure_file("source/linking/empty.cl" "source/linking/empty.cl" COPYONLY)
@@ -132,7 +138,6 @@ configure_file("source/include-macro.h" "source/include-macro.h" COPYONLY)
132138
configure_file("source/include-nested.h" "source/include-nested.h" COPYONLY)
133139
configure_file("source/source1.s" "source/source1.s" COPYONLY)
134140
configure_file("source/source1.hip" "source/source1.hip" COPYONLY)
135-
configure_file("source/source2.hip" "source/source2.hip" COPYONLY)
136141
configure_file("source/name-expression.hip" "source/name-expression.hip" COPYONLY)
137142
configure_file("source/rocm56slice.b" "source/rocm56slice.b" COPYONLY)
138143
configure_file("source/rocm57slice.b" "source/rocm57slice.b" COPYONLY)
@@ -177,10 +182,6 @@ if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
177182
PRIVATE -D_CRT_SECURE_NO_WARNINGS)
178183
endif()
179184

180-
if (DEFINED HIP_COMPILER AND "${HIP_COMPILER}" STREQUAL "clang")
181-
target_compile_definitions("${name}"
182-
PRIVATE -DHIP_COMPILER=${HIP_COMPILER})
183-
endif()
184185
target_link_libraries("${name}"
185186
amd_comgr)
186187
add_dependencies("${name}"
@@ -231,18 +232,8 @@ add_comgr_test(multithread_test cpp)
231232
add_comgr_test(nested_kernel_test c)
232233
add_comgr_test(map_elf_virtual_address_test c)
233234
add_comgr_test(compile_source_to_executable c)
234-
235-
# Test : Compile HIP tests only if HIP-Clang is installed.
236-
if (DEFINED HIP_COMPILER AND "${HIP_COMPILER}" STREQUAL "clang")
237-
add_test_input_bitcode_bundle(square source/square.hip source/square.bc)
238-
add_test_input_object_file_bundle(double source/double.hip source/double.o)
239-
240-
add_test_input_bitcode_bundle(cube source/cube.hip source/cube.bc)
241-
add_test_archive(cube_archive cube source/cube.bc source/cube.a)
242-
243-
add_comgr_test(compile_hip_test c)
244-
add_comgr_test(unbundle_hip_test c)
245-
add_comgr_test(compile_hip_to_relocatable c)
246-
add_comgr_test(mangled_names_hip_test c)
247-
add_comgr_test(name_expression_map_test c)
248-
endif()
235+
add_comgr_test(name_expression_map_test c)
236+
add_comgr_test(compile_hip_test c)
237+
add_comgr_test(compile_hip_to_relocatable c)
238+
add_comgr_test(mangled_names_hip_test c)
239+
add_comgr_test(unbundle_hip_test c)

‎amd/comgr/test/compile_hip_test.c‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -48,13 +48,13 @@ int main(int Argc, char *Argv[]) {
4848
amd_comgr_action_info_t ActionInfo;
4949
amd_comgr_status_t Status;
5050

51-
SizeSource = setBuf(TEST_OBJ_DIR "/source2.hip", &BufSource);
51+
SizeSource = setBuf(TEST_OBJ_DIR "/source1.hip", &BufSource);
5252

5353
Status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_SOURCE, &DataSrc);
5454
checkError(Status, "amd_comgr_create_data");
5555
Status = amd_comgr_set_data(DataSrc, SizeSource, BufSource);
5656
checkError(Status, "amd_comgr_set_data");
57-
Status = amd_comgr_set_data_name(DataSrc, "source2.hip");
57+
Status = amd_comgr_set_data_name(DataSrc, "source1.hip");
5858
checkError(Status, "amd_comgr_set_data_name");
5959
Status = amd_comgr_create_data_set(&DataSetSrc);
6060
checkError(Status, "amd_comgr_create_data_set");

‎amd/comgr/test/compile_hip_to_relocatable.c‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ int main(int argc, char *argv[]) {
5050
size_t CodeGenOptionsCount =
5151
sizeof(CodeGenOptions) / sizeof(CodeGenOptions[0]);
5252

53-
SizeSource = setBuf(TEST_OBJ_DIR "/source2.hip", &BufSource);
53+
SizeSource = setBuf(TEST_OBJ_DIR "/source1.hip", &BufSource);
5454

5555
Status = amd_comgr_create_data_set(&DataSetIn);
5656
checkError(Status, "amd_comgr_create_data_set");

‎amd/comgr/test/compile_source_to_executable.c‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -110,7 +110,8 @@ int main(int argc, char *argv[]) {
110110
free(BufSource);
111111
}
112112

113-
#ifdef HIP_COMPILER
113+
// Re-enable post https://github.com/llvm/llvm-project/pull/85672
114+
#if 0
114115
// HIP
115116
{
116117
char *BufSource;

‎amd/comgr/test/mangled_names_hip_test.c‎

Lines changed: 17 additions & 114 deletions
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,7 @@ int main(int argc, char *argv[]) {
9191
&DataBc);
9292
checkError(Status, "amd_comgr_action_data_get_data");
9393

94-
#if 0
94+
#if 1
9595
// write bitcode
9696
{
9797
size_t bytes_size = 0;
@@ -122,9 +122,8 @@ int main(int argc, char *argv[]) {
122122
Status = amd_comgr_populate_mangled_names(DataBc, &numNames);
123123
checkError(Status, "amd_comgr_populate_mangled_names");
124124

125-
const char *bcNames[] = {"_Z7source1Pi"};
126-
size_t bcNumNames = 1;
127-
bool bcFound[1] = {false};
125+
char *mangledSubstr = "__hip_cuid_";
126+
bool bcFound = false;
128127

129128
for (size_t I = 0; I < numNames; ++I) {
130129
size_t Size;
@@ -135,22 +134,18 @@ int main(int argc, char *argv[]) {
135134
Status = amd_comgr_get_mangled_name(DataBc, I, &Size, mName);
136135
checkError(Status, "amd_comgr_get_mangled_name");
137136

138-
for (size_t J = 0; J < bcNumNames; ++J) {
139-
if (!strcmp(mName, bcNames[J])) {
140-
bcFound[J] = true;
141-
}
137+
if (strstr(mName, mangledSubstr)) {
138+
bcFound = true;
142139
}
143140

144141
free(mName);
145142
}
146143

147-
for (size_t I = 0; I < bcNumNames; I++) {
148-
if (!bcFound[I]) {
149-
printf("amd_get_mangled_name from bc Failed: "
150-
"(expected '%s')\n",
151-
bcNames[I]);
152-
exit(1);
153-
}
144+
if (!bcFound) {
145+
printf("amd_get_mangled_name from bc Failed: "
146+
"(expected '%s*')\n",
147+
mangledSubstr);
148+
exit(1);
154149
}
155150

156151
Status = amd_comgr_create_data_set(&DataSetLinked);
@@ -218,9 +213,7 @@ int main(int argc, char *argv[]) {
218213

219214
Status = amd_comgr_populate_mangled_names(DataExec, &numNames);
220215

221-
const char *execNames[] = {"_Z7source1Pi", "_Z7source1Pi.kd"};
222-
size_t execNumNames = 2;
223-
bool execFound[2] = {false, false};
216+
bool execFound = false;
224217

225218
for (size_t I = 0; I < numNames; ++I) {
226219
size_t Size;
@@ -231,106 +224,20 @@ int main(int argc, char *argv[]) {
231224
Status = amd_comgr_get_mangled_name(DataExec, I, &Size, mName);
232225
checkError(Status, "amd_comgr_get_mangled_name");
233226

234-
for (size_t J = 0; J < execNumNames; ++J) {
235-
if (!strcmp(mName, execNames[J])) {
236-
execFound[J] = true;
237-
}
227+
if (strstr(mName, mangledSubstr)) {
228+
execFound = true;
238229
}
239230

240231
free(mName);
241232
}
242233

243-
for (size_t I = 0; I < execNumNames; I++) {
244-
if (!execFound[I]) {
245-
printf("amd_get_mangled_name from bc Failed: "
246-
"(expected '%s')\n",
247-
execNames[I]);
248-
exit(1);
249-
}
250-
}
251-
252-
//
253-
// Test AMD_COMGR_ACTION_COMPILE_SOURCE_TO_RELOCATABLE
254-
//
255-
256-
Status = amd_comgr_create_data_set(&DataSetReloc2);
257-
checkError(Status, "amd_comgr_create_data_set");
258-
259-
Status = amd_comgr_do_action(AMD_COMGR_ACTION_COMPILE_SOURCE_TO_RELOCATABLE,
260-
DataAction, DataSetIn, DataSetReloc2);
261-
checkError(Status, "amd_comgr_do_action");
262-
263-
Status = amd_comgr_action_data_count(DataSetReloc2,
264-
AMD_COMGR_DATA_KIND_RELOCATABLE, &Count);
265-
checkError(Status, "amd_comgr_action_data_count");
266-
267-
if (Count != 1) {
268-
printf("AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE Failed: "
269-
"produced %zu source objects (expected 1)\n",
270-
Count);
234+
if (!execFound) {
235+
printf("amd_get_mangled_name from exec Failed: "
236+
"(expected '%s*')\n",
237+
mangledSubstr);
271238
exit(1);
272239
}
273240

274-
Status = amd_comgr_create_data_set(&DataSetExec2);
275-
checkError(Status, "amd_comgr_create_data_set");
276-
277-
Status = amd_comgr_action_info_set_option_list(DataAction, NULL, 0);
278-
checkError(Status, "amd_comgr_action_info_set_option_list");
279-
280-
Status = amd_comgr_do_action(AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE,
281-
DataAction, DataSetReloc2, DataSetExec2);
282-
checkError(Status, "amd_comgr_do_action");
283-
284-
Status = amd_comgr_action_data_count(DataSetExec2,
285-
AMD_COMGR_DATA_KIND_EXECUTABLE, &Count);
286-
checkError(Status, "amd_comgr_action_data_count");
287-
288-
if (Count != 1) {
289-
printf("AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE Failed: "
290-
"produced %zu executable objects (expected 1)\n",
291-
Count);
292-
exit(1);
293-
}
294-
295-
// Get Mangled Names
296-
amd_comgr_data_t DataExec2;
297-
298-
Status = amd_comgr_action_data_get_data(
299-
DataSetExec2, AMD_COMGR_DATA_KIND_EXECUTABLE, 0, &DataExec2);
300-
301-
Status = amd_comgr_populate_mangled_names(DataExec2, &numNames);
302-
303-
for (size_t I = 0; I < execNumNames; ++I) {
304-
execFound[I] = false;
305-
}
306-
307-
for (size_t I = 0; I < numNames; ++I) {
308-
size_t Size;
309-
Status = amd_comgr_get_mangled_name(DataExec, I, &Size, NULL);
310-
checkError(Status, "amd_comgr_get_mangled_name");
311-
312-
char *mName = calloc(Size, sizeof(char));
313-
Status = amd_comgr_get_mangled_name(DataExec, I, &Size, mName);
314-
checkError(Status, "amd_comgr_get_mangled_name");
315-
316-
for (size_t J = 0; J < execNumNames; ++J) {
317-
if (!strcmp(mName, execNames[J])) {
318-
execFound[J] = true;
319-
}
320-
}
321-
322-
free(mName);
323-
}
324-
325-
for (size_t I = 0; I < execNumNames; I++) {
326-
if (!execFound[I]) {
327-
printf("amd_get_mangled_name from bc Failed: "
328-
"(expected '%s')\n",
329-
execNames[I]);
330-
exit(1);
331-
}
332-
}
333-
334241
Status = amd_comgr_release_data(DataSource);
335242
checkError(Status, "amd_comgr_release_data");
336243
Status = amd_comgr_release_data(DataBc);
@@ -347,10 +254,6 @@ int main(int argc, char *argv[]) {
347254
checkError(Status, "amd_comgr_destroy_data_set");
348255
Status = amd_comgr_destroy_data_set(DataSetExec);
349256
checkError(Status, "amd_comgr_destroy_data_set");
350-
Status = amd_comgr_destroy_data_set(DataSetReloc2);
351-
checkError(Status, "amd_comgr_destroy_data_set");
352-
Status = amd_comgr_destroy_data_set(DataSetExec2);
353-
checkError(Status, "amd_comgr_destroy_data_set");
354257
Status = amd_comgr_destroy_action_info(DataAction);
355258
checkError(Status, "amd_comgr_destroy_action_info");
356259
free(BufSource);

‎amd/comgr/test/source/cube.hip‎

Lines changed: 2 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -20,17 +20,6 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
2020
THE SOFTWARE.
2121
*/
2222

23-
#include "hip/hip_runtime.h"
24-
25-
extern __device__ bool test_flag;
26-
27-
template <typename T>
28-
__global__ void vector_cube(T* C_d, const T* A_d, size_t N) {
29-
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
30-
size_t stride = hipBlockDim_x * hipGridDim_x;
31-
32-
if (test_flag)
33-
for (size_t i = offset; i < N; i += stride) {
34-
C_d[i] = A_d[i] * A_d[i] * A_d[i];
35-
}
23+
void cube(int *j) {
24+
*j = *j * *j * *j;
3625
}

‎amd/comgr/test/source/double.hip‎

Lines changed: 2 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -20,21 +20,6 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
2020
THE SOFTWARE.
2121
*/
2222

23-
#include "hip/hip_runtime.h"
24-
25-
/*
26-
* Double each element in the array A and write to array C.
27-
*/
28-
29-
__device__ bool test_flag = true;
30-
31-
template <typename T>
32-
__global__ void vector_double(T* C_d, const T* A_d, size_t N) {
33-
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
34-
size_t stride = hipBlockDim_x * hipGridDim_x;
35-
36-
if (test_flag)
37-
for (size_t i = offset; i < N; i += stride) {
38-
C_d[i] = A_d[i] * 2;
39-
}
23+
void doubles(int *j) {
24+
*j = *j * 2;
4025
}

‎amd/comgr/test/source/source1.hip‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,8 +32,7 @@
3232
* THE SOFTWARE.
3333
*
3434
*******************************************************************************/
35-
#include "hip/hip_runtime.h"
3635

37-
__global__ void source1(int *j) {
36+
void source1(int *j) {
3837
*j += 2;
3938
}

‎amd/comgr/test/source/source2.hip‎

Lines changed: 0 additions & 48 deletions
This file was deleted.

0 commit comments

Comments
 (0)