@@ -14,9 +14,9 @@ __kernel void cai_dot_product
14
14
const int FNumBs ,
15
15
const int FSize ,
16
16
int ActFN ,
17
- __global __read_only float * FInputBufferAs ,
18
- __global __read_only float * FInputBufferBs ,
19
- __global float * FResultBuffer
17
+ __global float * FInputBufferAs ,
18
+ __global float * FInputBufferBs ,
19
+ __global float * FResultBuffer
20
20
)
21
21
{
22
22
const int a_id = get_global_id (0 );
@@ -118,9 +118,9 @@ __kernel void cai_dot_product2
118
118
const int FNumBs ,
119
119
const int FSize ,
120
120
int ActFN ,
121
- __global __read_only float * FInputBufferAs ,
122
- __global __read_only float * FInputBufferBs ,
123
- __global float * FResultBuffer
121
+ __global float * FInputBufferAs ,
122
+ __global float * FInputBufferBs ,
123
+ __global float * FResultBuffer
124
124
)
125
125
{
126
126
const int a_id = get_global_id (0 );
@@ -252,8 +252,8 @@ __kernel void simpleGEMMT(
252
252
const int FThreadCount ,
253
253
const int M , const int N , const int K ,
254
254
int ActFN ,
255
- __global __read_only float * A ,
256
- __global __read_only float * B ,
255
+ __global float * A ,
256
+ __global float * B ,
257
257
__global float * C ) {
258
258
259
259
// Thread identifiers
@@ -304,8 +304,8 @@ __kernel void cai_dot_product_simple
304
304
const int FNumBs ,
305
305
const int FSize ,
306
306
int ActFN ,
307
- __global __read_only float16 * FInputBufferAs ,
308
- __global __read_only float16 * FInputBufferBs ,
307
+ __global float16 * FInputBufferAs ,
308
+ __global float16 * FInputBufferBs ,
309
309
__global float * FResultBuffer
310
310
)
311
311
{
@@ -355,8 +355,8 @@ __kernel void myGEMM5(
355
355
const int FThreadCount ,
356
356
const int M , const int N , const int K ,
357
357
int ActFN ,
358
- __global __read_only float * A ,
359
- __global __read_only float * B ,
358
+ __global float * A ,
359
+ __global float * B ,
360
360
__global float * C ) {
361
361
362
362
// Thread identifiers
@@ -423,8 +423,8 @@ __kernel void myGEMM6(
423
423
const int FThreadCount ,
424
424
const int M , const int N , const int K ,
425
425
int ActFN ,
426
- __global __read_only float * A ,
427
- __global __read_only float * B ,
426
+ __global float * A ,
427
+ __global float * B ,
428
428
__global float * C ) {
429
429
// Thread identifiers
430
430
const int tidm = get_local_id (0 ); // Local row ID (max: TSM/WPTM)
0 commit comments