Skip to content

Commit c108e59

Browse files
rocm-cirocm-ci
authored
SWDEV-347688 - port compiler folder Dtest (#2858)
Change-Id: I79d505881b23d2f12429609cc8cbf478d19bf944
1 parent ca2ed9b commit c108e59

File tree

4 files changed

+463
-0
lines changed

4 files changed

+463
-0
lines changed

‎tests/catch/unit/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@ add_subdirectory(texture)
3131
add_subdirectory(streamperthread)
3232
add_subdirectory(kernel)
3333
add_subdirectory(multiThread)
34+
add_subdirectory(compiler)
3435
if(HIP_PLATFORM STREQUAL "amd")
3536
add_subdirectory(clock)
3637
endif()

‎tests/catch/unit/compiler/CMakeLists.txt‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
# Common Tests - Test independent of all platforms
2+
set(TEST_SRC
3+
hipClassKernel.cc
4+
)
5+
6+
hip_add_exe_to_target(NAME CompilerTest
7+
TEST_SRC ${TEST_SRC}
8+
TEST_TARGET_NAME build_tests)

‎tests/catch/unit/compiler/hipClassKernel.cc‎

Lines changed: 220 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,220 @@
1+
/*
2+
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
3+
4+
Permission is hereby granted, free of charge, to any person obtaining a copy
5+
of this software and associated documentation files (the "Software"), to deal
6+
in the Software without restriction, including without limitation the rights
7+
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8+
copies of the Software, and to permit persons to whom the Software is
9+
furnished to do so, subject to the following conditions:
10+
11+
The above copyright notice and this permission notice shall be included in
12+
all copies or substantial portions of the Software.
13+
14+
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15+
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16+
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17+
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18+
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19+
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20+
THE SOFTWARE.
21+
*/
22+
/* HIT_START
23+
* BUILD: %t %s ../test_common.cpp
24+
* TEST: %t
25+
* HIT_END
26+
*/
27+
#include "hipClassKernel.h"
28+
29+
__global__ void
30+
ovrdClassKernel(bool* result_ecd){
31+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
32+
testOvrD tobj1;
33+
result_ecd[tid] = (tobj1.ovrdFunc1() == 30);
34+
}
35+
36+
__global__ void
37+
ovldClassKernel(bool* result_ecd){
38+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
39+
testFuncOvld tfo1;
40+
result_ecd[tid] = (tfo1.func1(10) == 20)
41+
&& (tfo1.func1(10,10) == 30);
42+
}
43+
44+
TEST_CASE("Unit_hipClassKernel_Overload_Override") {
45+
bool *result_ecd, *result_ech;
46+
result_ech = AllocateHostMemory();
47+
result_ecd = AllocateDeviceMemory();
48+
49+
hipLaunchKernelGGL(ovrdClassKernel,
50+
dim3(BLOCKS),
51+
dim3(THREADS_PER_BLOCK),
52+
0,
53+
0,
54+
result_ecd);
55+
56+
VerifyResult(result_ech,result_ecd);
57+
FreeMem(result_ech,result_ecd);
58+
59+
result_ech = AllocateHostMemory();
60+
result_ecd = AllocateDeviceMemory();
61+
62+
hipLaunchKernelGGL(ovldClassKernel,
63+
dim3(BLOCKS),
64+
dim3(THREADS_PER_BLOCK),
65+
0,
66+
0,
67+
result_ecd);
68+
69+
VerifyResult(result_ech,result_ecd);
70+
FreeMem(result_ech,result_ecd);
71+
}
72+
73+
// check for friend
74+
__global__ void
75+
friendClassKernel(bool* result_ecd){
76+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
77+
testFrndB tfb1;
78+
result_ecd[tid] = (tfb1.showA() == 10);
79+
}
80+
81+
TEST_CASE("Unit_hipClassKernel_Friend") {
82+
bool *result_ecd;
83+
result_ecd = AllocateDeviceMemory();
84+
hipLaunchKernelGGL(friendClassKernel,
85+
dim3(BLOCKS),
86+
dim3(THREADS_PER_BLOCK),
87+
0,
88+
0,
89+
result_ecd);
90+
}
91+
92+
// check sizeof empty class is 1
93+
__global__ void
94+
emptyClassKernel(bool* result_ecd) {
95+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
96+
testClassEmpty ob1,ob2;
97+
result_ecd[tid] = (sizeof(testClassEmpty) == 1)
98+
&& (&ob1 != &ob2);
99+
}
100+
101+
TEST_CASE("Unit_hipClassKernel_Empty") {
102+
bool *result_ecd, *result_ech;
103+
result_ech = AllocateHostMemory();
104+
result_ecd = AllocateDeviceMemory();
105+
106+
hipLaunchKernelGGL(emptyClassKernel,
107+
dim3(BLOCKS),
108+
dim3(THREADS_PER_BLOCK),
109+
0,
110+
0,
111+
result_ecd);
112+
113+
VerifyResult(result_ech,result_ecd);
114+
FreeMem(result_ech,result_ecd);
115+
}
116+
117+
// tests for classes >8 bytes
118+
__global__ void
119+
sizeClassBKernel(bool* result_ecd) {
120+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
121+
result_ecd[tid] = (sizeof(testSizeB) == 12)
122+
&& (sizeof(testSizeC) == 16)
123+
&& (sizeof(testSizeP1) == 6)
124+
&& (sizeof(testSizeP2) == 13)
125+
&& (sizeof(testSizeP3) == 8);
126+
}
127+
128+
TEST_CASE("Unit_hipClassKernel_BSize") {
129+
bool *result_ecd, *result_ech;
130+
result_ech = AllocateHostMemory();
131+
result_ecd = AllocateDeviceMemory();
132+
133+
hipLaunchKernelGGL(sizeClassBKernel,
134+
dim3(BLOCKS),
135+
dim3(THREADS_PER_BLOCK),
136+
0,
137+
0,
138+
result_ecd);
139+
140+
VerifyResult(result_ech,result_ecd);
141+
FreeMem(result_ech,result_ecd);
142+
}
143+
144+
__global__ void
145+
sizeClassKernel(bool* result_ecd) {
146+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
147+
result_ecd[tid] = (sizeof(testSizeA) == 16)
148+
&& (sizeof(testSizeDerived) == 24)
149+
&& (sizeof(testSizeDerived2) == 20);
150+
}
151+
152+
TEST_CASE("Unit_hipClassKernel_Size") {
153+
bool *result_ecd, *result_ech;
154+
result_ech = AllocateHostMemory();
155+
result_ecd = AllocateDeviceMemory();
156+
157+
hipLaunchKernelGGL(sizeClassKernel,
158+
dim3(BLOCKS),
159+
dim3(THREADS_PER_BLOCK),
160+
0,
161+
0,
162+
result_ecd);
163+
164+
VerifyResult(result_ech,result_ecd);
165+
FreeMem(result_ech,result_ecd);
166+
}
167+
168+
__global__ void
169+
sizeVirtualClassKernel(bool* result_ecd) {
170+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
171+
result_ecd[tid] = (sizeof(testSizeDV) == 16)
172+
&& (sizeof(testSizeDerivedDV) == 16)
173+
&& (sizeof(testSizeVirtDerPack) == 24)
174+
&& (sizeof(testSizeVirtDer) == 24)
175+
&& (sizeof(testSizeDerMulti) == 48) ;
176+
}
177+
178+
TEST_CASE("Unit_hipClassKernel_Virtual") {
179+
bool *result_ecd, *result_ech;
180+
result_ech = AllocateHostMemory();
181+
result_ecd = AllocateDeviceMemory();
182+
183+
hipLaunchKernelGGL(sizeVirtualClassKernel,
184+
dim3(BLOCKS),
185+
dim3(THREADS_PER_BLOCK),
186+
0,
187+
0,
188+
result_ecd);
189+
190+
VerifyResult(result_ech,result_ecd);
191+
FreeMem(result_ech,result_ecd);
192+
}
193+
194+
// check pass by value
195+
__global__ void
196+
passByValueKernel(testPassByValue obj, bool* result_ecd) {
197+
int tid = threadIdx.x + blockIdx.x * blockDim.x;
198+
result_ecd[tid] = (obj.exI == 10)
199+
&& (obj.exC == 'C');
200+
}
201+
202+
TEST_CASE("Unit_hipClassKernel_Value") {
203+
bool *result_ecd,*result_ech;
204+
result_ech = AllocateHostMemory();
205+
result_ecd = AllocateDeviceMemory();
206+
207+
testPassByValue exObj;
208+
exObj.exI = 10;
209+
exObj.exC = 'C';
210+
hipLaunchKernelGGL(passByValueKernel,
211+
dim3(BLOCKS),
212+
dim3(THREADS_PER_BLOCK),
213+
0,
214+
0,
215+
exObj,
216+
result_ecd);
217+
218+
VerifyResult(result_ech,result_ecd);
219+
FreeMem(result_ech,result_ecd);
220+
}

0 commit comments

Comments
 (0)