Skip to content

Commit 0e653b5

Browse files
davidAlgisdavidAlgis
committed
♻️ Refactor sample cuda to add a cuda header on kernel caller to make it cleaner, I add .cuh in premake
1 parent 680a05f commit 0e653b5

9 files changed

+55
-51
lines changed

‎Plugin/SampleBasic/include/action_sample_texture.h

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,9 @@
11
#pragma once
2+
#include "sample_kernels.cuh"
23
#include "action.h"
4+
#include "texture.h"
35
#include "cuda_include.h"
46

5-
class Texture;
6-
77
namespace SampleBasic
88
{
99
class ActionSampleTexture : public Action

‎Plugin/SampleBasic/include/action_sample_texture_array.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,9 @@
11
#pragma once
2+
#include "sample_kernels.cuh"
23
#include "action.h"
4+
#include "texture.h"
35
#include "cuda_include.h"
46

5-
class Texture;
67

78
namespace SampleBasic {
89
class ActionSampleTextureArray : public Action {

‎Plugin/SampleBasic/include/action_sample_vertex_buffer.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,8 @@
11
#pragma once
2+
#include "sample_kernels.cuh"
23
#include "action.h"
4+
#include "vertex_buffer.h"
35

4-
class VertexBuffer;
56

67
namespace SampleBasic {
78
class ActionSampleVertexBuffer: public Action {

‎Plugin/SampleBasic/include/sample_kernels.cuh

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
#pragma once
2+
#include "cuda_include.h"
3+
#include "log.h"
4+
#include "math_constants.h"
5+
#include <device_launch_parameters.h>
6+
#include <vector_functions.h>
7+
8+
9+
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock,
10+
cudaSurfaceObject_t inputSurfaceObj,
11+
const float t, const int width, const int height);
12+
void kernelCallerWriteTextureArray(const dim3 dimGrid, const dim3 dimBlock,
13+
cudaSurfaceObject_t inputSurfaceObj,
14+
const float time, const int width,
15+
const int height, const int depth);
16+
17+
void kernelCallerWriteBuffer(const dim3 dimGrid, const dim3 dimBlock,
18+
float4 *vertexPtr, const int size,
19+
const float time);

‎Plugin/SampleBasic/src/action_sample_texture.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,7 @@
11
#pragma once
22
#include "action_sample_texture.h"
3-
#include "texture.h"
43
#include "unity_plugin.h"
54

6-
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock,
7-
cudaSurfaceObject_t inputSurfaceObj,
8-
const float t, const int width, const int height);
9-
105
namespace SampleBasic
116
{
127

‎Plugin/SampleBasic/src/action_sample_texture_array.cpp

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,16 +1,7 @@
11
#pragma once
22
#include "action_sample_texture_array.h"
3-
#include "texture.h"
43
#include "unity_plugin.h"
54

6-
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock,
7-
cudaSurfaceObject_t inputSurfaceObj,
8-
const float t, const int width, const int height);
9-
void kernelCallerWriteTextureArray(const dim3 dimGrid, const dim3 dimBlock,
10-
cudaSurfaceObject_t inputSurfaceObj,
11-
const float time, const int width,
12-
const int height, const int depth);
13-
145
namespace SampleBasic
156
{
167

‎Plugin/SampleBasic/src/action_sample_vertex_buffer.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,7 @@
11
#pragma once
22
#include "action_sample_vertex_buffer.h"
33
#include "unity_plugin.h"
4-
#include "vertex_buffer.h"
54

6-
void kernelCallerWriteBuffer(const dim3 dimGrid, const dim3 dimBlock, float4* vertexPtr, const int size, const float time);
75

86
namespace SampleBasic {
97

‎Plugin/SampleBasic/src/main.cu renamed to ‎Plugin/SampleBasic/src/sample_kernels.cu

Lines changed: 28 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -1,24 +1,22 @@
1-
#include <device_launch_parameters.h>
2-
#include <vector_functions.h>
3-
#include "math_constants.h"
4-
#include "log.h"
1+
#include "sample_kernels.cuh"
52

6-
__global__ void writeTex(cudaSurfaceObject_t surf, int width, int height, float time) {
3+
__global__ void writeTex(cudaSurfaceObject_t surf, int width, int height,
4+
float time)
5+
{
76
const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
87
const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
98

10-
if (x < width && y < height) {
9+
if (x < width && y < height)
10+
{
1111

12-
float4 t =
13-
make_float4(0, abs(cos(time)),
14-
0, 1.0f);
12+
float4 t = make_float4(0, abs(cos(time)), 0, 1.0f);
1513

1614
surf2Dwrite(t, surf, sizeof(float4) * x, y);
1715
}
1816
}
1917

20-
__global__ void writeTexArray(cudaSurfaceObject_t surf, int width, int height, int depth,
21-
float time)
18+
__global__ void writeTexArray(cudaSurfaceObject_t surf, int width, int height,
19+
int depth, float time)
2220
{
2321
const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
2422
const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -27,46 +25,45 @@ __global__ void writeTexArray(cudaSurfaceObject_t surf, int width, int height, i
2725
if (x < width && y < height)
2826
{
2927

30-
float4 t =
31-
make_float4(z%2, abs(cos(time)),
32-
0, 1.0f);
28+
float4 t = make_float4(z % 2, abs(cos(time)), 0, 1.0f);
3329

3430
surf3Dwrite(t, surf, sizeof(float4) * x, y, z);
3531
}
3632
}
3733

38-
39-
40-
__global__ void writeVertexBuffer(float4* pos, int size, float time)
34+
__global__ void writeVertexBuffer(float4 *pos, int size, float time)
4135
{
4236
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
43-
37+
4438
// write output vertex
4539
if (x < size)
4640
{
47-
pos[x] = make_float4( cos(2* CUDART_PI_F *time/x), sin(2* CUDART_PI_F *time/x), 0.0f,1.0f);
41+
pos[x] = make_float4(cos(2 * CUDART_PI_F * time / x),
42+
sin(2 * CUDART_PI_F * time / x), 0.0f, 1.0f);
4843
}
4944
}
5045

51-
52-
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock, cudaSurfaceObject_t inputSurfaceObj, const float time, const int width, const int height)
46+
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock,
47+
cudaSurfaceObject_t inputSurfaceObj,
48+
const float time, const int width,
49+
const int height)
5350
{
54-
writeTex << <dimGrid, dimBlock >> > (inputSurfaceObj, width, height, time);
55-
51+
writeTex<<<dimGrid, dimBlock>>>(inputSurfaceObj, width, height, time);
5652
}
5753

58-
5954
void kernelCallerWriteTextureArray(const dim3 dimGrid, const dim3 dimBlock,
60-
cudaSurfaceObject_t inputSurfaceObj,
61-
const float time, const int width,
62-
const int height, const int depth)
55+
cudaSurfaceObject_t inputSurfaceObj,
56+
const float time, const int width,
57+
const int height, const int depth)
6358
{
64-
// writeTexArray<<<dimGrid, dimBlock>>>(inputSurfaceObj, width, height,depth, time);
59+
// writeTexArray<<<dimGrid, dimBlock>>>(inputSurfaceObj, width,
60+
// height,depth, time);
6561
writeTex<<<dimGrid, dimBlock>>>(inputSurfaceObj, width, height, time);
6662
}
6763

68-
69-
void kernelCallerWriteBuffer(const dim3 dimGrid, const dim3 dimBlock, float4* vertexPtr,const int size, const float time)
64+
void kernelCallerWriteBuffer(const dim3 dimGrid, const dim3 dimBlock,
65+
float4 *vertexPtr, const int size,
66+
const float time)
7067
{
71-
writeVertexBuffer << <dimGrid, dimBlock >> > (vertexPtr, size, time);
68+
writeVertexBuffer<<<dimGrid, dimBlock>>>(vertexPtr, size, time);
7269
}

‎Plugin/premake5.lua

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -167,8 +167,10 @@ project "SampleBasic"
167167
SourceThirdPartyGLDir .. "**.c",
168168
SourceDir .. "**.cpp",
169169
IncludeDir .. "**.h",
170+
IncludeDir .. "**.cuh",
170171
IncludeDir .. "**.hpp",
171172
IncludeSubDir .. "**.h",
173+
IncludeSubDir .. "**.cuh",
172174
IncludeSubDir .. "**.hpp"
173175
}
174176

0 commit comments

Comments
 (0)