#include #include // config #define BLOCK_X 32 #define BLOCK_Y 1 #define PER_THREAD 8 #ifdef NDEBUG #define assert_assume(x) __builtin_assume(x) #else #define assert_assume(x) __builtin_assume(x) #endif __device__ float relu(float a) { return a > 0.f ? a : 0.f; } __global__ void kernel(const float* __restrict__ input, float* __restrict__ output, size_t n, size_t m) { assert_assume(reinterpret_cast(input) % sizeof(float4) == 0); assert_assume(reinterpret_cast(output) % sizeof(float4) == 0); size_t base_i = threadIdx.x * PER_THREAD + blockIdx.x * blockDim.x * PER_THREAD; size_t j = threadIdx.y + blockIdx.y * blockDim.y; assert_assume(j < m); assert_assume(base_i + (PER_THREAD-1) < n); if (PER_THREAD % 4 == 0) { float4 buffer[PER_THREAD/4]; const float4* input4 = reinterpret_cast(&input[j*n]); float4* output4 = reinterpret_cast(&output[j*n]); for (size_t item = 0; item < PER_THREAD; item += 4) { size_t i = base_i/4 + item/4; float4 lu = input4[i]; lu.x = relu(lu.x); lu.y = relu(lu.y); lu.z = relu(lu.z); lu.w = relu(lu.w); buffer[item/4] = lu; } for (size_t item = 0; item < PER_THREAD; item += 4) { size_t i = base_i/4 + item/4; output4[i] = buffer[item/4]; } return; } for (size_t item = 0; item < PER_THREAD; item++) { size_t i = base_i + item; float lu = input[j*n + i]; output[j*n + i] = relu(lu); } } // Note: input and output are all device pointers to float arrays extern "C" void solution(float* input, float* output, size_t n, size_t m) { assert(reinterpret_cast(input) % sizeof(float4) == 0); assert(reinterpret_cast(output) % sizeof(float4) == 0); dim3 blockSize(BLOCK_X, BLOCK_Y); assert(n % (blockSize.x*PER_THREAD) == 0); assert(m % blockSize.y == 0); dim3 dataDim(n, m); kernel<<>>(input, output, n, m); }

Test Case	Runtime	GFLOPS
4096x4096	50.89 μs	329.68
6144x4096	72.79 μs	345.73
7168x4096	83.85 μs	350.17
8192x4096	94.81 μs	353.93
9216x4096	105.54 μs	357.69