Skip to content

Commit 2703099

Browse files
tangjj11tangjj11
authored
[SYCLomatic #1345] Add 5 bindless image cases. (#574)
Signed-off-by: Tang, Jiajun jiajun.tang@intel.com
1 parent b7fe107 commit 2703099

11 files changed

Lines changed: 2562 additions & 6 deletions

‎features/config/TEMPLATE_image.xml‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,8 +5,5 @@
55
<files>
66
<file path="feature_case/image/${testName}.cu" />
77
</files>
8-
<rules>
9-
<optlevelRule excludeOptlevelNameString="cuda" />
10-
</rules>
118
</test>
129

‎features/config/TEMPLATE_image_before_12.xml‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
<?xml version="1.0" encoding="UTF-8"?>
2+
3+
<test driverID="test_feature" name="TEMPLATE">
4+
<description>test</description>
5+
<files>
6+
<file path="feature_case/image/${testName}.cu" />
7+
</files>
8+
<rules>
9+
<platformRule OSFamily="Linux" kit="CUDA12.0" kitRange="LATER_OR_EQUAL" runOnThisPlatform="false"/>
10+
<platformRule OSFamily="Windows" kit="CUDA12.0" kitRange="LATER_OR_EQUAL" runOnThisPlatform="false"/>
11+
</rules>
12+
</test>
13+

‎features/config/TEMPLATE_image_skip_cpu_gpu.xml‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
<?xml version="1.0" encoding="UTF-8"?>
2+
3+
<test driverID="test_feature" name="TEMPLATE">
4+
<description>test</description>
5+
<files>
6+
<file path="feature_case/image/${testName}.cu" />
7+
</files>
8+
<rules>
9+
<optlevelRule excludeOptlevelNameString="cpu" />
10+
<optlevelRule excludeOptlevelNameString="gpu" />
11+
</rules>
12+
</test>
13+

‎features/config/TEMPLATE_image_skip_cuda.xml‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
<?xml version="1.0" encoding="UTF-8"?>
2+
3+
<test driverID="test_feature" name="TEMPLATE">
4+
<description>test</description>
5+
<files>
6+
<file path="feature_case/image/${testName}.cu" />
7+
</files>
8+
<rules>
9+
<optlevelRule excludeOptlevelNameString="cuda" />
10+
</rules>
11+
</test>
12+

‎features/feature_case/image/text_experimental_build_only.cu‎

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,120 @@
1+
// ===-------- text_experimental_build_only.cu ------- *- CUDA -* ---------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//
8+
// ===---------------------------------------------------------------------===//
9+
10+
__global__ void
11+
CppLanguageExtensions_TextureFunctions(cudaTextureObject_t tex) {
12+
int i = 1, t = 1;
13+
float j = 1, k = 1, l = 1;
14+
tex1Dfetch<float4>(tex, i);
15+
tex1D<short2>(tex, i);
16+
tex2D<int1>(tex, j, k);
17+
// tex3D<char4>(tex, j, k, l); // TODO: need support.
18+
// tex1DLayered<uchar2>(tex, i, t); // TODO: need support.
19+
// tex2DLayered<uint2>(tex, j, k, t); // TODO: need support.
20+
}
21+
22+
void Runtime_MemoryManagement() {
23+
cudaChannelFormatDesc d;
24+
cudaExtent e;
25+
unsigned int u;
26+
cudaArray_t a = nullptr;
27+
cudaPitchedPtr p;
28+
size_t s = 1;
29+
void *v;
30+
cudaMemcpyKind k = cudaMemcpyDefault;
31+
cudaMemcpy3DParms pm;
32+
int i = 1;
33+
cudaArrayGetInfo(&d, &e, &u, a);
34+
// cudaFreeArray(a); // TODO: need support.
35+
cudaMalloc3D(&p, e);
36+
cudaMalloc3DArray(&a, &d, e, u);
37+
cudaMallocArray(&a, &d, s, s, u);
38+
cudaMallocPitch(&v, &s, s, s);
39+
cudaMemcpy2D(v, s, v, s, s, s, k);
40+
cudaMemcpy2DArrayToArray(a, s, s, a, s, s, s, s, k);
41+
cudaMemcpy2DAsync(v, s, v, s, s, s, k);
42+
cudaMemcpy2DFromArray(v, s, a, s, s, s, s, k);
43+
cudaMemcpy2DFromArrayAsync(v, s, a, s, s, s, s, k);
44+
cudaMemcpy2DToArray(a, s, s, v, s, s, s, k);
45+
cudaMemcpy2DToArrayAsync(a, s, s, v, s, s, s, k);
46+
cudaMemcpy3D(&pm);
47+
cudaMemcpy3DAsync(&pm);
48+
cudaMemset2D(v, s, i, s, s);
49+
cudaMemset2DAsync(v, s, i, s, s);
50+
cudaMemset3D(p, i, e);
51+
cudaMemset3DAsync(p, i, e);
52+
}
53+
54+
void Runtime_TextureObjectManagement() {
55+
int i = 1;
56+
cudaChannelFormatKind k = cudaChannelFormatKindSigned;
57+
cudaTextureObject_t o{0}; // TODO: need not "{0}".
58+
cudaResourceDesc r;
59+
cudaTextureDesc t;
60+
// cudaResourceViewDesc v; // TODO: need support.
61+
cudaArray_t a = nullptr;
62+
cudaChannelFormatDesc d;
63+
cudaCreateChannelDesc(i, i, i, i, k);
64+
cudaCreateTextureObject(&o, &r, &t, nullptr /*&v*/);
65+
cudaDestroyTextureObject(o);
66+
cudaGetChannelDesc(&d, a);
67+
cudaGetTextureObjectResourceDesc(&r, o);
68+
cudaGetTextureObjectTextureDesc(&t, o);
69+
}
70+
71+
void Driver_MemoryManagement() {
72+
CUarray a;
73+
CUDA_ARRAY_DESCRIPTOR D;
74+
CUDA_MEMCPY2D C2;
75+
CUstream cs;
76+
CUDA_MEMCPY3D C3;
77+
size_t s;
78+
CUdeviceptr d;
79+
void *v;
80+
cuArrayCreate(&a, &D);
81+
cuArrayDestroy(a);
82+
cuMemcpy2D(&C2);
83+
cuMemcpy2DAsync(&C2, cs);
84+
cuMemcpy3D(&C3);
85+
cuMemcpy3DAsync(&C3, cs);
86+
cuMemcpyAtoA(a, s, a, s, s);
87+
cuMemcpyAtoD(d, a, s, s);
88+
cuMemcpyAtoH(&v, a, s, s);
89+
cuMemcpyAtoHAsync(&v, a, s, s, cs);
90+
cuMemcpyDtoA(a, s, d, s);
91+
cuMemcpyDtoD(d, d, s);
92+
cuMemcpyDtoDAsync(d, d, s, cs);
93+
cuMemcpyDtoH(v, d, s);
94+
cuMemcpyDtoHAsync(v, d, s, cs);
95+
cuMemcpyHtoA(a, s, v, s);
96+
cuMemcpyHtoAAsync(a, s, v, s, cs);
97+
cuMemcpyHtoD(d, v, s);
98+
cuMemcpyHtoDAsync(d, v, s, cs);
99+
}
100+
101+
void Driver_TextureObjectManagement() {
102+
CUtexObject o{0}; // TODO: need not "{0}".
103+
CUDA_RESOURCE_DESC R;
104+
CUDA_TEXTURE_DESC T;
105+
// CUDA_RESOURCE_VIEW_DESC V; // TODO: need support.
106+
// cuTexObjectCreate(&o, &R, &T, nullptr /*&V*/); // TODO: need support.
107+
cuTexObjectDestroy(o);
108+
// cuTexObjectGetResourceDesc(&R, o); // TODO: need support.
109+
// cuTexObjectGetTextureDesc(&T, o); // TODO: need support.
110+
}
111+
112+
int main() {
113+
Runtime_MemoryManagement();
114+
Runtime_TextureObjectManagement();
115+
Driver_MemoryManagement();
116+
Driver_TextureObjectManagement();
117+
cudaTextureObject_t tex{0}; // TODO: need not "{0}".
118+
CppLanguageExtensions_TextureFunctions<<<1, 1>>>(tex);
119+
return 0;
120+
}

‎features/feature_case/image/text_experimental_build_only_before12.cu‎

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
// ===----- text_experimental_build_only_before12.cu ----- *- CUDA -* -----===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//
8+
// ===---------------------------------------------------------------------===//
9+
10+
static texture<float4, 2> r;
11+
12+
void Runtime_MemoryManagement() {
13+
cudaArray_t a = nullptr;
14+
size_t s = 1;
15+
cudaMemcpyKind k = cudaMemcpyDefault;
16+
void *v = nullptr;
17+
cudaMemcpyArrayToArray(a, s, s, a, s, s, s, k);
18+
cudaMemcpyFromArray(v, a, s, s, s, k);
19+
cudaMemcpyFromArrayAsync(v, a, s, s, s, k);
20+
cudaMemcpyToArray(a, s, s, v, s, k);
21+
cudaMemcpyToArrayAsync(a, s, s, v, s, k);
22+
}
23+
24+
void Runtime_TextureReferenceManagement() {
25+
size_t s = 1;
26+
void *v;
27+
cudaChannelFormatDesc d;
28+
cudaArray_t a = nullptr;
29+
// cudaBindTexture(&s, &r, v, &d); // TODO: need support.
30+
// cudaBindTexture2D(&s, &r, v, &d, s, s, s); // TODO: need support.
31+
cudaBindTextureToArray(&r, a, &d);
32+
cudaUnbindTexture(&r);
33+
}
34+
35+
void Driver_TextureReferenceManagement() {
36+
CUaddress_mode am;
37+
CUtexref r;
38+
int i = 1;
39+
CUfilter_mode fm;
40+
unsigned int u;
41+
size_t s = 1;
42+
CUdeviceptr d;
43+
CUDA_ARRAY_DESCRIPTOR D;
44+
CUarray a;
45+
CUarray_format f;
46+
cuTexRefGetAddressMode(&am, r, i);
47+
cuTexRefGetFilterMode(&fm, r);
48+
cuTexRefGetFlags(&u, r);
49+
cuTexRefSetAddress(&s, r, d, s);
50+
cuTexRefSetAddress2D(r, &D, d, s);
51+
cuTexRefSetAddressMode(r, i, am);
52+
cuTexRefSetArray(r, a, u);
53+
cuTexRefSetFilterMode(r, fm);
54+
cuTexRefSetFlags(r, u);
55+
cuTexRefSetFormat(r, f, i);
56+
}
57+
58+
int main() {
59+
Runtime_MemoryManagement();
60+
Runtime_TextureReferenceManagement();
61+
Driver_TextureReferenceManagement();
62+
return 0;
63+
}

0 commit comments

Comments
 (0)