Commit 5f37917f authored by Shucai Xiao's avatar Shucai Xiao
Browse files

clang format

parent f50bcff2
...@@ -11,23 +11,24 @@ namespace device { ...@@ -11,23 +11,24 @@ namespace device {
__global__ void add_kernel(__half* a, __half* b, __half* r, int n) __global__ void add_kernel(__half* a, __half* b, __half* r, int n)
{ {
int tid = blockIdx.x * blockDim.x + threadIdx.x; int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) if(tid < n)
{ {
r[tid] = a[tid] + b[tid%768]; r[tid] = a[tid] + b[tid % 768];
} }
} }
void add(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2) void add(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
{ {
auto s2 = arg2.get_shape(); auto s2 = arg2.get_shape();
if (s2.element_space() == 768 and s2.type() == shape::half_type) if(s2.element_space() == 768 and s2.type() == shape::half_type)
{ {
auto elem_num = s2.elements(); auto elem_num = s2.elements();
int block_size = 1024; int block_size = 1024;
int block_num = (elem_num + block_size - 1) / block_size; int block_num = (elem_num + block_size - 1) / block_size;
add_kernel<<<block_num, block_size>>>(reinterpret_cast<__half*>(arg1.data()), add_kernel<<<block_num, block_size>>>(reinterpret_cast<__half*>(arg1.data()),
reinterpret_cast<__half*>(arg2.data()), reinterpret_cast<__half*>(arg2.data()),
reinterpret_cast<__half*>(result.data()), elem_num); reinterpret_cast<__half*>(result.data()),
elem_num);
} }
else else
{ {
......
...@@ -11,24 +11,24 @@ namespace device { ...@@ -11,24 +11,24 @@ namespace device {
__global__ void mul_kernel(__half* a, __half* b, __half* r, int n) __global__ void mul_kernel(__half* a, __half* b, __half* r, int n)
{ {
int tid = blockIdx.x * blockDim.x + threadIdx.x; int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) if(tid < n)
{ {
r[tid] = a[tid] * b[tid%768]; r[tid] = a[tid] * b[tid % 768];
} }
} }
void mul(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2) void mul(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
{ {
auto s2 = arg2.get_shape(); auto s2 = arg2.get_shape();
if (s2.element_space() == 768 and s2.type() == shape::half_type) if(s2.element_space() == 768 and s2.type() == shape::half_type)
{ {
auto elem_num = s2.elements(); auto elem_num = s2.elements();
int block_size = 1024; int block_size = 1024;
int block_num = (elem_num + block_size - 1) / block_size; int block_num = (elem_num + block_size - 1) / block_size;
mul_kernel<<<block_num, block_size>>>(reinterpret_cast<__half*>(arg1.data()), mul_kernel<<<block_num, block_size>>>(reinterpret_cast<__half*>(arg1.data()),
reinterpret_cast<__half*>(arg2.data()), reinterpret_cast<__half*>(arg2.data()),
reinterpret_cast<__half*>(result.data()), elem_num); reinterpret_cast<__half*>(result.data()),
elem_num);
} }
else else
{ {
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment