Skip to content
Projects
Groups
Snippets
Help
Loading...
Help
Support
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
O
OpenXG-RAN
Project overview
Project overview
Details
Activity
Releases
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Issues
0
Issues
0
List
Boards
Labels
Milestones
Merge Requests
1
Merge Requests
1
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Operations
Operations
Metrics
Environments
Analytics
Analytics
CI / CD
Repository
Value Stream
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
OpenXG
OpenXG-RAN
Commits
d31ac580
Commit
d31ac580
authored
Sep 27, 2021
by
frtabu
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
add openCL board detection
parent
29c7a3aa
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
86 additions
and
19 deletions
+86
-19
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
+86
-19
No files found.
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
View file @
d31ac580
/*! \file PHY/CODING/nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
/*
* \brief LDPC cuda support BG1 all length
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* \author NCTU OpinConnect Terng-Yin Hsu,WEI-YING,LIN
* contributor license agreements. See the NOTICE file distributed with
* \email tyhsu@cs.nctu.edu.tw
* this work for additional information regarding copyright ownership.
* \date 13-05-2020
* The OpenAirInterface Software Alliance licenses this file to You under
* \version
* the OAI Public License, Version 1.0 (the "License"); you may not use this file
* \note
* except in compliance with the License.
* \warning
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
*/
/*! \file nrLDPC_decoder_CL.c
* \brief ldpc decoder, openCL implementaion
* \author
* \date 2021
* \version 1.0
* @ingroup
*/
/* uses HW component id for log messages ( --log_config.hw_log_level <warning| info|debug|trace>) */
#include <stdio.h>
#include <stdio.h>
#include <unistd.h>
#include <unistd.h>
#include <cuda_runtime.h>
#include <cuda_runtime.h>
...
@@ -20,7 +39,7 @@
...
@@ -20,7 +39,7 @@
#define MAX_ITERATION 2
#define MAX_ITERATION 2
#define MC 1
#define MC 1
#define MAX_OCLDEV 10
typedef
struct
{
typedef
struct
{
char
x
;
char
x
;
...
@@ -28,13 +47,23 @@ typedef struct{
...
@@ -28,13 +47,23 @@ typedef struct{
short
value
;
short
value
;
}
h_element
;
}
h_element
;
typedef
struct
{
cl_uint
max_CU
;
cl_uint
max_WID
;
size_t
*
max_WIS
;
cl_context
context
;
cl_program
program
;
cl_kernel
kernel
;
cl_command_queue
queue
;
}
ocldev_t
;
typedef
struct
{
ocldev_t
ocldev
[
MAX_OCLDEV
];
}
ocl_t
;
ocl_t
ocl
;
void
init_LLR_DMA
(
t_nrLDPC_dec_params
*
p_decParams
,
int8_t
*
p_llr
,
int8_t
*
p_out
){
void
init_LLR_DMA
(
t_nrLDPC_dec_params
*
p_decParams
,
int8_t
*
p_llr
,
int8_t
*
p_out
){
...
@@ -57,26 +86,64 @@ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out
...
@@ -57,26 +86,64 @@ void init_LLR_DMA(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out
// cudaDeviceSynchronize();
// cudaDeviceSynchronize();
}
}
cl_error_callback
(
const
char
*
errinfo
,
const
void
*
private_info
,
size_t
cb
,
void
*
user_data
)
{
LOG_E
(
HW
,
"OpenCL accelerator error %s
\n
"
,
errinfo
);
}
char
*
clutil_getstrdev
(
int
intdev
)
{
static
char
retstring
[
255
]
=
""
;
char
*
retptr
=
retstring
;
retptr
+=
sprintf
(
retptr
,
"0x%08x: "
,(
uint32_t
)
intdev
);
if
(
intdev
&
CL_DEVICE_TYPE_CPU
)
retptr
+=
sprintf
(
retptr
,
"%s"
,
"cpu "
);
if
(
intdev
&
CL_DEVICE_TYPE_GPU
)
retptr
+=
sprintf
(
retptr
,
"%s"
,
"gpu "
);
if
(
intdev
&
CL_DEVICE_TYPE_ACCELERATOR
)
retptr
+=
sprintf
(
retptr
,
"%s"
,
"acc "
);
return
retstring
;
}
/* from here: entry points in decoder shared lib */
/* from here: entry points in decoder shared lib */
int
ldpc_autoinit
(
void
)
{
// called by the library loader
int
ldpc_autoinit
(
void
)
{
// called by the library loader
cl_platform_id
platforms
[
10
];
cl_platform_id
platforms
[
10
];
cl_int
num_platforms_found
;
cl_
u
int
num_platforms_found
;
cl_int
rt
=
clGetPlatformIDs
(
sizeof
(
platforms
)
/
sizeof
(
cl_platform_id
),
platforms
,
&
num_platforms_found
);
cl_
u
int
rt
=
clGetPlatformIDs
(
sizeof
(
platforms
)
/
sizeof
(
cl_platform_id
),
platforms
,
&
num_platforms_found
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetPlatformIDs error %d
\n
"
,
(
int
)
rt
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetPlatformIDs error %d
\n
"
,
(
int
)
rt
);
AssertFatal
(
num_platforms_found
>
0
,
"clGetPlatformIDs: no cl compatible platform found
\n
"
);
AssertFatal
(
num_platforms_found
>
0
,
"clGetPlatformIDs: no cl compatible platform found
\n
"
);
for
(
int
i
=
0
;
i
<
(
int
)
num_platforms_found
;
i
++
)
{
for
(
int
i
=
0
;
i
<
(
int
)
num_platforms_found
;
i
++
)
{
char
stringval
[
255
];
rt
=
clGetPlatformInfo
(
platforms
[
i
],
CL_PLATFORM_PROFILE
,
sizeof
(
stringval
),
stringval
,
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetPlatformInfo PROFILE error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Platform %i, OpenCL profile %s
\n
"
,
i
,
stringval
);
rt
=
clGetPlatformInfo
(
platforms
[
i
],
CL_PLATFORM_VERSION
,
sizeof
(
stringval
),
stringval
,
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetPlatformInfo VERSION error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Platform %i, OpenCL version %s
\n
"
,
i
,
stringval
);
cl_device_id
devices
[
20
];
cl_device_id
devices
[
20
];
cl_int
num_devices_found
;
cl_
u
int
num_devices_found
;
rt
=
clGetDeviceIDs
(
platforms
[
i
],
CL_DEVICE_TYPE_ALL
,
sizeof
(
devices
)
/
sizeof
(
cl_device_id
),
devices
,
&
num_devices_found
);
rt
=
clGetDeviceIDs
(
platforms
[
i
],
CL_DEVICE_TYPE_ALL
,
sizeof
(
devices
)
/
sizeof
(
cl_device_id
),
devices
,
&
num_devices_found
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceIDs error %d
\n
"
,
(
int
)
rt
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceIDs error %d
\n
"
,
(
int
)
rt
);
for
(
int
j
=
0
;
j
<
num_devices_found
;
j
++
)
{
for
(
int
j
=
0
;
j
<
num_devices_found
;
j
++
)
{
cl_device_type
devtype
;
cl_bool
abool
;
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_AVAILABLE
,
sizeof
(
abool
),
&
abool
,
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo DEVICE_AVAILABLE error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Device %i is %s available
\n
"
,
j
,
(
abool
==
CL_TRUE
?
""
:
"not"
));
cl_device_type
devtype
;
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_TYPE
,
sizeof
(
cl_device_type
),
&
devtype
,
NULL
);
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_TYPE
,
sizeof
(
cl_device_type
),
&
devtype
,
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo error %d
\n
"
,
(
int
)
rt
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo DEVICE_TYPE error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Device %i, type %d
\n
"
,
j
,(
int
)
devtype
);
LOG_I
(
HW
,
"Device %i, type %d = %s
\n
"
,
j
,(
int
)
devtype
,
clutil_getstrdev
(
devtype
));
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_MAX_COMPUTE_UNITS
,
sizeof
(
ocl
.
ocldev
[
j
].
max_CU
),
&
(
ocl
.
ocldev
[
j
].
max_CU
),
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo MAX_COMPUTE_UNITS error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Device %i, number of Compute Units: %d
\n
"
,
j
,
ocl
.
ocldev
[
j
].
max_CU
);
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
,
sizeof
(
ocl
.
ocldev
[
j
].
max_WID
),
&
(
ocl
.
ocldev
[
j
].
max_WID
),
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d
\n
"
,
(
int
)
rt
);
LOG_I
(
HW
,
"Device %i, max Work Items dimension: %d
\n
"
,
j
,
ocl
.
ocldev
[
j
].
max_WID
);
ocl
.
ocldev
[
j
].
max_WIS
=
(
size_t
*
)
malloc
(
ocl
.
ocldev
[
j
].
max_WID
*
sizeof
(
size_t
));
rt
=
clGetDeviceInfo
(
devices
[
j
],
CL_DEVICE_MAX_WORK_ITEM_SIZES
,
sizeof
(
ocl
.
ocldev
[
j
].
max_WID
)
*
sizeof
(
size_t
),
ocl
.
ocldev
[
j
].
max_WIS
,
NULL
);
AssertFatal
(
rt
==
CL_SUCCESS
,
"clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d
\n
"
,
(
int
)
rt
);
for
(
int
k
=
0
;
k
<
ocl
.
ocldev
[
j
].
max_WID
;
k
++
)
LOG_I
(
HW
,
"Device %i, max Work Items size for dimension: %d %u
\n
"
,
j
,
k
,(
uint32_t
)
ocl
.
ocldev
[
j
].
max_WIS
[
k
]);
}
}
}
}
return
0
;
return
0
;
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment