Skip to content

Commit 7f87ef7

Browse files
author
david.algis
committed
✨ Add comment code to improve performances with textures
1 parent 84b6de3 commit 7f87ef7

File tree

11 files changed

+202
-109
lines changed

11 files changed

+202
-109
lines changed

‎InteropUnityCUDA/Assets/Plugin/PluginInteropUnityCUDA.dll

-1 KB
Binary file not shown.

‎InteropUnityCUDA/Assets/Plugin/PluginInteropUnityCUDA.exp

438 Bytes
Binary file not shown.

‎InteropUnityCUDA/Assets/Plugin/PluginInteropUnityCUDA.lib

756 Bytes
Binary file not shown.

‎InteropUnityCUDA/Assets/Plugin/SampleBasic.dll

83.5 KB
Binary file not shown.

‎InteropUnityCUDA/Assets/Plugin/Utilities.dll

0 Bytes
Binary file not shown.

‎Plugin/PluginInteropUnityCUDA/src/Texture/texture.cpp

Lines changed: 45 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -12,17 +12,11 @@ Texture::Texture(void* textureHandle, int textureWidth, int textureHeight, int t
1212
// set a default size of grid and block to avoid calculating it each time
1313
// TODO : update this for texture depth
1414
_dimBlock = { 8, 8, 1 };
15-
_dimGrid = calculateDimGrid(
16-
_dimBlock, {(unsigned int)textureWidth, (unsigned int)textureHeight, 1},
15+
_dimGrid = calculateDimGrid(_dimBlock,
16+
{(unsigned int)textureWidth,
17+
(unsigned int)textureHeight,
18+
(unsigned int)_textureDepth},
1719
false);
18-
dim3 dimGrid = { (textureWidth + _dimBlock.x - 1) / _dimBlock.x,
19-
(textureHeight + _dimBlock.y - 1) / _dimBlock.y, 1};
20-
Log::log().debugLog("(" + std::to_string(dimGrid.x) + ", " +
21-
std::to_string(dimGrid.y) + ", " +
22-
std::to_string(dimGrid.z) + ")");
23-
Log::log().debugLog("(" + std::to_string(_dimGrid.x) + ", " +
24-
std::to_string(_dimGrid.y) + ", " +
25-
std::to_string(_dimGrid.z) + ")");
2620
_pGraphicsResource = nullptr;
2721
}
2822

@@ -35,7 +29,7 @@ Texture::Texture(void* textureHandle, int textureWidth, int textureHeight, int t
3529
cudaSurfaceObject_t Texture::mapTextureToSurfaceObject(int indexInArray)
3630
{
3731
// map the resource to cuda
38-
cudaGraphicsMapResources(1, &_pGraphicsResource);
32+
CUDA_CHECK(cudaGraphicsMapResources(1, &_pGraphicsResource));
3933
// cuda array on which the resources will be sended
4034
cudaArray* arrayPtr;
4135
//https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__INTEROP.html#group__CUDART__INTEROP_1g0dd6b5f024dfdcff5c28a08ef9958031
@@ -48,17 +42,57 @@ cudaSurfaceObject_t Texture::mapTextureToSurfaceObject(int indexInArray)
4842
resDesc.res.array.array = arrayPtr;
4943
cudaSurfaceObject_t inputSurfObj = 0;
5044
CUDA_CHECK(cudaCreateSurfaceObject(&inputSurfObj, &resDesc));
45+
CUDA_CHECK(cudaGetLastError());
5146
return inputSurfObj;
5247
}
5348

5449

50+
cudaTextureObject_t Texture::mapTextureToTextureObject(int indexInArray)
51+
{
52+
// map the resource to cuda
53+
CUDA_CHECK(cudaGraphicsMapResources(1, &_pGraphicsResource));
54+
// cuda array on which the resources will be sended
55+
cudaArray *arrayPtr;
56+
// https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__INTEROP.html#group__CUDART__INTEROP_1g0dd6b5f024dfdcff5c28a08ef9958031
57+
CUDA_CHECK(cudaGraphicsSubResourceGetMappedArray(
58+
&arrayPtr, _pGraphicsResource, indexInArray, 0));
59+
60+
// Wrap the cudaArray in a surface object
61+
cudaResourceDesc resDesc;
62+
memset(&resDesc, 0, sizeof(resDesc));
63+
resDesc.resType = cudaResourceTypeArray;
64+
resDesc.res.array.array = arrayPtr;
65+
66+
struct cudaTextureDesc texDesc;
67+
memset(&texDesc, 0, sizeof(texDesc));
68+
texDesc.addressMode[0] = cudaAddressModeWrap;
69+
texDesc.addressMode[1] = cudaAddressModeWrap;
70+
texDesc.filterMode = cudaFilterModeLinear;
71+
texDesc.readMode = cudaReadModeElementType;
72+
texDesc.normalizedCoords = 1;
73+
74+
cudaTextureObject_t texObj = 0;
75+
CUDA_CHECK(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL));
76+
CUDA_CHECK(cudaGetLastError());
77+
return texObj;
78+
}
79+
80+
81+
5582
void Texture::unMapTextureToSurfaceObject(cudaSurfaceObject_t& inputSurfObj)
5683
{
5784
CUDA_CHECK(cudaGraphicsUnmapResources(1, &_pGraphicsResource));
5885
CUDA_CHECK(cudaDestroySurfaceObject(inputSurfObj));
5986
CUDA_CHECK(cudaGetLastError());
6087
}
6188

89+
void Texture::unMapTextureToTextureObject(cudaTextureObject_t &texObj)
90+
{
91+
CUDA_CHECK(cudaGraphicsUnmapResources(1, &_pGraphicsResource));
92+
CUDA_CHECK(cudaDestroyTextureObject(texObj));
93+
CUDA_CHECK(cudaGetLastError());
94+
}
95+
6296
dim3 Texture::getDimBlock() const
6397
{
6498
return _dimBlock;

‎Plugin/SampleBasic/include/action_sample_texture.h

Lines changed: 18 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1,23 +1,29 @@
11
#pragma once
22
#include "action.h"
3+
#include "cuda_include.h"
34

45
class Texture;
56

6-
namespace SampleBasic {
7-
class ActionSampleTexture : public Action {
8-
public:
9-
ActionSampleTexture(void* texturePtr, int width, int height);
7+
namespace SampleBasic
8+
{
9+
class ActionSampleTexture : public Action
10+
{
11+
public:
12+
ActionSampleTexture(void *texturePtr, int width, int height);
1013

11-
inline int Start() override;
12-
inline int Update() override;
13-
inline int OnDestroy() override;
14+
inline int Start() override;
15+
inline int Update() override;
16+
inline int OnDestroy() override;
1417

15-
private:
16-
Texture* _texture;
17-
};
18+
private:
19+
Texture *_texture;
20+
cudaSurfaceObject_t _surf;
21+
};
1822
} // namespace SampleBasic
1923

20-
extern "C" {
24+
extern "C"
25+
{
2126

22-
UNITY_INTERFACE_EXPORT SampleBasic::ActionSampleTexture* UNITY_INTERFACE_API createActionSampleTextureBasic(void* texturePtr, int width, int height);
27+
UNITY_INTERFACE_EXPORT SampleBasic::ActionSampleTexture *UNITY_INTERFACE_API
28+
createActionSampleTextureBasic(void *texturePtr, int width, int height);
2329
}

‎Plugin/SampleBasic/include/action_sample_texture_array.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,19 +1,20 @@
11
#pragma once
22
#include "action.h"
3+
#include "cuda_include.h"
34

45
class Texture;
56

67
namespace SampleBasic {
78
class ActionSampleTextureArray : public Action {
89
public:
910
ActionSampleTextureArray(void* texturePtr, int width, int height, int depth);
10-
1111
inline int Start() override;
1212
inline int Update() override;
1313
inline int OnDestroy() override;
1414

1515
private:
1616
Texture* _texture;
17+
cudaSurfaceObject_t _surf;
1718
};
1819
} // namespace SampleBasic
1920

‎Plugin/SampleBasic/src/action_sample_texture.cpp

Lines changed: 46 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -1,49 +1,55 @@
11
#pragma once
22
#include "action_sample_texture.h"
3-
#include "unity_plugin.h"
43
#include "texture.h"
4+
#include "unity_plugin.h"
55

6-
7-
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock, cudaSurfaceObject_t inputSurfaceObj, const float t, const int width, const int height);
8-
9-
namespace SampleBasic {
10-
11-
ActionSampleTexture::ActionSampleTexture(void* texturePtr, int width, int height) : Action()
12-
{
13-
_texture = CreateTextureInterop(texturePtr, width, height, 1);
14-
}
15-
16-
17-
inline int ActionSampleTexture::Start()
18-
{
19-
_texture->registerTextureInCUDA();
20-
return 0;
21-
}
22-
23-
int ActionSampleTexture::Update()
24-
{
25-
26-
cudaSurfaceObject_t surf = _texture->mapTextureToSurfaceObject();
27-
kernelCallerWriteTexture(_texture->getDimGrid(), _texture->getDimBlock(), surf, GetTime(), _texture->getWidth(), _texture->getHeight());
28-
cudaDeviceSynchronize();
29-
_texture->unMapTextureToSurfaceObject(surf);
30-
return 0;
31-
}
32-
33-
inline int ActionSampleTexture::OnDestroy()
34-
{
35-
_texture->unRegisterTextureInCUDA();
36-
return 0;
37-
}
6+
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock,
7+
cudaSurfaceObject_t inputSurfaceObj,
8+
const float t, const int width, const int height);
9+
10+
namespace SampleBasic
11+
{
12+
13+
ActionSampleTexture::ActionSampleTexture(void *texturePtr, int width,
14+
int height)
15+
: Action()
16+
{
17+
_texture = CreateTextureInterop(texturePtr, width, height, 1);
18+
_surf = 0;
19+
}
20+
21+
inline int ActionSampleTexture::Start()
22+
{
23+
_texture->registerTextureInCUDA();
24+
_surf = _texture->mapTextureToSurfaceObject();
25+
return 0;
26+
}
27+
28+
int ActionSampleTexture::Update()
29+
{
30+
kernelCallerWriteTexture(_texture->getDimGrid(), _texture->getDimBlock(),
31+
_surf, GetTime(), _texture->getWidth(),
32+
_texture->getHeight());
33+
cudaDeviceSynchronize();
34+
return 0;
35+
}
36+
37+
inline int ActionSampleTexture::OnDestroy()
38+
{
39+
_texture->unMapTextureToSurfaceObject(_surf);
40+
_texture->unRegisterTextureInCUDA();
41+
return 0;
42+
}
3843

3944
} // namespace SampleBasic
4045

46+
extern "C"
47+
{
4148

42-
extern "C" {
43-
44-
UNITY_INTERFACE_EXPORT SampleBasic::ActionSampleTexture* UNITY_INTERFACE_API createActionSampleTextureBasic(void* texturePtr, int width, int height)
45-
{
46-
return (new SampleBasic::ActionSampleTexture(texturePtr, width, height));
47-
}
49+
UNITY_INTERFACE_EXPORT SampleBasic::ActionSampleTexture *UNITY_INTERFACE_API
50+
createActionSampleTextureBasic(void *texturePtr, int width, int height)
51+
{
52+
return (
53+
new SampleBasic::ActionSampleTexture(texturePtr, width, height));
54+
}
4855
} // extern "C"
49-

‎Plugin/SampleBasic/src/action_sample_texture_array.cpp

Lines changed: 55 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -1,52 +1,71 @@
11
#pragma once
22
#include "action_sample_texture_array.h"
3-
#include "unity_plugin.h"
43
#include "texture.h"
4+
#include "unity_plugin.h"
55

6-
void kernelCallerWriteTexture(const dim3 dimGrid, const dim3 dimBlock, cudaSurfaceObject_t inputSurfaceObj, const float t, const int width, const int height);
7-
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);
813

9-
namespace SampleBasic {
14+
namespace SampleBasic
15+
{
1016

11-
ActionSampleTextureArray::ActionSampleTextureArray(void* texturePtr, int width, int height, int depth) : Action()
12-
{
13-
_texture = CreateTextureInterop(texturePtr, width, height, depth);
14-
}
17+
ActionSampleTextureArray::ActionSampleTextureArray(void *texturePtr, int width,
18+
int height, int depth)
19+
: Action()
20+
{
21+
_texture = CreateTextureInterop(texturePtr, width, height, depth);
22+
_surf = 0;
23+
}
1524

25+
inline int ActionSampleTextureArray::Start()
26+
{
27+
_texture->registerTextureInCUDA();
28+
//_surf = _texture->mapTextureToSurfaceObject();
29+
return 0;
30+
}
1631

17-
inline int ActionSampleTextureArray::Start()
18-
{
19-
_texture->registerTextureInCUDA();
20-
return 0;
21-
}
32+
int ActionSampleTextureArray::Update()
33+
{
34+
/* kernelCallerWriteTextureArray(
35+
_texture->getDimGrid(), _texture->getDimBlock(),
36+
_surf, GetTime(), _texture->getWidth(),
37+
_texture->getHeight(), _texture->getDepth());*/
2238

23-
int ActionSampleTextureArray::Update()
24-
{
25-
for (int i = 0; i < _texture->getDepth(); i++)
26-
{
27-
cudaSurfaceObject_t surf = _texture->mapTextureToSurfaceObject(i);
28-
kernelCallerWriteTexture(_texture->getDimGrid(), _texture->getDimBlock(), surf, GetTime()+2*i, _texture->getWidth(), _texture->getHeight());
29-
_texture->unMapTextureToSurfaceObject(surf);
30-
}
39+
for (int i = 0; i < _texture->getDepth(); i++)
40+
{
41+
cudaSurfaceObject_t surf = _texture->mapTextureToSurfaceObject(i);
42+
kernelCallerWriteTexture(_texture->getDimGrid(),
43+
_texture->getDimBlock(), surf, GetTime()+2*i, _texture->getWidth(),
44+
_texture->getHeight()); _texture->unMapTextureToSurfaceObject(surf);
45+
}
3146

32-
cudaDeviceSynchronize();
33-
return 0;
34-
}
47+
cudaDeviceSynchronize();
48+
return 0;
49+
}
3550

36-
inline int ActionSampleTextureArray::OnDestroy()
37-
{
38-
_texture->unRegisterTextureInCUDA();
39-
return 0;
40-
}
51+
inline int ActionSampleTextureArray::OnDestroy()
52+
{
53+
//_texture->unMapTextureToSurfaceObject(_surf);
54+
_texture->unRegisterTextureInCUDA();
55+
return 0;
56+
}
4157

4258
} // namespace SampleBasic
4359

60+
extern "C"
61+
{
4462

45-
extern "C" {
46-
47-
UNITY_INTERFACE_EXPORT SampleBasic::ActionSampleTextureArray* UNITY_INTERFACE_API createActionSampleTextureArrayBasic(void* texturePtr, int width, int height, int depth)
48-
{
49-
return (new SampleBasic::ActionSampleTextureArray(texturePtr, width, height, depth));
50-
}
63+
UNITY_INTERFACE_EXPORT SampleBasic::ActionSampleTextureArray
64+
*UNITY_INTERFACE_API
65+
createActionSampleTextureArrayBasic(void *texturePtr, int width,
66+
int height, int depth)
67+
{
68+
return (new SampleBasic::ActionSampleTextureArray(texturePtr, width,
69+
height, depth));
70+
}
5171
} // extern "C"
52-

0 commit comments

Comments
 (0)